如何在 CUDA 中对浮点值使用 atomicMax？

Posted 2023-04-15

【中文标题】如何在 CUDA 中对浮点值使用 atomicMax？

【英文标题】：How do I use atomicMax on floating-point values in CUDA?

【发布时间】：2013-06-28 06:11:55

【问题描述】：

我已经使用atomicMax()在CUDA内核中找到最大值：

__global__ void global_max(float* values, float* gl_max)

    int i=threadIdx.x + blockDim.x * blockIdx.x;
    float val=values[i];

    atomicMax(gl_max, val);

它抛出以下错误：

错误：没有重载函数“atomicMax”的实例与参数列表匹配

参数类型为：(float *, float)。

【问题讨论】：

不支持，但您可以创建自己的

【参考方案1】：

atomicMax 不适用于浮点类型。但是你可以通过atomicCAS实现它：

__device__ static float atomicMax(float* address, float val)

    int* address_as_i = (int*) address;
    int old = *address_as_i, assumed;
    do 
        assumed = old;
        old = ::atomicCAS(address_as_i, assumed,
            __float_as_int(::fmaxf(val, __int_as_float(assumed))));
     while (assumed != old);
    return __int_as_float(old);

【讨论】：

要实现浮点 atomicMin 版本，只需将 fmaxf 替换为 fminf。

我不确定这是一个好的解决方案：atomicCAS 的参数没有以“原子”方式处理：因此在评估 ::fmax(val, ...) 时可能会出现竞争条件。我尝试使用此实现，但结果是错误的输出。我怀疑是因为这个“非原子”::fmax。 Xiaojing An's solution 似乎效果更好。

@Shai 为什么会有fmax 的竞争条件？两者都是局部变量，只有当地址上的值是那里存储的值和val 中的最大值时，while 循环才会退出。话虽如此，另一种解决方案可能是更好的选择，因为它使用单个原子指令并且很可能更快

【参考方案2】：

基于CUDA Toolkit Documentation v9.2.148，浮点数没有原子操作。 但是我们可以通过将 atomicMax 和 atomicMin 与有符号和无符号整数转换混合来实现它！

这是一个浮点原子最小值：

__device__ __forceinline__ float atomicMinFloat (float * addr, float value) 
        float old;
        old = (value >= 0) ? __int_as_float(atomicMin((int *)addr, __float_as_int(value))) :
             __uint_as_float(atomicMax((unsigned int *)addr, __float_as_uint(value)));

        return old;

这是一个浮点原子最大值：

__device__ __forceinline__ float atomicMaxFloat (float * addr, float value) 
    float old;
    old = (value >= 0) ? __int_as_float(atomicMax((int *)addr, __float_as_int(value))) :
         __uint_as_float(atomicMin((unsigned int *)addr, __float_as_uint(value)));

    return old;

【讨论】：

是否也应该对*addr 进行测试是否为阴性？如果value 和*addr 的符号不同会怎样？

只知道值的符号就足够了。例如，在 atomicMinFloat 中，当 value >= 0 时，我们使用 atomicMin 作为有符号整数：如果 *addr =0，则比较 *addr 和 value 之间的最小值。

我认为这不能正确处理 float negative zero 的情况。您可以通过在使用之前将零添加到 value 来轻松解决此问题。

在我看来，对于atomicMinFloat：如果比较-0 和0，它会将*addr 设置为-0；否则，正确。您认为这里的第一个案例是不正确的吗？如果没有，我真的很感激用一个反例来澄清。 :)

【参考方案3】：

您需要将 float 映射到 orderedIntFloat 才能使用 atomicMax！

__device__ __forceinline__ int floatToOrderedInt( float floatVal ) 
 int intVal = __float_as_int( floatVal );
 return (intVal >= 0 ) ? intVal : intVal ^ 0x7FFFFFFF;

__device__ __forceinline__ float orderedIntToFloat( int intVal ) 
 return __int_as_float( (intVal >= 0) ? intVal : intVal ^ 0x7FFFFFFF);

【讨论】：

你怎么能以这种方式映射*address - 映射不是原子的？

【参考方案4】：

简短的回答是你不能。从atomic function documentation 可以看出，atomicMax 仅支持整数参数，而 64 位整数参数仅在计算能力 3.5 设备上支持。

【参考方案5】：

这是 Atomic MAX 的语法

int atomicMax(int* address,int val);

但也有像 atomicAdd 这样支持浮点数的异常。