CUDA - atomicAdd 只加起来 16777216答案

【问题标题】：CUDA - atomicAdd only adds up to 16777216CUDA - atomicAdd 只加起来 16777216
【发布时间】：2013-09-30 00:34:53
【问题描述】：

在运行以下内核时，我遇到了以下容易重现的问题，除了浮点数的 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，顺便说一下 0x1000000 十六进制。问题如果OUT_ITERATIONS 就不会出现问题，也就是说如果我改成16777000为例，果然输出是16777000.0！

系统：NVidia-Titan、CUDA 5.5、Windows7

【问题讨论】：

标签： cuda

【解决方案1】：

此问题是由于 float 类型的精度有限造成的。

float 只有 24 位二进制精度。如果您将两个数字相加，其中一个比另一个大 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版本atomicAdd()的实现请参考cuda doc

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

【讨论】：

愚蠢的我。非常感谢！
没有对双精度原子添加的硬件支持。
@Archaea cuda doc 提供了一个使用 atomicCAS() 的实现，如上面的链接所示。
@Archaea 我认为您的链接和我的链接中的代码都提供了相同的功能，可以原子地为现有的双精度数据添加值。他们都使用 atomicCAS() 来实现它。而且我的链接中提供的实现似乎更有效。
当然只是用double替换不能解决问题。这就是为什么我提供了一个关于如何实现atomicAdd() 的double 版本的链接。

【解决方案2】：

20M 不适合float 中的可用整数精度。

float 数量没有 32 位尾数（通过观察“顺便说一下 0x1000000 十六进制”，您发现有多少尾数位），因此它不能以与 @987654323 相同的方式表示所有整数@或unsigned int可以。

16777216 是可以可靠存储在float 中的最大整数。

将您的存储范围限制在适合float 的范围内，或者使用其他一些表示形式，例如unsigned int 或double，如果您想可靠地将20M 存储为整数。

这并不是真正的 CUDA 问题。尝试在主机代码中的 float 中存储大整数时也会遇到类似的困难。

【讨论】：