库达- atomicAdd加起来只有16777216

Question

在运行以下内核时，我遇到了以下容易重复的问题，它只执行浮点数的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的十六进制。如果< 16777216，也就是说，如果我将其更改为16777000，足够确定输出为16777000.0，则不会出现问题！

系统: NVidia-Titan，CUDA 5.5，Windows7

Answer 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

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

Answer 2

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

一个float数量不包含32位尾数(您发现有多少尾数位与您观察的“附带的0x1000000十六进制”一起存在)，所以它不能以与int或unsigned int相同的方式表示所有整数。

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

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

这其实不是数据自动化系统的问题。在主机代码中将大整数存储在float中也会遇到类似的困难。