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

atomicMaxfloat 型には使用できません。ただし、次の方法で実装できます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

atomicMaxを使用するには、float を OrderedIntFloatにマップする必要があります。

__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 3

短い答えは、できないということです。アトミック関数のドキュメント からわかるように、整数引数のみがサポートされてatomicMaxおり、64 ビット整数引数は計算機能 3.5 デバイスでのみサポートされています。

Answer 4

これは Atomic MAX の構文です

int atomicMax(int* address,int val);

ただし、フロートをサポートするatomicAddのような例外があります。