cuda add float array

Question

I'm new to Cuda. I'm trying to add the float elements of an array in the kernel, but the final result is wrong. because I need to do it atomically, but in the other hand atomicAdd is only used for integers...any ideas?

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

score 1 · Answer 1 · answered Aug 31 '13 at 13:48

You can use atomicAdd for float and double as well. as follow:

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

or find the file "derived_atomic_functions.h" and add in your project as your header file.

dreamcrash · Answer 2 · 2018-11-02T11:16:18.230

Try to use this function to do the AtomicAdd for floats.

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

For (capability > 2.0) you have a native float atomicAdd(float* address, float val);

As pointed out in the comments, if you want a more efficient implementation you can use a parallel reduction. NVIDIA itself provides an efficient implementation of one, which can be found here and a detailed description of the implementation can be found here

cuda add float array

2 Answers2