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

Question

当我float atomicAdd(float *address, float val)用来添加小于约的浮点值时.1e-39to 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手册中所述:

浮点运算.add是单精度,32位运算.atom.add.f32舍入到最近的偶数并刷新次正常输入,结果为符号保留零.

因此,即使普通的浮点数添加不应该将非正规数刷新为零(如果你指定-ftz=false这是默认值nvcc),对全局内存的浮点原子添加操作将刷新为零(总是).