并行归约

Question

我阅读了Mark Harris的《在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();

优化代码:在两个负载和第一个减少的加法中:

// 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

我无法理解第二行；如果我有256个元素，而我选择128作为块大小，那么为什么要乘以2？请说明如何确定块大小？

- robot

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位置。

- dreamcrash

Answer 2

在优化的代码中，块的大小比非优化实现中的块小一半。

让我们把非优化代码中块的大小称为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个元素，并且使用大小为256的2个块（work=256，unit=128）运行未经优化的代码，则经过优化的代码现在具有2个大小为128的块。第二个块中的第一个线程（blockIdx=1）执行g_idata[2*128] + g_idata[2*128+128]。

- P Marecki

---

@P Marecki: 非常感谢您。您的答案确实帮助我理解了解决方案，但如果您能让我知道第一段中总共有多少元素，那就太好了。如果是256个，那么如何用单个块占用256个元素？对于第二段512个元素和仅有2个块的128个线程的问题也是同样的。 - robot

@机器人: 在第一个段落中，g_idata 中元素的总数为256，在第二个段落中为512。正确：在优化的代码中，sdata 要小2倍（你只有128个元素，或者在第二个段落中是2 * 128），但这已足够用于减少的目的。 - P Marecki

@P Marecki：感谢您的回复。但是如果有256个元素，那么我们必须处理256个元素，如何将元素减少到128个？您的意思是有256个元素，我们使用块大小为128的1个块来处理256个元素吗？ - robot

@P Marecki：难道不应该是这样吗：在优化代码中，块的数量应该减半，与非优化代码相比。您写道对于512个元素，最初有256个线程/块的块，在优化代码中有128个线程/块；那么我们在哪里将块的数量减半呢？ - robot

@P Marecki：实际上写的是“减半块数”，而不是“减半块大小”。那么处理512个元素时，应该将块数从2块减少到1块，每块256个线程/块。 - robot