2

私はCudaが初めてです。カーネルに配列の float 要素を追加しようとしていますが、最終結果が間違っています。アトミックに行う必要があるためですが、一方で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

float と double にもatomicAddを使用できます。次のように:

__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

この関数を使用して、float の 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) の場合、ネイティブ フロートatomicAdd(float* address, float val); があります。

コメントで指摘されているように、より効率的な実装が必要な場合は、並列削減を使用できます。NVIDIA 自体が効率的な実装を提供しています。これはここで見つけることができ、実装の詳細な説明はここで見つけることができます。

于 2012-11-10T21:53:20.140 に答える