Chúng ta không thể sử dụng các phép toán nguyên tử cho các biến số dấu chấm động trong CUDA?

Question

Tôi đã sử dụng atomicMax() để tìm giá trị lớn nhất trong các hạt nhân 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); 
} 

Nó được ném các lỗi sau:

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

Các loại đối số là: (float *, float).

Answer 1

Câu trả lời ngắn gọn là không. Như bạn thấy từ atomic function documentation, chỉ đối số nguyên được hỗ trợ cho các đối số số nguyên atomicMax và 64 bit chỉ được hỗ trợ trên các thiết bị có khả năng tính toán 3.5.

Answer 2

Đây là cú pháp cho nguyên tử MAX

int atomicMax(int* address,int val); 

Nhưng cũng có những ngoại lệ như atomicAdd có hỗ trợ nổi.

Answer 3

atomicMax không khả dụng cho các loại phao. Nhưng bạn có thể thực hiện nó qua 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

Bạn cần để ánh xạ phao để orderedIntFloat sử dụng 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); 
}