并行还原

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

优化代码：随着2个负荷和减少1ST ADD：

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

它基本上执行在下面的图片中的运行表明：

此代码基本上是“说”的线程中有一半将性能从全局内存和写入共享内存，如阅读在图片中显示。

您执行一个Kernell，现在您想要减少一些值，您将以上代码的访问权限限制为仅运行线程总数的一半。想象您有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中，所以现在有一半的线程，并且共享内存中的数组是两倍。在第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]。