CUDA：Mark Harris 的并行缩减样本不就是对每个线程块求和吗？

CUDA: Isn't Mark Harris's parallel reduction sample just summing each thread block?

本文关键字：线程 求和 Harris Mark 并行 样本 CUDA      更新时间：2023-10-16

他的幻灯片链接：

http://docs.nvidia.com/cuda/samples/6_Advanced/reduction/doc/reduction.pdf

这是他为第一个版本的并行归约编写的代码：

__global__ void reduce0(int *g_idata, int *g_odata) {
extern __shared__ int sdata[];
// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_idata[i];
__syncthreads();
// do reduction in shared mem
for(unsigned int s=1; s < blockDim.x; s *= 2) {
if (tid % (2*s) == 0) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];
}

他后来对其进行了优化。这怎么不只是对每个线程块的所有int求和，并将答案放在另一个向量中？这就是它的本意吗？*g_odata本身不是一个向量吗？因为它将和放在向量中的每个"blockIdx.x"点上？如何将向量g_idata求和为一个数字？

这怎么不只是对每个线程块的所有int求和，并将答案放在另一个向量中？

它正是这样做的。

这就是它的本意吗？

是的。

g_odata本身不是一个向量吗？因为它把和放在向量中的每个"blockIdx.x"点上？

是的，它是包含块级和的向量。

如何将向量g_idata求和为一个数字？

调用内核两次。在原始数据集上一次，在上一次调用的矢量输出上一次（块级总和）。请注意，第二步只使用一个块，并且要求每个块可以启动足够的线程来覆盖整个向量，即上一步的每个和一个线程。如果您查看旨在伴随您链接的演示的cuda示例代码，您将发现这样的调用序列，例如在reduction.cpp的第304和333行。对reduce<T>的第二次调用执行对部分块和求和的归约，如第324行的注释所示：

304:reduce<T>(n, numThreads, numBlocks, whichKernel, d_idata, d_odata);
    // check if kernel execution generated an error
    getLastCudaError("Kernel execution failed");
    if (cpuFinalReduction)
    {
        // sum partial sums from each block on CPU
        // copy result from device to host
        checkCudaErrors(cudaMemcpy(h_odata, d_odata, numBlocks*sizeof(T), cudaMemcpyDeviceToHost));
        for (int i=0; i<numBlocks; i++)
        {
            gpu_result += h_odata[i];
        }
        needReadBack = false;
    }
    else
    {
324:    // sum partial block sums on GPU
        int s=numBlocks;
        int kernel = whichKernel;
        while (s > cpuFinalThreshold)
        {
            int threads = 0, blocks = 0;
            getNumBlocksAndThreads(kernel, s, maxBlocks, maxThreads, blocks, threads);
333:        reduce<T>(s, threads, blocks, kernel, d_odata, d_odata);

注意来自线304处的第一次缩减的输出CCD_ 4作为输入传递到线333上的第二次缩减。

还要注意，内核分解的必要性和这种方法在您在幻灯片3-5上链接的演示中有介绍。

相关文章：