C-03 处方里反复出现 barrier但没回答哪一层、多贵、何时上 grid。本章用空同步 micro-bench 对照__syncwarp/__syncthreads/this_grid().sync()并用同载荷phases钉死「grid ≠ 自动更快」。本机结论断层在 grid vs SM 内~17×不是 warp vs block。TL;DR工程结论口径RTX 5090 /sm_120CUDA eventmedian。完整表见docs/results/C-04_sync_layers.md。主断层 grid vs block同配置空同步grid/block ≈ 17.2×nwarps8nblocks170。warp 与 block 只差约1.1×1.23×——处方是别把跨 SM 会合当便宜的 block sync不是「warp ≪ block」神话。__syncwarp≠ 免费也不能冒充 block同 warp 经内存交接要会合可见性跨 warp SMEM 交接走__syncthreadscorrectness有 sync 已过。Block 默认__syncthreads/block.sync()整 block 阶段分界、SMEM 交接用硬件 block barrier。C-02大 tile CG sync 勿当 block 默认。Grid 有门槛、随装填变贵、未必更快必须 cooperative launch本机 coop_max≈1020。sweep_grid12 blocks/SM 几乎平坦6/SM 时 0.63 ms相对 1 block ~4×phases单核0.0197 msvs 两 kernel~0.02 msratio≈1——为状态复用上 grid不是为加速launch 拆解 → C-06。判停看grid/blocksweep_grid装填跳变block/warp只作次级形状。禁止把 ncu 附着墙钟当结论。1. 问题barrier 写进处方之后问题本章交付三层空同步各多贵warp/block/gridsweepgrid 随规模怎么涨sweep_grid缺 block sync 会怎样correctness有 sync 断言无 sync UB 不冒充稳定失败grid 是否该替代双 kernelphases总墙钟一行边界章节已覆盖本章不重复C-01mask /*_sync不重测 shfl vs SMEMC-02CG 分组 / tile 悬崖不重测抽象税this_grid留给本章C-03原子争用 / stagingbarrier 只当承接不重测 atomic 曲线C-05 / C-06fusion / launch·Graph不做寄存器融合账不拆 launchB-08 / C-09mbarrier / named barrier仅 §7 钩子上warp 内__syncwarp。中block 内__syncthreadsSMEM 交接。下grid 内this_grid().sync()且必须 cooperative launch。2. 物理模型可见性范围不是「同一种 bar」隐式 barrier kernel 边界多 kernel / stream 同步 │ ├─ warp __syncwarp(mask) ← 同 warp 会合 memory fence ├─ block__syncthreads() ← 同 CTASMEM 交接默认这一层 └─ grid this_grid().sync() ← 跨 SM要 coop launch 可常驻层典型 API可见性直觉本机空同步直觉warp__syncwarp同 warp 经内存读写~0.01 ms 量级与 block 接近block__syncthreads/block.sync()同 block含 SMEM略贵于 warp~1.1×默认 SMEM 交接gridgrid.sync()全 grid贵一个数量级~17× block随 blocks/SM 跳隐式两次 kernel全局经 kernel 边界phases与 grid 单核同量级见 C-06文献形状Zhang et al.IISWC’24grid 延迟更相关每 SM block 数。本机sweep_grid12/SM 平坦满 coop6/SM跳到 ~0.63 ms——形状对齐warp vs block 的「数量级差」在本机空同步上不成立以grid/block为主结论。3. API 清单本章用到的层代表注意Warp__syncwarp()/__syncwarp(mask)参与集一致勿用它「假装」block barrierBlock__syncthreads()SMEM handoff 默认与cg::sync(this_thread_block())同层Gridcg::this_grid().sync()禁止只靠要用cudaLaunchCooperativeKernelHostcudaDevAttrCooperativeLaunch occupancygrid 规模夹紧到可常驻上限否则 launch 失败// grid sync与示例同构__global__voidk(intiters,unsignedlonglong*out){autogridcooperative_groups::this_grid();for(inti0;iiters;i)grid.sync();// ...}// host:cudaLaunchCooperativeKernel((void*)k,gridDim,blockDim,args,0,stream);论坛常见坑启动后this_grid().is_valid()0再sync→ launch failure。示例启动时打印CooperativeLaunch与coop_max_grid。4. 决策表信号建议同 warp、经 SMEM/global 交接__syncwarp或已含会合的*_sync集体同 block、SMEM 生产者–消费者 / 阶段分界__syncthreads跨 block 多阶段且要保 SMEM/寄存器状态grid sync cooperative launch先查 coop / occupancy跨 block 只是「算完再开下一核」多 kernel隐式 barrier先跑phases看总墙钟子集 warp 流水 / async copy 交接named barrier /cuda::barrier→ B-08 / C-09本章不做想靠「少 sync」省时间先确认少的是正确层乱删 block sync 是正确性事故5. 实验怎么设计mode问题进主结论warp/block/grid单层空同步 median定点sweepnwarps上block/warp次级本机仅 ~1.1×防「warp≪block」误读sweep_gridnblocks上 grid 时延 / 装填跳变主曲线correctness有__syncthreads的 SMEM 交接定点phasesgrid 单核 vs 两 kernel定点不拆 launchmodes定点全表主数字grid/block写结果用./bin/03_compute_primitives_04_sync_layers--modesweep ./bin/03_compute_primitives_04_sync_layers--modesweep_grid ./bin/03_compute_primitives_04_sync_layers--modemodes防 DCE空同步循环累加clock64()写回。iters默认 256。6. 实测RTX 5090完整表docs/results/C-04_sync_layers.md。口径CUDA eventmedian空同步iters256。6.1 Sweepblock/warpvsnwarpsnwarpswarp_msblock_msblock/warp10.01560.01240.80噪声勿采信20.01220.01331.1040.01160.01331.1580.01160.01351.16160.01330.01611.21320.01500.01851.23次级曲线差距有限别据此说「能 syncwarp 就别 syncthreads」。6.2 Sweep grid空grid.syncvsnblocksnblocksgrid_ms相对 nblocks110.1591.00×1701/SM0.1701.06×3402/SM0.1841.16×10206/SMcoop_max0.6323.97×6.3 Modes / phases定点路径median_ms相对warp0.0089—block0.0098block/warp1.11×grid0.1686grid/block17.2×phases_grid0.0197verify OKphases_two_k~0.02verify OKphases ratio≈1.0grid/two_kernel不自动更快怎么读先看17×再决定要不要上 gridphases≈1说明同载荷下 coop 单核并不白捡加速。7. 扩展阅读不抢后续 ModuleAsync arrive-wait / mbarrier见 §10 与 B-08全 block/warp 仍优先__syncthreads/__syncwarp官方口径。Named barrier / warp specialization候选 C-07/C-09。Launch 墙钟与 CUDA Graph→C-06本章phases只给总时延一行。SyncMicrobenchmark 原始数据与测量法见 §10-C。8. 误区与 SOP误区纠正this_grid().sync()改cudaLaunchCooperativeKernel查is_valid/ 属性用__syncwarp同步整 block SMEM改__syncthreadsgrid 规模开到「理论最大 block」按 occupancy 夹紧否则 coop launch 失败看见 grid sync 就期待加速先跑phases慢或持平都可能正确把 ncu 附着 ms 写进结论只信裸跑 event medianSOP最短定可见性范围warp / block / grid / 多 kernel。Block SMEM 交接 →__syncthreads先跑correctness。要测代价 →--mode sweep--mode sweep_grid。犹豫 grid vs 双核 →--mode phases再决定是否上 coop。launch / Graph 优化 → C-06不在本章拆。9. 小结与下一章同步不是一种 bar范围错了是正确性层选贵了是性能。空同步曲线回答「多贵」phases回答「该不该上 grid」。下一章C-05 Kernel fusion 代价边界fuse 少了 launch/同步点却可能打爆寄存器与 occupancy——不再复读本章三层 sync 曲线。10. 参考文献A. 官方CUDA Programming Guide — Synchronization /__syncthreadsUsing CUDA Warp-Level Primitives__syncwarpCUDA Programming Guide — Cooperative Groupsthis_grid/ coop launchCUDA Runtime —cudaLaunchCooperativeKernelAsync Barriers边界全 block/warp 仍建议经典 syncB. 工程NVIDIA 论坛this_grid→ invalid / launch failure本仓库 C-02 实测大 tile CG sync ≠ block 默认 barrierC. 实证Zhang et al., IPDPS’20 / arXiv:2004.05371SyncMicrobenchmarkIISWC’24Characterizing CUDA and OpenMP syncsPDF本仓库C-04_sync_layers.md5090 主结论D. 前沿 / 扩展Programmatic Dependent Launchsm_90 PG→ C-06 / ENamed barrier / warp specialization 工程笔记 → C-07/C-09