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

Question

我曾经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).

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

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

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

Answer 4

The short answer is that you can't. As you can see from the atomic function documentation, only integer arguments are supported for atomicMax and 64 bit integer arguments are only supported on compute capability 3.5 devices.

Answer 5

这是 Atomic MAX 的语法

int atomicMax(int* address,int val);

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