c++ - 平行减少

Question

我读过 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？请解释如何确定块大小？

Answer 1

基本上，它正在执行下图所示的操作：

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

你执行了一个内核，现在你想减少一些值，你限制了对上面代码的访问，只有运行的线程总数的一半。假设您有 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块上的线程不会计算g_idata已计算的位置。

Answer 2

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

让我们在非优化代码中调用块的大小work，让这个大小的一半被调用unit，并且让这些大小对于优化代码也具有相同的数值。

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

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

示例：如果您使用blockDim=256=work(single block, 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]。