CUDA缩减优化示例

我正试图在CUDA Reduction上实现优化，并且要成功完成，直到第6部分。感谢大家的帮助。为了获得CUDA的完整感受，我还需要完成最终优化，如幻灯片＃31中所述，称为算法级联。CUDA缩减优化示例

这个想法本质上是每个线程有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'的重要性。访问内存的适当方式是什么，以便我们获得更好的性能？这是一个跨越访问的例子吗？

如果您有任何其他问题，请提前在下面提供帮助和评论。

来源

2016-05-17 Rahul Gunjan

所有这些缩减代码的完整工作示例在相应的[CUDA示例代码]（http://docs.nvidia.com/cuda/cuda-samples/index.html#cuda-parallel-reduction）中提供。你不应该猜测任何参数。我怀疑你提供了足够的信息来解释你的观察。如果您的第6部分的实施表现不佳，您可能应该运行CUDA示例代码并研究差异。 –

这是一个合并访问的例子。最好的gridDim取决于你的硬件。根据每个线程的寄存器和每个块的最大线程数，该值应该是硬件上可用多处理器数量的某个乘数。如果你的问题足够大，8倍的多处理器计数对于开普勒来说是个好选择，而对于麦克斯韦来说是16倍。

来源

2016-05-17 06:04:53

假设你有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内存的最低效的方式。

来源

2016-05-17 08:50:37 kangshiyin

@RobertCrovella我认为代码是Nvidia演示，而不是Rahul的方式。这一个是：'我尝试了一种方法，我的每个线程都从内存中访问连续的512个数字。' – kangshiyin

对不起，我误解了你的答案。删除了我以前的评论。 –

CUDA缩减优化示例

回答

相关问题