Cuda 扫描 - 使用共享内存时，Debug 和 Release 模式下的不同结果答案

【问题标题】：Cuda scan - different results in Debug and Release modes when using shared memoryCuda 扫描 - 使用共享内存时，Debug 和 Release 模式下的不同结果
【发布时间】：2018-10-18 12:02:12
【问题描述】：

我正在尝试在 cuda 中编写分段扫描，其中段的长度等于经线的长度 (32)。这是我的内核：

__global__ void kernel(int totalSize, unsigned short* result)
{
    __shared__ unsigned short s_data[1024];

    const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int intraWarpThreadId = threadIdx.x & 31;

    if (tid >= totalSize)
        return;

    s_data[threadIdx.x] = result[tid];
    __syncthreads();

    IntraWarpScan(s_data, threadIdx.x, intraWarpThreadId);
    __syncthreads();

    result[tid] = s_data[threadIdx.x];
}

__device__ void IntraWarpScan(unsigned short* s_data, unsigned int intraBlockThreadId, unsigned int& intraWarpThreadId)
{
    if (intraWarpThreadId >= 1)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 1];

    if (intraWarpThreadId >= 2)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 2];

    if (intraWarpThreadId >= 4)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 4];

    if (intraWarpThreadId >= 8)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 8];

    if (intraWarpThreadId >= 16)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 16];
}

我想我在共享内存中有一些竞争条件，但我无法弄清楚它们为什么会发生。由于每个片段都在一个扭曲中扫描，我不需要在 IntraWarpScan 过程中进行任何同步，对吧？但是如果没有在 IntraWarpScan 中的每个 if 指令之后进行同步，我在 Release 构建中会得到错误的结果。在调试中，我得到了正确的结果。

另一方面，如果我决定不使用共享内存而只使用设备内存，那么我在两个构建中都会得到正确的结果，如下所示：

__global__ void kernel(int totalSize, unsigned short* result)
{
    const unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int intraWarpThreadId = threadIdx.x & 31;

    if (tid >= totalSize)
        return;

    IntraWarpScan(result, tid, intraWarpThreadId);
    __syncthreads();
}

__device__ void IntraWarpScan(unsigned short* s_data, unsigned int intraBlockThreadId, unsigned int& intraWarpThreadId)
{
    if (intraWarpThreadId >= 1)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 1];

    if (intraWarpThreadId >= 2)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 2];

    if (intraWarpThreadId >= 4)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 4];

    if (intraWarpThreadId >= 8)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 8];

    if (intraWarpThreadId >= 16)
        s_data[intraBlockThreadId] += s_data[intraBlockThreadId - 16];
}

但它显然更慢，所以我更愿意了解在我的第一个内核中发生了什么导致发布构建中的错误结果。如有任何建议，我将不胜感激。

【问题讨论】：

尝试将__shared__ 内存声明为volatile。可能是编译优化（这个效果在编程指南中有描述）
谢谢，嗯，现在我觉得很傻。事实上，添加 volatile 解决了这个问题。
也许您可以添加一个简短的答案以供将来参考

标签： cuda

【解决方案1】：

感谢 talonmies 指出问题。解决方案是在共享内存声明中添加 volatile。基本上在发布模式下，编译器被允许将值保存在寄存器中，并且不支持共享内存的中间存储。如果线程访问由其他线程修改的共享内存位置，则应将内存声明为 volatile 以关闭此类优化。

非常好的和深入的解释可以在这里找到：When to use volatile with shared CUDA Memory

【讨论】：