1

次のカーネルを実行すると、次の簡単に再現可能な問題が発生します。これは、フロートのatomicAdds以外は何もしません。

#define OUT_ITERATIONS 20000000
#define BLOCKS 12
#define THREADS 192

__global__ void testKernel(float* result) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    float bias = 1.0f;
    int n = 1;

    while (i < OUT_ITERATIONS) {
        atomicAdd(result, bias);
        i += BLOCKS * THREADS;
    }
}

カーネルは、結果を OUT_ITERATIONS 回、つまり 20M ずつインクリメントすることになっています。この標準コードでカーネルを呼び出します。

int main() {
cudaError_t cudaStatus;
float* result;
float* dev_result;

// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
    goto Error;
}

result = new float;
cudaStatus = cudaMalloc((void**)&dev_result, sizeof(float));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}
cudaStatus = cudaMemset(dev_result, 0, sizeof(float));
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemset failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}

// Launch a kernel on the GPU with one thread for each element.
testKernel<<<BLOCKS, THREADS>>>(dev_result);

// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}

// cudaDeviceSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
    goto Error;
}

cudaStatus = cudaMemcpy(result, dev_result, sizeof(float), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMemcpy failed: %s\n", cudaGetErrorString(cudaStatus));
    goto Error;
}

printf("Result: %f\n", *result);

ただし、最後に出力される結果は 16777216.0 で、これは偶然にも 16 進数で 0x1000000 です。OUT_ITERATIONS < 16777216の場合、つまり、たとえば 16777000 に変更すると、問題は発生しません。出力は 16777000.0 です!

システム: NVidia-Titan、CUDA 5.5、Windows7

4

2 に答える 2

7

この問題は、型の精度が限られていることが原因floatです。

float24ビットのバイナリ精度しかありません。一方が他方の倍以上大きい 2 つの数値を加算する2^24-1と、結果は大きい方の数値とまったく同じになります。

16777216.0(=2^24) のような大きな数値と 1.0 のような小さな数値を加算すると、精度が失われ、結果は 16777216.0 のままになります。標準の C プログラムでも同じ状況が発生します。

float a=16777216.0f;
float b=1.0f;
printf("%f\n",a+b);

floatこの問題を置き換えるdoubleかint、解決することができます。

doubleのバージョンの実装については、cuda doc を参照してください。atomicAdd()

http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions

于 2013-09-27T16:52:39.197 に答える
4

20M は、a で使用可能な整数精度に収まりませんfloat。

数量には 32 ビットの仮数がありません( float「ついでに 16 進数で 0x1000000」を観察して仮数ビットがいくつあるかを発見しましたint) unsigned int。

16777216 は、確実に格納できる最大の整数ですfloat。

ストレージ範囲を に収まるように制限するか、または20M を整数として確実に格納したい場合floatなど、他の表現を使用してください。unsigned intdouble

これは実際には CUDA の問題ではありません。floatホストコードに大きな整数を格納しようとすると、同様の問題が発生します。

于 2013-09-27T16:48:55.140 に答える