如何用两个u32缓冲区正确模拟u64的atomicAdd操作？

我需要对u64类型执行原子加法操作，但WebGPU原生不支持该操作，于是我将u64拆分为两个u32缓冲区来存储数值。但目前的实现存在线程安全问题，不确定如何正确模拟u64的atomicAdd效果，避免线程在加载和存储值期间修改内存。

当前的非线程安全实现 #

const BASE: u32 = 1u << 32; // 对应u64的低32位和高32位拆分基准

fn to_u64_parts(value: u32) -> vec2u {
    return vec2u(value / BASE, value % BASE);
}

fn add_u64_parts(a: vec2u, b: vec2u) -> vec2u {
    let sum_low = a.y + b.y;
    let carry = sum_low / BASE;
    let sum_high = a.x + b.x + carry;
    return vec2u(sum_high, sum_low % BASE);
}

fn main() {
    // .....

    // 将u32值转换为u64的双u32表示
    let b: vec2u = to_u64_parts(value);
    // 从缓冲区读取旧值
    var a = vec2u(0); 
    a.x = atomicLoad(&buffer[index]);
    a.y = atomicLoad(&buffer[index+1]);
    // 计算新值
    let result = add_u64_parts(a, b);
    // 写回缓冲区
    atomicStore(&buffer[index], result.x);
    atomicStore(&buffer[index+1], result.y);
}

这个实现仅在无其他线程修改同一索引的缓冲区时有效，否则会出现线程安全问题：线程1修改buffer[index+1]时，线程2可能读取到buffer[index]的旧值和buffer[index+1]的新值，导致计算出错误的结果。

CUDA的参考实现思路 #

CUDA中可以基于atomicCAS（比较并交换）实现任意原子操作，官方给出过double类型的atomicAdd示例，核心是通过循环重试，保证整个操作的原子性：

#if __CUDA_ARCH__ < 600
__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;

    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));

    // 注：使用整数比较避免NaN导致的死循环（因为NaN不等于自身）
    } while (assumed != old);

    return __longlong_as_double(old);
}
#endif

提问 #

请问这个基于atomicCAS的思路能否适配到WebGPU中？WebGPU里如何实现类似的u64 atomicAdd模拟？

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