multithreading - 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中未定义的行为。

避免此问题的常用方法是重新设计算法。只是没有多个线程写入同一内​​存位置。有一种非常广泛描述的共享内存缩减技术，可用于从共享内存阵列中累加总和而无内存争用。

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