CUDA 多线程写入共享变量

Question

我是CUDA初学者。我这里有一个由 2 个线程执行的内核。所有线程都应该将它们的结果存储到一个共享变量中。三个都完成后，

sum

的结果应该是12，但我得到6！

__global__ void kernel (..)
{
    int i=blockDim.x*blockIdx.x+threadIdx.x;

    __shared__ double sum;
        
        ...

    if(i==0)
        sum=0.0;
    __syncthreads();

    if(i<=1)
        sum+= 2.0*3.0;
    __syncthreads();

    //sum should be 12 here, but I get 6. Why?
}

被

叫

test<<<1,2>>>(..);

Answer 1

您的代码中存在内存竞赛。这个：

sum+= 2.0*3.0;

潜在允许多个线程同时累加求和。在您的示例中，两个线程都试图同时在同一地址加载和存储。这是 CUDA 中的未定义行为。

避免这个问题的通常方法是重新设计算法。只是不要让多个线程写入同一内存位置。有一种广泛描述的共享内存缩减技术，您可以使用它来从共享内存阵列中累加总和，而不会出现内存竞争。

或者，可以使用原子内存访问原语来序列化内存访问。您的示例是双精度浮点数，我相当确定没有内在的原子加法函数。编程指南包括一个 user space atomic add for double precision 的示例。根据您的硬件，它可能会或可能不会在共享内存变量上使用，因为 64 位共享内存原子操作仅在计算能力 2.x 和 3.x 设备上受支持。在任何情况下，原子内存操作都应该谨慎使用，因为序列化内存访问会大大降低性能。