【发布时间】：2015-09-15 21:12:41
【问题描述】：
我是 CUDA 的新手，目前我正在研究减和样本，这与我的最终目标相关。
提供的文档描述了如何优化内核以快速减少跨块的大型数组。 reduction_kernel.cu 中的宿主函数reduce 使用模板在编译时优化各种内核。
template <class T>
void reduce(int size, int threads, int blocks,
int whichKernel, T *d_idata, T *d_odata)
{
//
// Long list with switch statement to have all optimized functions at compile-time
//
// amongst which (for instance):
case 32:
reduce5<T, 32><<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
break;
编辑：内核reduce5 用d_idata 的部分和填充d_odata。更具体地说，它将g_idata 的元素与索引2*blockSize*blockIdx.x 相加，直到2*blockSize*(blockIdx.x + 1)（不包括在内）并将结果存储在g_odata[blockIdx.x] 中。 （编辑结束）
总和是通过跨块减少直到剩下一个块来获得的。主机代码用于通过在缩减阵列上重复启动内核来跨“级别”同步内核。 reduction.cpp中的相关代码：
template <class T>
T benchmarkReduce(int n, numThreads, numBlocks, /* more args */,
T *h_odata, T *d_idata, T *d_odata) {
// first kernel launch
reduce<T>(n, numThreads, numBlocks, whichKernel, d_idata, d_odata);
// repeated kernel launches
int s=numBlocks;
int kernel = whichKernel;
while (s > cpuFinalThreshold)
{
int threads = 0, blocks = 0;
getNumBlocksAndThreads(kernel, s, maxBlocks, maxThreads, blocks, threads);
reduce<T>(s, threads, blocks, kernel, d_odata, d_odata);
if (kernel < 3)
s = (s + threads - 1) / threads;
else
s = (s + (threads*2-1)) / (threads*2);
}
}
我对第一个内核调用很满意，它将d_idata 的部分和存储在d_odata 中。我担心第二次内核启动（在 while 循环内）：即，内核将对d_odata进行读写，这可能导致数据竞争。
例如，第二个块可以在第一个块读取其原始值之前将其部分和写入d_odata[1]；这是第一个块的部分总和所必需的。
我错过了一个细节吗？
【问题讨论】：
-
这取决于您没有提供的归约内核的实现。可以在没有数据竞争的情况下就地减少。
-
是的，这是可能的，但在示例代码中不会发生。例如，
reduce5（由reduce包裹）读取d_idata[blockIdx.x*(blockSize*2) + threadIdx.x]和d_idata[blockIdx.x*(blockSize*2) + threadIdx.x + blockSize]并存储在d_odata[blockIdx.x]。 -
正在讨论的代码是cuda sample reduction code。 OP 没有明确说明这一点。
-
我相信这是一个潜在的竞争条件。我认为您没有遗漏任何细节。在实践中没有观察到会导致问题，因为： 1. 对于某些选定的缩减内核，此时只会启动 1 个块。 2. 对于其他一些配置，此时的所有块可以同时调度，在某些 GPU 上， 3. 默认块调度顺序倾向于先调度编号较低的
blockIdx.x，然后再安排编号较高的blockIdx.x我已经提交了一个错误NVIDIA 调查并可能改进代码。我目前没有任何进一步的细节。