c++ - CUDA 归约和样本 : data racing?

Question

我是 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 的部分和所必需的。

我是否遗漏了一个细节？

Answer 1

这已在 CUDA 8.0 包中修复。 CUDA 8.0 应该很快就会可用。