平行减少答案

【问题标题】：Parallel Reduction平行减少
【发布时间】：2012-11-17 06:46:37
【问题描述】：

我阅读了 Mark Harris 的文章 Optimizing Parallel Reduction in CUDA，我发现它非常有用，但有时我仍然无法理解 1 或 2 个概念。它写在第 18 页：

//First add during load

// 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();

优化代码：有 2 次加载和第一次添加减少：

// perform first level of reduction,

// reading from global memory, writing to shared memory
unsigned int tid = threadIdx.x;                                    ...1

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;          ...2

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];                   ...3

__syncthreads();                                                   ...4

我无法理解第 2 行；如果我有 256 个元素，并且如果我选择 128 作为我的块大小，那么为什么我将它乘以 2？请解释如何确定块大小？

【问题讨论】：

标签： c++ c cuda parallel-processing gpu

【解决方案1】：

在优化的代码中，您运行内核的块大小是未优化实现的一半。

让我们将未优化代码中的块大小称为work，将此大小的一半称为unit，并让这些大小对于优化代码也具有相同的数值。

在未优化的代码中，您使用与work 一样多的线程运行内核，即blockDim = 2 * unit。每个块中的代码只是将g_idata 的一部分复制到共享内存中的一个数组，大小为2 * unit。

在优化代码blockDim = unit中，现在有1/2的线程，共享内存中的数组小了2倍。在第 3 行中，第一个和来自偶数单位，而第二个来自奇数单位。这样一来，减少所需的所有数据都被考虑在内。

示例：如果您使用blockDim=256=work（单块，unit=128）运行未优化的内核，则优化的代码具有blockDim=128=unit 的单块。由于这个块得到blockIdx=0，*2 无关紧要；第一个线程执行g_idata[0] + g_idata[0 + 128]。

如果您有 512 个元素，并且使用 2 个大小为 256 的块（work=256、unit=128）运行未优化，则优化后的代码有 2 个块，但现在大小为 128。第二个块中的第一个线程（ blockIdx=1) 确实是 g_idata[2*128] + g_idata[2*128+128]。

【讨论】：

@P Marecki：非常感谢。您的回答确实有助于我理解解决方案，但是如果您可以在第一段中让我知道，总元素是多少。如果是 256，那么 256 个元素将如何被单个块占用？第 2 段 512 个元素的相同问题，只有 2 个具有 128 个线程的块。
@robot ：g_idata 中的元素总数在第 1 段为 256，在第二段为 512。是的：在优化代码中sdata 比2x 更小（那里只有128 元素，或第二段中的2 * 128），但这足以减少。
@P Marecki：感谢您的回复。但是如果有 256 个元素，那么我们必须处理 256 个元素，如何将元素减少到 128 个元素？你的意思是有 256 个元素，我们有 1 个块大小为 128 的块来处理 256 个元素？
@P Marecki：不应该是这样：与非优化代码相比，优化代码中的块数应该减半。您已经为 512 个元素编写了该代码，最初有 256 个线程/块的块，而在优化代码中有 128 个线程/块；那么我们在哪里将区块数量减半？
@robot ：我想我理解了混乱的根源。在编辑后的答案中，现在有变量work 和unit，这两种实现的数值相同，而blockDim 在两者中不同。

【解决方案2】：

基本上就是执行下图所示的操作：

这段代码基本上说一半的线程将执行从全局内存读取和写入共享内存，如图所示。

你执行一个内核，现在你想减少一些值，你限制了对上面代码的访问，只有运行的线程总数的一半。假设您有 4 个块，每个块有 512 个线程，您将上面的代码限制为只能由前两个块执行，并且您有一个 g_idate[4*512]：

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;  

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];

所以：

thread 0 of block = 0  will copy the position 0 and 512,  
thread 1 of block = 0 position 1 and 513;
thread 511 of block = 0 position 511 and 1023;
thread 0 of block 1 position 1024 and 1536
thread 511 of block = 1 position 1535 and 2047

之所以使用blockDim.x*2，是因为每个线程都会访问i和i+blockDim.x的位置，所以你需要乘以2来保证下一个id块上的线程不会计算@的位置987654330@ 已经计算好了。

【讨论】：