Leoda

4-同步原子与内存模型

3 min read

管什么：线程之间怎么等、怎么原子地改同一块内存，以及”别的线程 / 别的卡什么时候能看到我写的数据”。

线程同步

block 内同步

跨 block 与跨卡同步

原子操作

atom 与 red

  • atom：读-改-写，返回旧值，要等结果回来。
  • red：同样的读-改-写交给 L2 原子单元做，不返回旧值，发出去就不用等，比 atom 便宜。
  • atomicAdd(float*, float) 不用返回值时，编译器一般会自动生成 red.global.add.f32。
red{.sem}{.scope}{.global|.shared}.add.type [addr], val;

__device__ __forceinline__ void red_add_f16x2(__half2* p, __half2 v) {
  unsigned int u = *reinterpret_cast<unsigned int*>(&v);
  asm volatile("red.relaxed.gpu.global.add.noftz.f16x2 [%0], %1;"
               :: "l"(p), "r"(u) : "memory");
}

坑：

  • 浮点 red 累加顺序不固定，结果不能逐位复现。
  • 目标地址在别的卡上（P2P / ICN 映射）时，依赖互联支持远程原子；ICN 是否支持未确认。

各类型支持

按 NVIDIA PTX 规范；810E 按 sm80 算。

类型最低 smsm80
.u32 .s32 .u64 .f32早就有能
.f64sm60能
.noftz.f16x2sm60能
.noftz.f16sm70能
.bf16 / .bf16x2sm90不能
向量版 .v2/.v4.f32 等sm90不能

sm90 的规约扩展

  • cp.reduce.async.bulk：TMA 整块写回并累加。
  • multimem.red：在 NVSwitch 上做规约（NVLS）。
  • red.async：往 cluster 里别的 CTA 的 smem 累加，配 mbarrier。

内存模型

语义 sem

  • .relaxed：只保证这次操作本身是原子的，不和别的读写排先后。适合纯计数。
  • “先写数据再发信号”：信号用 .release，读端用 .acquire，对端才保证看到新数据。

作用域 scope

.cta / .gpu / .sys：这次操作要对多大范围的线程可见、可排序。跨卡通信一般要 .sys。

fence

fence.acq_rel.gpu/.sys：单独插一道硬件内存屏障。编译器层面的 "memory" clobber 管不了硬件乱序，见 volatile 与 memory clobber。

Graph View