Non possiamo usare le operazioni atomiche per variabili in virgola mobile in CUDA?

Question

Ho usato atomicMax() per trovare il valore massimo nel kernel 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); 
} 

Si sta gettando il seguente errore:

error: no instance of overloaded function "atomicMax" matches the argument list

I tipi di argomenti sono: (float *, float).

Answer 1

La risposta breve è no. Come potete vedere dallo atomic function documentation, solo gli argomenti interi sono supportati per atomicMax e gli argomenti interi a 64 bit sono supportati solo su dispositivi di capacità 3.5.

Answer 2

Questa è la sintassi per Atomic MAX

int atomicMax(int* address,int val); 

Ma ci sono eccezioni come atomicAdd che supportano galleggianti.

Answer 3

atomicMax non è disponibile per i tipi di galleggiante. Ma è possibile implementare tramite 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 4

è necessario mappare galleggiante orderedIntFloat utilizzare 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); 
}