cuda::atomic 到底是做什么的？

Question

我刚刚发现 libcu++ 库并尝试使用这些cuda::atomic变量。我编写了以下程序，但它给了我意想不到的结果：

#include <atomic>
#include <cuda/atomic>
#include <stdio.h>


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}


__global__ void atomic_test()
{
    cuda::atomic<int, cuda::thread_scope_block> x{0};
    x.fetch_add(1, cuda::memory_order_seq_cst);
    __syncthreads();
    int y = x.load(cuda::memory_order_acquire);
    printf("(%d %d) - Value of x is %d\n", blockIdx.x, threadIdx.x, y); 
}

int main()
{
    atomic_test<<<2, 32>>>();
    gpuErrchk( cudaDeviceSynchronize() );
    return 0;
}

__syncthreads()由于后面有x.fetch_add()，我希望块中的每个线程都从 读取相同的值x。但是当我运行这个程序时，除了线程 31 之外的每个线程都打印0并且线程 31 打印32。有人可以解释一下我做错了什么以及为什么我会看到这个输出吗？

我尝试浏览 libcu++ 网站。但除了trie.cu和并发_hash_table.cu之外，我无法获得任何完整的示例。不幸的是，这些例子对我来说似乎有点太复杂了。

我还找到了使用semaphores 的堆栈溢出解决方案。但它按照我的预期工作。

Answer 1

在您的代码中，您在本地内存中声明一个原子，这意味着每个原子对于一个线程来说都是唯一的：它不能同时被多个线程访问。因此，将局部变量声明为原子变量没有任何价值。

为什么我会看到这个输出？

作为原子的一般优化，编译器将每个线程上的原子增量 1 替换为主导线程上的原子增量 32（warp 中的线程数）。鉴于此优化对于本地内存无效，此特定实例很可能是编译器错误。

实现此目的的正确方法是将原子存储在共享内存中。请注意，共享内存没有初始化机制。您可以通过在块中选择一个线程来初始化共享内存变量来完成此操作。下面这个例子是有效的：

__global__ void atomic_test()
{
    __shared__ cuda::atomic<int, cuda::thread_scope_block> x;
    if (threadIdx.x == 0)
    {
        x = 0;
    }
    __syncthreads();

    x.fetch_add(1, cuda::memory_order_seq_cst);
    __syncthreads();
    int y = x.load(cuda::memory_order_acquire);
    printf("(%d %d) - Value of x is %d\n", blockIdx.x, threadIdx.x, y);
}

并产生预期结果：

(1 0) - Value of x is 32
(1 1) - Value of x is 32
...
(0 30) - Value of x is 32
(0 31) - Value of x is 32