参考
首先列出参考文献：
代码部分：会了么的个人空间-会了么个人主页-哔哩哔哩视频 (bilibili.com)
图片及部分理解部分：Cuda C编程权威指南1.并行规约分化+循环展开-CSDN博客
背景
cuda的执行模型
逻辑概念上，grid>block>thread
也就是，网格>线程块>线程
每个核函数的启动都对应着一个grid，grid中的所有block共享全局内存，每个block又是由许多线程构成的（block内的线程共享共享内存）。
CUDA中SIMD的基本单位是一个warp（线程束一般是由32个线程组成，共享寄存器）。
threadIdx.x：当前线程在线程块中的索引
blockIdx.x：当前线程块的索引
blockDim.x：每个block的线程数
一般在main函数中用下面代码定义：
// initialization int size = 1 << 24; // total number of elements to reduce printf(" with array size %d ", size); // execution configuration int blocksize = 1024; // initial block sizeif(argc > 1) { blocksize = atoi(argv[1]); // block size from command line argument } dim3 block (blocksize, 1); dim3 grid ((size + block.x - 1) / block.x, 1); printf("grid %d block %d\n", grid.x, block.x);
以上代码每个block的thread数初始值是1024，但是可以在终端设置大于1的数。
此外，
tid：当前thread在这个block中的id索引
idx：当前thread在全局中的索引
ideta：每一个block计算的部分data的数据指针起点；进行了循环展开后相当于间隔八个data块处理一次
在核函数定义：
//set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockDim.x*blockIdx.x*8+threadIdx.x;
//boundary check
if (tid >= n) return;
//convert global data pointer to the
int *idata = g_idata + blockIdx.x*blockDim.x*8;相邻配对和交错配对
提示：有的也叫邻域并行计算和间域并行计算
邻域：
// Neighbored Pair Implementation with divergence
__global__ void reduceNeighbored (int *g_idata, int *g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if (idx >= n) return;
// in-place reduction in global memory
for (int stride = 1; stride < blockDim.x; stride *= 2)
{
if ((tid % (2 * stride)) == 0)
{
idata[tid] += idata[tid + stride];
}
// synchronize within threadblock
__syncthreads();
}间域：
__global__ void reduceInterleaved (int *g_idata, int *g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if(idx >= n) return;
// in-place reduction in global memory
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
{
if (tid < stride)
{
idata[tid] += idata[tid + stride];
}
__syncthreads();
}上面两种是基础情况，比较容易理解（尤其看着图）
循环展开
这点东西我看了一整天还是晕乎乎
由于核函数中的每个线程块都只处理对应数据块的数据内容，那么可以通过代码让一个block处理多个data块。
也就是在数据求和前，使用每个thread做一次加法，从多个（示例是8个）data块中取数据
最后再归约相加
循环展开，先贴个代码：（8倍block）
__global__ void reduceUnrolling8 (int *g_idata, int *g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;
// convert global data pointer to the local pointer of this block
int *idata = g_idata + blockIdx.x * blockDim.x * 8;
// unrolling 8
if (idx + 7 * blockDim.x < n)
{
int a1 = g_idata[idx];
int a2 = g_idata[idx + blockDim.x];
int a3 = g_idata[idx + 2 * blockDim.x];
int a4 = g_idata[idx + 3 * blockDim.x];
int b1 = g_idata[idx + 4 * blockDim.x];
int b2 = g_idata[idx + 5 * blockDim.x];
int b3 = g_idata[idx + 6 * blockDim.x];
int b4 = g_idata[idx + 7 * blockDim.x];
g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;
}
__syncthreads();
示意图长这样：
一个block处理8个data的数据，比如256个线程处理512*8个数据（用一个block将8个data求和成一个data）
然后对每个block里的data的求和操作：（示意图如下）
①首先是最简单的情况，确定了blocksize是512，不需要再写条件语句
//in-place reduction in global memory
for (int stride = blockDim.x/2; stride>32; stride >>=1)
{
if (tid <stride)
{
idata[tid] += idata[tid + stride];
}
//synchronize within block
__syncthreads();
}
//write result for this block to global mem
这里注意三种情况都需要用__syncthreads();来同步指令，这是因为每一层的操作都是并行的，不同速度的thread要进行等待来进行block内的同步，防止出现层之间的交叉导致逻辑不合理。
②如果想在终端传入参数
此时可使用blockDim书写条件语句
//in-place reduction in global memory
if(blockDim.x>=1024 && tid <512)
idata[tid]+=idata[tid+512];
__syncthreads();
if(blockDim.x>=512 && tid <256)
idata[tid]+=idata[tid+256];
__syncthreads();
if(blockDim.x>=256 && tid <128)
idata[tid]+=idata[tid+128];
__syncthreads();
if(blockDim.x>=128 && tid <64)
idata[tid]+=idata[tid+64];
__syncthreads();③也可以使用模版函数
在核函数前加一条语句
template <unsigned int iBlockSize>
在核函数中加入，根据Switch的传入的值，分别计算idata的值
//in-place reduction in global memory
if(iBlockSize>=1024 && tid <512)
idata[tid]+=idata[tid+512];
__syncthreads();
if(iBlockSize>=512 && tid <256)
idata[tid]+=idata[tid+256];
__syncthreads();
if(iBlockSize>=256 && tid <128)
idata[tid]+=idata[tid+128];
__syncthreads();
if(iBlockSize>=128 && tid <64)
idata[tid]+=idata[tid+64];
__syncthreads();在main函数里需要加一个Switch函数，根据在终端传入的size参数值分别处理
//kernel 4:reduceCompleteUnroll
CHECK(cudaMemcpy(idata_dev, idata_host, bytes, cudaMemcpyHostToDevice));
CHECK(cudaDeviceSynchronize());
iStart = cpuSecond();
switch(blocksize)
{
case 1024:
reduceCompleteUnroll <1024><< <grid.x/8, block >> >(idata_dev, odata_dev, size);
break;
case 512:
reduceCompleteUnroll <512><< <grid.x/8, block >> >(idata_dev, odata_dev, size);
break;
case 256:
reduceCompleteUnroll <256><< <grid.x/8, block >> >(idata_dev, odata_dev, size);
break;
case 128:
reduceCompleteUnroll <128><< <grid.x/8, block >> >(idata_dev, odata_dev, size);
break;
}----------------
以上第2、3种情况，都需要对[0,32]之间的thread进行代码展开
if(tid < 32) {
volatile int *vmem = idata;
vmem[tid] *= vmem[tid + 32];
vmem[tid] *= vmem[tid + 16];
vmem[tid] *= vmem[tid + 8];
vmem[tid] *= vmem[tid + 4];
vmem[tid] *= vmem[tid + 2];
vmem[tid] *= vmem[tid + 1];
}
从64个数合并到一个数，首先将前32个数，按照步长为32，进行并行加法，前32个tid得到64个数字的两两和，存在前32个数字中
然后这32个数加上步长为16的变量，得到16个数，这16个数的和就是最后这个块的归约结果
当执行 tid+32的时候，这32个线程都在执行这步，而不会有任何本线程束内的线程会进行到下一句
使用volatile修饰data就是为了防止编译器优化全局内存或共享内存的读取，防止一层求和完成后下一层求和会出现不同的thread读同一位置的情况。
OK结束。
7507
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
