CUDA Compute Capability 5.2下如何实现__half与float的原子加法？

核心思路 #

利用__half在内存中以16位数据存储的特性，通过原子操作16位整数模拟半精度原子加法——本质是把半精度值的读写转换为对其底层16位整数的原子比较交换（CAS）操作，全程保证原子性。

代码实现 #

__device__ void atomicAddHalf(__half* addr, float val) {
#if __CUDA_ARCH__ >= 700
    // Volta及以上架构直接调用原生半精度原子加法
    atomicAdd(addr, __float2half(val));
#else
    // Kepler及以下架构，用16位整数CAS模拟原子加法
    unsigned short* addr_short = reinterpret_cast<unsigned short*>(addr);
    unsigned short old_val, new_val;
    do {
        // 原子读取当前半精度值的底层16位整数
        old_val = *addr_short;
        // 转换为float进行加法运算
        float current_float = __half2float(*reinterpret_cast<const __half*>(&old_val));
        float new_float = current_float + val;
        // 结果转回半精度的底层16位整数
        new_val = *reinterpret_cast<const unsigned short*>(&__float2half(new_float));
        // CAS尝试更新，失败则循环重试（保证原子性）
    } while (atomicCAS(addr_short, old_val, new_val) != old_val);
#endif
}

关键注意点 #

• 依赖CUDA中__half的内存布局：始终以16位无符号整数存储IEEE半精度格式，跨架构兼容性稳定
• CAS循环在高并发场景下会有轻微性能开销，但这是无原生半精度原子支持时，保证原子性的唯一可行方案
• 加法过程中的精度损失和直接使用半精度运算的特性一致，属于正常现象

内容的提问来源于stack exchange，提问作者Skip