CUDA - atomicAdd(float) 不会添加非常小的值

Question

当我使用 float atomicAdd(float *address, float val) 时添加一个小于约的浮点值。 1e-39至 0 ，加法不起作用，值在address保持为 0。

这是最简单的代码:

__device__ float test[6] = {0};
__global__ void testKernel() {
    float addit = sinf(1e-20);
    atomicAdd(&test[0], addit);
    test[1] += addit;
    addit = sinf(1e-37);
    atomicAdd(&test[2], addit);
    test[3] += addit;
    addit = sinf(1e-40);
    atomicAdd(&test[4], addit);
    test[5] += addit;
}

当我以 testKernel<<<1, 1>>>(); 运行上面的代码时并停止我看到的调试器:

test    0x42697800
    [0] 9.9999997e-21
    [1] 9.9999997e-21
    [2] 9.9999999e-38
    [3] 9.9999999e-38
    [4] 0            
    [5] 9.9999461e-41

注意 test[4] 和 test[5] 之间的区别。两者都做同样的事情，但是简单的加法起作用了，而原子的加法什么也没做。我在这里缺少什么？

更新:系统信息:CUDA 5.5.20、NVidia Titan 卡、驱动程序 331.82、Windows 7x64、Nsight 3.2.1.13309。

Answer 1

atomicAdd 是一种特殊指令，它不一定遵循您可能会得到的刷新和舍入行为，如果您指定例如 -ftz=true 或 -ftz=false 其他浮点运算(例如普通的 fp add)

如 PTX ISA manual 中所述:

The floating-point operation .add is a single-precision, 32-bit operation. atom.add.f32 rounds to nearest even and flushes subnormal inputs and results to sign-preserving zero.

因此，即使您指定 -ftz=false(我相信这是 nvcc 的默认设置)，普通浮点加法不应将非正规数刷新为零，浮点原子添加操作到全局内存将刷新为零(总是)。