CUDA，有原子阅读吗?

CUDA, is there an atomicRead?

本文关键字：CUDA      更新时间：2023-10-16

我正在开发一个 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 似乎不支持双打原子。 有什么办法吗？

CUDA 中是否只有一个原子读取选项或等效的东西。

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

因此，原子读取的概念（仅其本身）在这种情况下并没有真正的意义。 这只是一个操作。 在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 似乎不支持双打原子。有什么办法吗？

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);
}

作为一种良好的编程实践，原子应该谨慎使用，尽管开普勒全局 32 位原子非常快。 但是，当使用这些类型的自定义 64 位原子时，该建议特别适用;它们将明显比普通读写慢。

相关文章：