C-02 给了coalesced_threads聚合入口正确性已过。还缺一条轴撞同一地址有多贵staging / 手写聚合各值多少。本章用过滤计数 micro-bench 扫hit_rate本机结论与 Kepler 博客不同——以 5090 曲线为准。TL;DR工程结论口径RTX 5090 /sm_120CUDA eventmedian。完整表见docs/results/C-03_atomics_contention.md。争用决定墙钟同址atomicAdd堆得越满naive 越慢处方优先少撞 global 原子单元不是先换「更快的 atomic 助记符」。本机主处方 SMEM stagingblock 内先atomicAdd(SMEM)再每 block 一次 globalsmem/naive随 hit_rate 抬升1.0 → 6.31×modes6.56×。手写 warp-agg ≈ naive~1.0×同 warp、同地址、相同 1 时现代工具链/硬件多半已自动聚合coalesced写法仍值得保留可读、可移植别预期数量级加速。别硬套 Kepler Pro Tip当年「聚合 global ≫ SMEM」本机是staging ≫已等价的naive/agg global。数字用上表不抄 blog 20×。判停看smem/naive随 hit_rate 形状agg/naive≈1可接受。禁止把 ncu 附着墙钟当结论。1. 问题聚合入口有了争用账怎么算问题本章交付同址原子有多贵hit_rate争用曲线手写 warp 聚合还值不值aggvsnaiveblock staging 是否更赚smem/agg_smem和老博客数字怎么对齐本机校准表边界章节已覆盖本章不重复C-01ballot / elect 入口不重讲 maskC-02coalesced_threads形态 正确性不重测 CG 抽象税 / tile 悬崖C-04下同步分层__syncthreads只作 staging 一步Module DDeviceReduce / 直方图库不做生产 histogramARC / 渲染应用侧原子墙只作 §7 动机左每命中线程直接 global atomic。中warp 内聚合成一次 global。右先撞 SMEM再每 block 刷一次 global——本机主收益在右路。2. 物理模型少次数比「更快的 atomic」更先高争用同址计数 │ ├─ naive每命中线程 → atomicAdd(global) ← 请求挤在 L2 atomic 路径 ├─ agg coalesced 后每组一次 → global ← 少次数现代栈上常已被自动做掉 └─ smem atomicAdd(SMEM) → 每 block 一次 global ← 争用留在 SM 内路径全局原子次数量级本机角色naive~命中线程数基线墙钟随 hit_rate 涨agg~活跃 coalesced 组数与 naive 几乎同速smem~block 数有命中时主加速比agg_smem~block 数≈ smem直觉原子贵在冲突串行。能把冲突关进 SMEM、把 global 次数压到「每 block 一次」往往比在 global 上再抠 warp 聚合更赚——至少在本机同址计数形态如此。3. API / 模式本章用到的层代表作用Global atomicatomicAdd(ull*, 1)设备可见计数Shared atomicatomicAdd到__shared__block 内 stagingWarp 聚合cg::coalesced_threads() leader承接 C-02每组一次写回谓词控制in[i] thresh用hit_rate扫争用强度示例与 NVIDIA filtering Pro Tip 同构的简化版数组值域[0,999]thresh round(hit_rate*1000)命中则对同一全局计数器 1。// smem staging与示例同构__shared__unsignedlonglongblock_ctr;if(threadIdx.x0)block_ctr0;__syncthreads();if(in[i]thresh)atomicAdd(block_ctr,1ULL);__syncthreads();if(threadIdx.x0block_ctr)atomicAdd(out,block_ctr);agg路径继续用 C-02 的coalesced_threads不在本章重开 CG 教程。4. 决策表信号建议同 block 大量线程更新同一计数器本机SMEM staging→ 每 block 一次 global要可读的 warp 聚合写法 / 防旧工具链coalesced/ ballotelect别默认期待 ≫ naive多 bin、低冲突直方图先估冲突冲突低时原子本身可能不是主墙 → Module D / 专用库渲染梯度等多地址、变增量原子墙见 §7 ARC本章微基准不覆盖跨 block 同步 / grid barrierC-04整 device 规约Module D处方与示例同构谓词过滤计数扫 hit_rate │ ├─ naive命中 → atomicAdd(global) ├─ smem 命中 → atomicAdd(SMEM) → block flush ├─ agg 命中 → coalesced → leader atomicAdd(global) └─ agg_smemcoalesced → SMEM → block flush5. 实验怎么设计项路径代码examples/03_compute_primitives/03_atomics_contention.cu结果docs/results/C-03_atomics_contention.md·C-03_sweep.csv/C-03_modes.csv绘图python scripts/plot_c03_atomics_contention.py一条主命令主结论./bin/03_compute_primitives_03_atomics_contention--modesweep定点全表默认 hit_rate1./bin/03_compute_primitives_03_atomics_contention--modemodesmode问题进主结论naive/smem/agg定点时延对照agg_smemstaging聚合定点sweephit_rate∈{0.05…1.0} 的加速比形状主曲线modes定点一次跑齐写结果用证据eventmedian计数与 host 期望比对。NCU 可选默认无 profile shell。5.1 本机实测RTX 5090 / sm_120平台与完整表docs/results/C-03_atomics_contention.md。python scripts/plot_c03_atomics_contention.pySweep主结论hit_ratenaive_mssmem_msagg_msagg/naivesmem/naive0.050.03790.02630.03581.0601.4430.1250.05210.02670.05141.0151.9530.250.07900.02780.07901.0002.8390.50.13210.03080.13300.9934.2821.00.23350.03700.23351.0006.312Modeshit_rate1.0tagmedian_ms相对 naivenaive / agg0.2319 / 0.23101.004×smem / agg_smem0.0354 / 0.03546.56× / 6.55×怎么读agg 全程贴着 naive不是 verify 失败是差价已被吃掉。smem 才随争用拉开hit 越高staging 越值naive 近似线性变慢smem 几乎平坦。agg_smem ≈ smem已有 block staging 时再套 coalesced 几乎无额外收益。5.2 旁证本章未跑 NCU。若要补高 hit 下对比 naive vs smem 的 atomic / L2 相关 metric主结论仍以裸跑 median 为准。6. 工程边界项说明硬件atomicAddint/ull全架构常用本章主路径 ull 计数形态同址过滤计数多 bin / 变增量另测正确性各 mode hits 与 host 期望一致编译器朴素同址 1 可能被自动聚合——对照要诚实与 C-01/C-02聚合 API 已会本章只加争用轴与 D生产直方图 / Device 规约 → Module D7. 扩展阅读不抢后续章想继续去向Kepler 时代 warp-agg 故事NVIDIA Pro Tip见 §10-B数字勿直接当本机渲染梯度自适应原子ARC ASPLOS’25见 §10-Dgrid / block 同步分层C-04库级 histogram / DeviceReduceModule D/ CUB8. SOP 误区SOP确认是同址高争用还是多地址稀疏更新。同址计数先写SMEM staging再测--mode sweep。需要聚合写法时用coalesced/ ballot用agg/naive检查是否还有差价。与老 blog 数字冲突时以本机 sweep 为准。判停smem/naive形状合理 verify OK。误区误区正解手写 warp-agg 一定数倍于 naive本机 ≈1.0×先看工具链是否已聚合抄 Kepler「聚合 global 最优」本机 staging 更优重测再决策SMEM atomic 永远更快形态/架构相关用 sweep 说话本章 生产直方图只到同址计数处方库归 Module D把 ncu 附着 ms 当加速比禁止9. 小结与下一章同址争用本机先把原子关进 SMEM再每 block 刷 global手写 warp-agg 当保险写法不当加速神话。sweep回答「争用加重时谁在涨」同步原语全家桶留给下一章。下一章C-04 同步分层__syncwarp/ block barrier / grid sync——不再复读本章原子争用曲线。10. 参考文献A. 官方CUDA Programming Guide — Atomic FunctionsAdvanced Kernel Programming — scoped atomicsB. 工程NVIDIA, CUDA Pro Tip: Warp-Aggregated AtomicsNVIDIA, Cooperative Groupsaggregated atomic 段NVIDIA 论坛讨论shared vs global atomicTuring——SMEM 不总更快C. 实证本仓库C-03_atomics_contention.md5090 主结论C-02_cooperative_groups.mdcoalesced 入口D. 前沿 / 扩展Durvasula et al., ARC, ASPLOS’25, DOI:10.1145/3669940.3707238Garcia de Gonzalo et al., CGO’19自动 warp 原语 / 原子规约