【CUDA 】第3章 CUDA执行模型——3.5循环展开（1）

待解决的问题：

第三章 CUDA执行模型

3.5 循环展开

循环展开是一种循环优化的技术，通过减少分支出现频率+循环维护指令。
循环主体代码被多次编写，任何封闭的循环可以把迭代次数减少或完全删除。

循环展开因子：循环体的复制数量
迭代次数=原始迭代次数/循环展开因子
在顺序数组中，如果循环迭代的次数在执行前已经知道，循环展开是最有效提升性能的方法。

举一个简单的例子：

for(int i = 0; i < 100; i++){
	a[i] = b[i] + c[i];
}

for(int i = 0; i < 100; i += 2){
	a[i] = b[i] + c[i];
	a[i+1] = b[i+1] + c[i+1];
}

下面的循环迭代次数是上面的一半，这是因为重复操作了一次循环体。
这种提升是由于编译器执行循环展开时低级指令的改进和优化。

目标：通过减少指令消耗和增加更多的独立调度指令来提高性能，更多的并发操作被添加到流水线上，产生更高的指令和内存带宽，帮线程束调度器提供更多符合条件的线程束，隐藏指令或内存延迟。

3.5.1 展开的规约

用一个线程块手动处理两个数据块，每个线程作用域多个数据块，并处理每个块的一个元素

//展开规约1：手动展开两个块block的处理
__global__ void reduceUnrolling2(int *g_idata, int *g_odata, unsigned int n){
    //设置线程id
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;//这里乘2

    //把全局数据指针转换为块内局部指针
    int *idata = g_idata + blockIdx.x * blockDim.x * 2;//这里乘2

    //展开两个数据块
    if(idx + blockDim.x < n) g_idata[idx] += g_idata[idx + blockDim.x];//每个线程都加一个相邻块的元素，可以算一个迭代，该循环可在数据块间规约
    __syncthreads();

    //在全局内存中就地规约
    for(int stride = blockDim.x/2; stride > 0; stride >>= 1){
        if(tid < stride){
            idata[tid] += idata[tid + stride];
        }

        __syncthreads();
    }

    //把这个块的结果写回全局内容
    if(tid == 0) g_odata[blockIdx.x] = idata[0];
}

可以利用线程块的空闲资源，原本一个线程块中，只有前半个线程块在工作，后半个线程块闲置（交错规约）。把两个数据块一起处理后，后半个线程块也能利用起来。

关键点：
①索引的计算

unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;
int *idata = g_idata + blockIdx.x * blockDim.x * 2;

②数据规约

if (idx + blockDim.x < n) 
    g_idata[idx] += g_idata[idx + blockDim.x];

第一个块的内容与第二个块相加，数据存回第一个块。

③块内交错规约

for (int stride = blockDim.x/2; stride > 0; stride >>= 1) {
    if (tid < stride) {
        idata[tid] += idata[tid + stride];
    }
    __syncthreads();
}

完成数据块内的规约，交错规约。

对应的main函数代码也需要修改：网格大小减半，在核函数调用参数、数据传回主机和最后求和方面，都需要除2

//kernel4:reduceUnrolling2 展开规约2：一个线程块处理两个数据块
    cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice);
    cudaDeviceSynchronize();
    iStart = seconds();
    //因为现在每个线程块处理两个数据块，需要调整内核的执行配置，网格大小减为原来的一半
    reduceUnrolling2<<<grid.x/2, block>>>(d_idata, d_odata, size);//这里写错成了bytes
    cudaDeviceSynchronize();
    iElaps = seconds() - iStart;
    cudaMemcpy(h_odata, d_odata, grid.x/2*sizeof(int), cudaMemcpyDeviceToHost);//需要除2
    gpu_sum = 0;
    for(int i = 0; i<grid.x/2; i++){//需要除2
        gpu_sum += h_odata[i];
    }
    printf("gpu reduceUnrolling2    elapsed %f ms gpu_sum: %d <<<grid %d block %d>>>\n", iElaps, gpu_sum, grid.x/2, block.x);//需要除2

执行结果：

./3-3reduceInteger2 starting reduction atdevice 0: NVIDIA GeForce RTX 3090 
   with array size 16777216 grid 32768 block 512
cpu reduce         elapsed 0.039772 ms cpu_sum: 2139353471
gpu Warmup          elapsed 0.003663 ms gpu_sum: 0 <<<grid 32768 block 512>>>
gpu Neighbored          elapsed 0.001532 ms gpu_sum: 2139353471 <<<grid 32768 block 512>>>
gpu NeighboredLess      elapsed 0.001402 ms gpu_sum: 2139353471 <<<grid 32768 block 512>>>
gpu reduceInterleaved    elapsed 0.001371 ms gpu_sum: 2139353471 <<<grid 32768 block 512>>>
gpu reduceUnrolling2    elapsed 0.001357 ms gpu_sum: 2139353471 <<<grid 16384 block 512>>>

原始的交错规约时间reduceInterleaved0.001371 ms，二合一的交错规约reduceUnrolling20.001357 ms。

原因：一个线程中有更多的独立内存加载/存储操作性能更好，因为内存延迟可以被更好隐藏，可以查看设备内存吞吐量指标，内存吞吐量提高了很多。