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

【问题标题】：CUDA - atomicAdd(float) does not add very small valuesCUDA - atomicAdd(float) 不会添加非常小的值
【发布时间】：2023-03-08 23:56:01
【问题描述】：

当我使用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。

【问题讨论】：

1e-40 是单精度浮点格式的非正规数，超出 float 数据类型的精度范围。很可能编译器将其刷新为零。此外，如果内核在多个线程上启动，对同一内存位置的+= 操作将导致未定义的行为。
@sgar91：只有 1 个线程在运行。在第二种情况下，编译器不会将其刷新为 0，因此在第一种情况下没有理由将其刷新。此外，这是一个最小的代码示例，它的动机是一段更复杂的代码，其中 addit 是由于复杂的动态数值积分而分配的，因此编译器无法刷新它。
@sgar91：为了防止可能的编译器优化，我将 addit 的赋值更改为 addit = sinf(value) 而不是 addit = value。
不仅有 1 个线程在运行。一个完整的经线正在运行，即 32 个线程。
@Michael: docs.nvidia.com/cuda/cuda-c-programming-guide/… "在warp当前执行路径上的线程被称为活动线程，而不在当前路径上的线程是非活动的（禁用). 线程可能是非活动的，因为它们比它们的 warp 的其他线程更早退出，或者因为它们与当前由 warp 执行的分支路径位于不同的分支路径上，或者因为它们是最后一个线程一个块，其线程数不是扭曲大小的倍数。"

标签： cuda

【解决方案1】：

atomicAdd 是一条特殊指令，它不一定遵循在其他浮点运算（例如普通 fp 加法）上指定例如 -ftz=true 或 -ftz=false 时可能获得的相同刷新和舍入行为

如PTX ISA manual 中所述：

浮点运算 .add 是单精度 32 位运算。 atom.add.f32 舍入到最接近的偶数并将次正规输入和结果刷新为符号保持零。

因此，即使您指定-ftz=false（我相信，对于nvcc，这是默认值），普通浮点加法不应将非规范化值刷新为零，浮点原子加法操作到全局内存将刷新为零（总是）。

【讨论】：