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示例。取决于您的硬件，它可能对共享内存变量可用也可能不可用，因为64位共享内存原子操作仅在计算能力2.x和3.x设备上受支持。在任何情况下，应该谨慎使用原子内存操作，因为串行化内存访问会大大降低性能。