问CUDA缩减-和样本:数据竞赛？

EN

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;

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);  
    }
}