c++ - CUDA，有 atomicRead 吗？

Question

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

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 似乎不支持 double 原子。有什么办法解决这个问题吗？

Answer 1

Is there simply an atomicRead option in CUDA or something equivalent.

atomic 操作的想法是，它允许组合多个操作，而无需来自其他线程的干预操作。规范用途是用于读取-修改-写入。 RMW 操作的所有 3 个步骤都可以针对内存中的给定位置以原子方式执行，而无需其他线程的干预事件。

因此，原子读取 的概念(仅就其本身而言)在这种情况下并没有真正的意义。这只是一个操作。在 CUDA 中，所有正确对齐的基本类型(int、float、double 等)读取都是原子发生的，即全部在一个操作中，没有其他操作影响该读取或该读取的一部分的可能性。

根据您所展示的内容，您的用例的正确性似乎应该得到满足，而读取操作没有任何特殊行为。如果您只是想确保从全局值填充 current_dz 值，在任何线程有机会修改它之前，在 block 级别，这可以简单地用 __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 值，您可以通过普通读取来完成。您将获得“一致”的值(value)。在网格级别，一些 atomicMin 操作可能发生在该读取之前，一些可能发生在该读取之后，但它们都不会破坏读取，导致你阅读虚假的值(value)。您读取的值要么是那里的初始值，要么是通过 atomicMin 操作正确存放的某个值(基于您显示的代码)。

Also I just realized that cuda does not seem to support atomics on doubles. Any way around this?

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

__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 位原子时，这个建议尤其适用；它们将明显比普通的读写慢。