原子操作与归约：安全汇总并行结果

普通的读—改—写可能互相覆盖，形成 Race Condition。Atomic 操作保证一次修改不可分割，使并发 invocation 对共享整数的更新不会丢失。

@group(0) @binding(0)
var<storage, read_write> bins: array<atomic<u32>>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id: vec3u) {
  let bin = values[id.x] % 8u;
  atomicAdd(&bins[bin], 1u);
}

1. atomic 类型 WGSL 原子值使用 atomic<u32> 或 atomic<i32>，不能把普通 u32 当原子操作。
2. 地址传入 atomicAdd 接收目标原子变量的引用与增量。
3. 正确但可能争用 大量 invocation 命中同一地址时会序列化，分层归约通常更快。

所有 bin 的计数总和等于输入元素数量，不再因并发覆盖而随机变小。