C-01 钉死了__shfl_*_sync与 mask 纪律。手写 mask 能跑但分组一变复杂就难组合。本章把分组收成句柄thread_block_tile/coalesced_group并用本机曲线钉死「抽象税≈0」与「tile32 悬崖」。配套可复现Yapeng-Gao/AI-System-Performance-Lab文章 .cu 实测表。有用请 Star。 本章示例examples/03_compute_primitives/02_cooperative_groups.cu。TL;DR工程结论口径RTX 5090 /sm_120CUDA eventmedianreduce 用--reps 50放大原语代价。完整表见docs/results/C-02_cooperative_groups.md。CG 分组句柄不是更快的 shufflethread_block_tileNN≤32把 mask/rank 收进类型本机 tile32 / cg_reduce 相对手写 intrinsic ≈1.00× / 0.99×。该上静态 tile固定子 warp / 半 warp 集体、少手写 mask、要把 group 当参数传进__device__函数时用tiled_partitionN。tile 大小悬崖本机cg::reduce在 tile≤32 持平64→1.60×128→1.66×相对 tile32。大 tile 走软件多 warp sync——别当 block-wide 默认。coalesced_group发散后「此刻活跃线程」的安全集合适合聚合原子入口不是任意逻辑 mask 的替代与 C-01activemask纪律同族。判停看tile32/intrinsic≈1 且sweep在 32 抬升cluster 有则 verify OK。禁止把 ncu 附着墙钟当结论。1. 问题intrinsic 会了分组怎么写得可组合问题本章交付少写 mask还能同速thread_block_tilevs 手写 intrinsic固定子集集体怎么表达tiled_partitionN发散后谁该参与集体coalesced_threads大 tilecg::reduce能不能当 block reducesweep悬崖曲线Hopper 簇内共享怎么验cluster可选支线功能边界章节已覆盖本章不重复C-01*_sync/ shfl vs SMEM / redux不重测 smem 树手写 shfl 只作基线标签A-04Divergence / ITS不重讲机制课C-03下Atomics 争用coalesced 只给入口形态 正确性C-04下同步分层 / grid sync不做this_grid/ cooperative launch 全家桶Module DDeviceReduce不做多 block 全局规约B-08TMA / async不抢 DSMEMTMA 深挖左静态 tile——编译期固定大小N≤32 走 warp 原语对应tiled_partitionN。右动态活跃集——发散后此刻活跃 laneleader 一次聚合写回对应coalesced_threads()。2. 物理模型先有 group再有集体this_thread_block() ← 启动配置隐含 │ ├─ tiled_partitionN → thread_block_tileN 固定、可嵌套 ├─ coalesced_threads() → coalesced_group 动态活跃集 └─ this_cluster() → cluster_group sm_90可选路径集合怎么定典型集体本机角色手写__shfl_*_sync你填 maskshuffle / vote基线thread_block_tileN分区几何 编译期 Nshfl_*/reduce/sync主对照coalesced_group此刻收敛的活跃 lanesize/shfl/ leader atomic定点cluster_group启动时 cluster 维sync/map_shared_rank可选功能N≤32 的 tile 落在单 warp 内同步便宜N32 跨多个 warpCG 用更通用的软件路径论坛口径busy-wait on memory墙钟会抬——这就是主曲线要钉的悬崖。3. API 分层本章用到的层代表作用Implicitthis_thread_block()拿到 block 句柄再 partitionTiletiled_partitionN(block)静态子组N 为 2 的幂且 ≤ blockCoalescedcoalesced_threads()发散分支内的活跃集Collectivesg.shfl_down/cg::reduce/g.sync在 group 上做集体Cluster可选this_cluster()簇屏障 DSMEM 映射namespacecgcooperative_groups;cg::thread_block blockcg::this_thread_block();cg::thread_block_tile32tilecg::tiled_partition32(block);floatvcg::reduce(tile,local,cg::plusfloat());coalesced聚合入口与示例同构if(pred){cg::coalesced_group gcg::coalesced_threads();if(g.thread_rank()0)atomicAdd(out,(unsignedlonglong)g.size());}mask 细则仍归 C-01本章用类型把「谁参与」说清楚。grid 级this_grid()cudaLaunchCooperativeKernel留给 C-04。4. 决策表信号建议固定 N≤32 的 warp/子 warp 集体要少写 maskthread_block_tileN或cg::reduce要把 group 当参数传入 device 函数传thread_block_tileN/coalesced_group发散后只对活跃线程聚合原子coalesced_threads()→ C-03 深挖争用block-wide / 多 warp 规约要性能别用大 tilecg::reduce回 C-01 处方或CUB整 device 规约 / SoftmaxModule D跨 block 同步 / grid barrierC-04簇内 DSMEM 小验证sm_90本章cluster支线生产路径另测处方与示例同构每线程 grid-stride 局部累加 │ ├─ intrinsic__shfl_down_sync 树基线 ├─ tile32thread_block_tile32.shfl_down └─ cg_reducecg::reduce(tile, …) sweep 扫 N--reps只放大规约轮次一次 GMEM 载入把「分组原语差价」从访存墙里剥出来verify 仍走 reps1 精确比对。5. 实验怎么设计项路径代码examples/03_compute_primitives/02_cooperative_groups.cu结果docs/results/C-02_cooperative_groups.md·C-02_sweep.csv/C-02_modes.csv绘图python scripts/plot_c02_cooperative_groups.py一条主命令主结论./bin/03_compute_primitives_02_cooperative_groups--modesweep定点全表抽象税 / coalesced / cluster./bin/03_compute_primitives_02_cooperative_groups--modemodesmode问题进主结论intrinsic/tile32/cg_reduce抽象税 tile32定点sweeptile∈{8,16,32,64,128} 的cg::reduce形状主曲线coalesced聚合计数正确定点表clustersm_90 DSMEM 邻块读可选modes定点一次跑齐写结果用证据优先级CUDA eventmedian→ 相对 tile32 的 norm正确性失败非零退出。NCU 可选本章默认不加 profile shell。5.1 本机实测RTX 5090 / sm_120平台与完整表docs/results/C-02_cooperative_groups.md。重画python scripts/plot_c02_cooperative_groups.pySweep主结论tilemedian_msnorm(÷tile32)80.02750.993160.02580.931320.02771.000640.04441.6031280.04611.664Modes 定点tile32tagmedian_ms相对intrinsic / tile32 / cg_reduce0.0253 / 0.0254 / 0.02501.004× / 0.990×coalesced0.2354verify OKodd8388608cluster0.0069verify OKclusize2怎么读抽象税≈0三种写法在噪声内——CG 不额外收通行费tile≤32。tile16 的 norm 0.93 略快于 32属亚 warp 波动不改「≤32 持平」结论。悬崖在 32 之后64/128 相对 tile32 贵约六成与「大 tile 软件 sync」同向。coalesced / cluster功能旁证时延不与 reduce 加速比横比。cluster 核几乎无计算短时延别当性能故事。5.2 旁证本章未跑 NCU。若要补对比intrinsicvscg_reduce32的inst_executed应对齐以及128的 barrier/stall 是否抬升。主结论仍以裸跑 median 为准。6. 工程边界项说明硬件tile / coalesced全架构常用路径this_clustersm_90类型示例 float reduce int coalesced 计数生产类型按 CUB/CG 文档正确性partial 求和 / odd_count / DSMEM 邻块值与 host 期望比对大 tileCC≤7.5 的大 tile 可能需block_tile_memorysm_80 按官方说明与 C-01mask 纪律不改本章只换表达层与 D多 block DeviceReduce → Module D7. 扩展阅读不抢后续章想继续去向生产级 block reduceCUBBlockReduce见 §10-B别用大 tile CG 顶替warp 聚合原子争用曲线C-03见 §10-D ARCgrid sync / cooperative launchC-04见 §10-C SyncMicrobenchmarkCluster DSMEM 深挖官方 Clusters 节见 §10-A不抢 B-08 TMA8. SOP 误区SOP确认集合是固定 tile、发散活跃集还是必须跨 warp/block。N≤32优先thread_block_tile/cg::reduce要极致可控仍可用手写*_sync。N32 或整 block回 C-01 处方或 CUB先跑--mode sweep看悬崖。发散聚合coalesced_threads leader争用曲线留给 C-03。判停modes抽象税≈1 且sweep在 32 抬升再考虑库与下章。误区误区正解CG 比手写 shuffle 更快本机 ≈同速赚的是可读性与可组合性cg::reduce(tile128)当默认 block reduce悬崖本机 ~1.6×用 CUB / C-01 模式coalesced_threads 任意逻辑 mask只是此刻活跃集逻辑 mask 仍自己算C-01cluster 短时延 生产加速比功能验证核DSMEM 收益另测继续用 legacy 无_syncshuffleVolta 仍走 C-01 纪律9. 小结与下一章分组先写成句柄tile≤32 几乎零抽象税跨过 32 要付软件 sync 的账。sweep回答「大 tile 值不值」coalesced 只演示聚合入口争用留给下一章。下一章C-03 Atomics 与 contentionglobal vs shared、分层规约、warp-aggregated 深挖——不再复读本章 CG 分层表。10. 参考文献A. 官方CUDA Programming Guide — Cooperative GroupsDevice-Callable APIs — Cooperative GroupsCUDA C Programming Guide — Thread Block ClustersClusters 节B. 工程NVIDIA, Cooperative Groups: Flexible CUDA Thread ProgrammingNVIDIA 论坛CG vs CUB reducetile≤32 vs 32CCCL/CUB — Block/Warp collectivesC. 实证Zhang et al., arXiv:2004.05371SyncMicrobenchmarkgrid sync 成本旁证 → C-04本仓库 C-01 结果C-01_warp_primitives.md手写 shfl 基线D. 前沿 / 扩展Collange, Warp-synchronous programming with Cooperative GroupsDurvasula et al., ARC, ASPLOS’25, DOI:10.1145/3669940.3707238 →C-03本文配套代码与实测AI-System-Performance-Lab。觉得有用请 Star后续章更新更好找。