CUDA，有原子读吗？

Question

CUDA，有原子读吗？

3

我正在开发一个CUDA程序，所有的块和线程都需要动态确定迭代问题的最小步长。我希望块中的第一个线程负责将全局dz值读入共享内存，以便其余线程可以对其进行归约。同时，其他块中的其他线程可能会写入它。 CUDA是否有原子读取选项或等效选项？我想我可以使用原子加法与零或其他方法。或者这是必要的吗？

template<typename IndexOfRefractionFunct>
    __global__ void _step_size_kernel(IndexOfRefractionFunct n, double* dz, double z, double cell_size)
    {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        if(idx >= cells * cells)
            return;

        int idy = idx / cells;
        idx %= cells;

        double x = cell_size * idx;
        double y = cell_size * idy;

        __shared__ double current_dz;
        if(threadIdx.x == 0)
            current_dz = atomicRead(dz);

        ...

        atomicMin(dz, calculated_min);
    }

我刚刚意识到，cuda似乎不支持双精度原子操作。是否有任何解决方法？

- chasep255

不幸的是，CUDA 没有相当于原子加载的等效物。您可以使用 atomicAdd(0) 来解决这个问题，或在常规加载之前添加屏障。虽然语义确实不是原子加载的那些，但据我所知，结果是相同的。 - user703016

@VeronikaPrüssels，我们之前在使用你的另一个化名时进行过类似的讨论。也许你应该阅读一下我的回答。CUDA对于1、2、4、8或16字节的正确对齐类型的读取是“原子性”的，不需要任何额外的努力。 - Robert Crovella

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

在CUDA中是否有atomicRead选项或等效选项。

原子操作的概念是允许将多个操作组合在一起，而不会受到其他线程的干扰。经典用途是读-改-写。所有三个RMW操作步骤都可以在内存中的给定位置原子地执行，而不会受到其他线程的干扰。

因此，在这种情况下，仅具有原子读（单独）的概念实际上没有意义。它只是一个操作。在CUDA中，所有基本类型（int、float、double等）的正确对齐读取都是原子性的，即所有读取都在一个操作中进行，而没有其他操作影响该读取或该读取的部分。

根据您展示的内容，似乎您的用例的正确性应该在不需要读操作特殊行为的情况下得到满足。如果您只想确保current_dz值从全局值获取，然后任何线程都没有机会修改它，则可以使用__syncthreads()在块级别简单解决此问题：

    __shared__ double current_dz;
    if(threadIdx.x == 0)
        current_dz = dz;
    __syncthreads(); // no threads can proceed beyond this point until
                     // thread 0 has read the value of dz

    ...

    atomicMin(dz, calculated_min);

如果你需要确保这种行为在整个网格中得到执行，那么我的建议是先将dz的初始值用线程不写入的方式进行设置，然后在另一个位置执行atomicMin操作（即在内核级别上将写/输出与读/输入分离）。

但是，再次强调，我并不认为这对你的用例是必需的。如果你只想获取当前的dz值，你可以使用普通的读取。你将得到一个“一致”的值。在网格级别上，可能已经发生了若干个atomicMin操作，但是它们中的任何一个都不会破坏读取，导致你读取一个虚假的值。你读取的值将是初始值或者是由atomicMin操作正确存储的某个值（根据你展示的代码）。

另外，我刚意识到cuda似乎不支持double类型的原子操作。有什么解决办法吗？

CUDA支持一组有限的64位原子操作。特别地，有一个64位的atomicCAS操作。编程指南演示了如何在自定义函数中使用此操作来实现任意64位原子操作（例如，在double量上进行64位atomicMin）。编程指南中的示例描述了如何执行double atomicAdd操作。以下是在double上执行atomicMin和atomicMax的示例：

__device__ double atomicMax(double* address, double val)
{
  unsigned long long int* address_as_ull =(unsigned long long int*)address;
  unsigned long long int old = *address_as_ull, assumed;

  while(val > __longlong_as_double(old) ) {
    assumed = old;
    old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val));
  }

  return __longlong_as_double(old);
}
__device__ double atomicMin(double* address, double val)
{
  unsigned long long int* address_as_ull =(unsigned long long int*)address;
  unsigned long long int old = *address_as_ull, assumed;

  while(val < __longlong_as_double(old) ) {
    assumed = old;
    old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val));
  }

  return __longlong_as_double(old);
}

作为良好的编程实践，原子操作应该谨慎使用，尽管 Kepler 全局 32 位原子操作速度相当快。但是，在使用这些自定义的 64 位原子操作时，特别需要注意；它们的速度明显比普通的读写慢。