2

我是库达的新手。我正在尝试在内核中添加数组的浮点元素，但最终结果是错误的。因为我需要以原子方式进行，但另一方面 atomicAdd 仅用于整数......有什么想法吗？

__global__ void add_element(float *my_array, float *result_sum){

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    *result_sum += my_array[tid];
}

int main(int argc, char** argv){

    float   my_array[10];   
    float   result_sum = 0;
    float   *device_array, *device_sum;

    cudaMalloc((void**)&device_array, 10*sizeof(float) );
    cudaMalloc((void**)&device_sum, sizeof(float) );

    //  fill the array
    for (int i=0; i<10; i++){
        my_array[i] = (float)i/2;
    }

    cudaMemcpy(device_array, my_array, 10*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(device_sum, &result_sum, sizeof(float),cudaMemcpyHostToDevice);

    add_element<<<1,10>>>(device_array, device_sum);

    cudaMemcpy(&result_sum, device_sum, sizeof(float), cudaMemcpyDeviceToHost);

    for(int i=0; i<10; i++){
        printf(" %f \n", my_array[i]);
    }   
    printf("+\n----------\n %f\n", result_sum);

    cudaFree(device_array);
    cudaFree(device_sum);

    return 0;
}
4

2 回答 2

1

您也可以将 atomicAdd 用于 float 和 double 。如下：

__device__ float atomicAdd(float *address, float val) { return 0; }

__device__ __forceinline__ float atomicAdd(float *address, float val)
{
// Doing it all as longlongs cuts one __longlong_as_double from the inner loop
unsigned int *ptr = (unsigned int *)address;
unsigned int old, newint, ret = *ptr;
do {
    old = ret;
    newint = __float_as_int(__int_as_float(old)+val);
} while((ret = atomicCAS(ptr, old, newint)) != old);

return __int_as_float(ret);
}

或找到文件“derived_atomic_functions.h”并作为头文件添加到您的项目中。

于 2013-08-31T13:48:03.993 回答
0

尝试使用此函数为浮点数执行 AtomicAdd。

__device__ inline void atomicFloatAdd(float *address, float val)
{
    int tmp0 = *address; 
    int i_val = __float_as_int(val + __int_as_float(tmp0)); 
    int tmp1;
    // compare and swap v = (old == tmp0) ? i_val : old;
    // returns old

    while( (tmp1 = atomicCAS((int *)address, tmp0, i_val)) != tmp0)
    { 
     tmp0 = tmp1; 
      i_val = __float_as_int(val + __int_as_float(tmp1));
    }
}

对于 (capability > 2.0) 你有一个本机 float atomicAdd(float* address, float val);

正如评论中所指出的，如果您想要更有效的实现，您可以使用并行缩减。NVIDIA 本身提供了一个高效的实现，可以在这里找到，实现的详细描述可以在这里找到

于 2012-11-10T21:53:20.140 回答