在 CUB 类之后使用 cudaDeviceSynchronize

Posted 2023-02-16

【中文标题】在 CUB 类之后使用 cudaDeviceSynchronize

【英文标题】：Using cudaDeviceSynchronize after a CUB class

【发布时间】：2014-08-28 04:39:07

【问题描述】：

在从 CUDA 内核调用 CUB 类后是否需要调用 cudaDeviceSynchronize？当从设备中使用说 DeviceReduce::Sum() 时，有隐式内存副本会阻止设备继续运行，但是在使用 GPU 上调用的以下代码后遇到了一些不稳定：

__device__ void calcMonomerFlux(double* fluxes, double* lengths, double* dt) //temp2 temp1


    void *d_temp_storage = NULL;
    size_t temp_storage_bytes = 0;

    arrInitToLengths<<< numBlocks, numThreads >>>(lengths); 
    cudaDeviceSynchronize();
    arrMult<<< numBlocks, numThreads >>>(fluxes, lengths, lengths);
    cudaDeviceSynchronize();
    double sum = 0;

    cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, lengths, lengths, maxlength);
    //cudaDeviceSynchronize();

    cudaMalloc(&d_temp_storage, temp_storage_bytes);
    //cudaDeviceSynchronize();

    cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, lengths, lengths, maxlength);
    //cudaDeviceSynchronize();

    cudaFree(d_temp_storage);

【参考方案1】：

是的，每次 CUB 调用后都需要 cudaDeviceSynchronize()。请注意注释的同步调用在问题中的位置。我花了很多时间追查为什么我的总和计算不正确甚至不一致。最终，我在通过 NSIGHT 调试器前进时发现，只有当我在每个 CUB 函数之后放置断点时，计算才会正确。

【讨论】：

CUB，当从设备代码调用时，正在使用 CUDA 动态并行，即它正在生成子内核。与任何内核启动一样，这些内核启动与调用线程是异步的，并且不能保证在将控制权返回给调用线程时完成。因此，如果调用线程要求CUB调用产生的数据完整且准备就绪，则需要同步并等待子内核完成。这个想法通常适用于子内核通过 CUDA 动态并行性生成的数据。