1

我是 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>>>(..);
4

1 回答 1

9

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

sum+= 2.0*3.0;

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

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

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

于 2013-05-28T06:54:59.697 回答