我是 CUDA 的新手,目前我正在研究与我的最终目标相关的总和减少样本。
提供的文档描述了如何优化内核以快速减少跨 block 的大型数组。 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]
中。 (编辑结束)
总和是通过跨 block 减少直到剩下一个 block 来获得的。主机代码用于通过在简化阵列上重复启动内核来跨“级别”同步内核。 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
,这可能导致数据竞争。
例如,第二个 block 可以在第一个 block 读取其原始值之前将其部分和写入d_odata[1]
;这是第一个 block 的部分和所必需的。
我是否遗漏了一个细节?
最佳答案
这已在 CUDA 8.0 包中修复。 CUDA 8.0 应该很快就会可用。
关于c++ - CUDA 归约和样本 : data racing?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/31115128/