原子添加到浮点值在OpenCL为NVIDIA gpu ?

Question

OpenCL 3.0规范似乎没有为浮点值提供原子加法的内在/内置功能，只有整数值(这似乎是OpenCL 1中的情况)。X和2。X也是)。然而，CUDA提供浮点原子运算已经有一段时间了:

float  atomicAdd(float*  address, float  val); // since Fermi
double atomicAdd(double* address, double val); // since Pascal
__half atomicAdd(__half *address, __half val); // ?

当然，任何简单的原子操作都可以用比较与交换来模拟，这个在OpenCL中是可用的。但我的问题是:

1. NVIDIA在OpenCL中以某种方式暴露浮点原子吗?例如，通过供应商扩展?用语法吗?隐式?
2. 是否有比比较交换模拟更有效的机制，我可以考虑作为浮点原子的替代品?是NVIDIA gpu还是普通gpu ?

Answer 1

正如@ProjectPhysX在他们的回答中暗示的那样，当你用NVIDIA的驱动程序编译OpenCL时，它接受内联PTX汇编(当然这不是OpenCL的一部分，也不是公认的供应商扩展)。这让你基本上可以做CUDA提供的任何事情——在OpenCL中;这包括自动添加浮点值。

因此，这里有一些包装器函数，用于在全局和局部内存中自动添加单精度(32位)浮点值:

float atomic_add_float_global(__global float* p, float val)
{
float prev;
asm volatile(
"atom.global.add.f32 %0, [%1], %2;" 
: "=f"(prev) 
: "l"(p) , "f"(val) 
: "memory" 
);
return prev;
}
float atomic_add_float_local(__local float* p, float val)
{
float prev;
// Remember "local" in OpenCL means the same as "shared" in CUDA.
asm volatile(
"atom.shared.add.f32 %0, [%1], %2;"
: "=f"(prev) 
: "l"(p) , "f"(val) 
: "memory" 
);
return prev;
}

还可以通过检查OpenCL驱动程序是否是NVIDIA的(在这种情况下使用内联汇编)或非NVIDIA的(在这种情况下使用原子比较交换实现)来调整这一点。

Answer 2

本机浮点原子是OpenCL 3.0非常需要的扩展。到目前为止，它们仍然不可用。

1. 唯一可能的方法是使用内联PTX。
2. 。FP32和FP64的原子比较交换实现是目前最先进的，没有更好的方法。

2022年6月更新:浮点原子被添加到OpenCL 3.0标准中，但硬件供应商可能还需要一段时间才能采用。