CUDA 缩减优化示例答案

【问题标题】：CUDA Reduction Optimization ExampleCUDA 缩减优化示例
【发布时间】：2016-05-17 03:02:42
【问题描述】：

我正在尝试在CUDA Reduction 实施优化，并成功完成到第 6 部分。感谢大家帮助我。为了全面了解 CUDA，我还需要完成最终优化，如幻灯片 #31 中所述，称为算法级联。

这个想法本质上是每个线程有 512 个元素，并在执行归约之前按顺序对所有元素求和。

我尝试了一种方法，我的每个线程都从内存中访问连续的 512 个数字。不幸的是，它的表现最差。我猜是银行冲突的原因，但仍然没有弄清楚。你们中的任何人都可以提出这种行为的原因吗？

我还将在下面发布 Nvidia 提供的示例代码。

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockSize*2) + threadIdx.x;
unsigned int gridSize = blockSize*2*gridDim.x;
sdata[tid] = 0;
while (i < n) {
    sdata[tid] += g_idata[i] + g_idata[i+blockSize];
    i += gridSize;
}
__syncthreads();

有几个参数未定义。我可以推断 blockSize 等于每个块的线程数。但我无法推断变量“gridSize”的重要性。什么是访问内存的适当方式以便我们获得更好的性能？这是跨步访问的示例吗？

提前感谢您的帮助，如果您有任何其他问题，请在下方发表评论。

【问题讨论】：

在相应的CUDA sample code 中提供了所有这些缩减代码的完整示例。您不必猜测任何参数。我怀疑您是否提供了足够的信息来解释您的观察。如果您对第 6 部分的实现表现不佳，也许您应该运行 CUDA 示例代码并研究差异。

标签： optimization cuda nvidia reduction

【解决方案1】：

假设您每个块有blockDim.x = blockSize = 256 线程，并且网格中有gridDim.x = 32 块，并且您想要减少一个大数组g_idata[8,192,000]。

那么你总共有 8192 个线程。让我们使用

thread[x][y], x=0..31, y=0..255

来表示这些线程。

每个thread[x][y] 正在加载

g_idata[iter*512*x+y] and g_idata[iter*512*x+256+y], iter = 0 .. 999

到共享内存sdata。

对于每次迭代 iter，所有 8192 个 threads[x][y] 将从 GPU 内存中加载 gridSize = 16384 元素。

这是合并的内存访问，是访问 GPU 内存的正确方式。

但是，每个thread[x] 读取data[i*x*512 .. i*(x+1)*512-1], i=0... 的方式都不是一个好方法。实际上这是访问 GPU 内存的最低效的方式。

【讨论】：

@RobertCrovella 我认为代码是 Nvidia 演示，而不是 Rahul 的方式。这个是：I tried with a approach, where each thread of mine was accessing continuous 512 numbers from memory.
对不起，我误解了你的回答。删除了我之前的评论。

【解决方案2】：

这是合并访问的一个示例。最好的 gridDim 取决于您的硬件。根据每个线程的寄存器和每个块的最大线程数，这个值应该是硬件上可用的多处理器数量的乘数。如果您的问题足够大，8 倍的多处理器数量对于 Kepler 来说是一个不错的选择，而对于 Maxwell 来说是 16 倍。

【讨论】：