B-08. Hopper TMA：从单线程 Bulk Copy 到吞吐墙

B-07 解决的是 Ampere 路径cp.async/cuda::pipeline用多 stage 藏 GMEM 延迟。当 tile 更大、地址更不规则、算地址 发一堆 copy 本身吃指令带宽时下一站是把搬运卸给专用引擎。本章进入 HopperTMA单线程 issue、descriptor / tensor map、mbarrier交接——以及什么时候固定开销反而让你变慢。TL;DR工程结论TMA 是专用异步拷贝引擎GMEM↔SMEM亦可 cluster DSM与 B-07 的 SM 内cp.async、B-06 的 Host CE不是一层。收益来自单线程 issue 引擎搬运 ∥ compute不是「换 TMA 指令本身更快」。本机RTX 5090立刻 wait 的bulk1d/tensor2d约0.861.05×2-stagepipe2低 AI 可达约1.69×。该上需要大 tile / 多维搬运、或要把 copy 从计算 warp 卸掉warp-spec上之前先用sweep证明pipe2段 1。别上 / 先别上只换引擎立刻 wait、对齐搞不定、已 compute-bound本机fma≥128时pipe2已回落到 ~1.03×、B-07 pipeline 已够用。同步模型G2S 用mbarrierexpect_tx/arrive_txS2G 常走bulk async-group——两套完成模型不要混用。1. 问题B-07 之后还卡在哪B-07 的四层梯子里第 ④ 层只留了钩子① sync ② pipeline_memcpy ③ cuda::pipeline ④ TMA bulk gmem→reg→smem Ampere 异步拷贝 多 stage 流水线 专用引擎 ▲ 本章Amperecp.async已经旁路寄存器但每个线程仍要算自己的地址并发指令。tile 变大、维度变多时地址/predication 吃指令带宽、挤寄存器warp-spec 时 copy warp「发令」过贵。TMA 的回答一个线程把 descriptor 坐标交给硬件引擎其余线程去算或等 barrier。左每线程各自算地址、发一串小拷贝。右elect 线程发一条控制TMA 引擎整块搬 tile。层级API / 硬件谁算地址典型证据B-06 Host↔DeviceCE Stream驱动/CENSYSB-07 Devicecp.asyncLDGSTS / pipeline每线程intensity 曲线B-08 Device TMAbulk / tensor map描述符 引擎intensity 曲线本章2. 物理模型引擎、描述符、mbarrier┌─────────────┐ elect thread ───►│ 提交命令 │ (tensor_map, coords) 或 (ptr, bytes) │ │ 一条 issue │ │ └──────┬──────┘ │ │ 异步 │ ▼ GMEM ════════════════► TMA engine ════════════════► SMEM tile (HBM) │ (可旁路 L1) │ complete_tx按字节记账 ▼ mbarrier ──wait──► 全体线程开始读 SMEM / 计算时间线上两种用法差很多上bulk1d/tensor2d立刻 wait → copy 与 compute 串行本机 ≈ 0.91.0×。下pipe2算当前 tile 时预取下一块 → 才能藏延迟本机低 AI ≈ 1.7×。工程含义文献Luo et al.完整路径可有约170 cycle固定开销。本机4 KiB tile 立刻 wait 几乎不赚重叠才赚。对齐不满足时memcpy_async_tx/cp.async.bulk是UB——对齐必须写死。3. API 分层与同步心里记四层由浅到深① 1D cp.async.bulk / cuda::device::memcpy_async_tx ← 指针长度无 tensor map ② cuTensorMapEncode* cp.async.bulk.tensor ← 多维 tile 坐标 ③ cuda::memcpy_asyncHopper 满足对齐时也可走 TMA← 有回退对照实验慎用 ④ CuTe / CUTLASS PipelineTmaAsync ← 生产 GEMM扩展阅读3.1 1D bulk指针 字节数GMEM / SMEM16B 对齐size16 的倍数。用ptx::elect_sync或invoke_one选举 issue 线程避免if (threadIdx.x0)被插 peeling loop。memcpy_async_tx总是走 TMAtx 字节用barrier_arrive_tx/expect_tx显式声明。3.2 2D tensor map主机编码 __grid_constant__HostcuTensorMapEncodeTiled可用cudaGetDriverEntryPointByVersion。Kernelconst __grid_constant__ CUtensorMap。Devicecp.async.bulk.tensor 坐标目标 SMEM 128B 对齐。OOB 区域 G2S 零填充——硬件行为不是 bug。3.3 G2S vs S2G 完成模型本章 MVP只做 G2S。S2G 常见于 GEMM epilogue完成模型不同方向完成跟踪GMEM→SMEM本章mbarrier tx 字节SMEM→GMEM扩展cp.async.bulk.commit_groupwait_group_readS2G 写回前fence_proxy_async→__syncthreads()否则引擎可能读到旧 SMEM。3.4 和 B-07 pipeline 的关系本章 mode对应 B-07时间线见图本机结果bulk1d/tensor2d≈async1立刻 wait串行≈ 0.91.0×pipe2≈pipe2prefetch ∥ compute低 AI 最高约 1.69×生产级还会拆 producer warp只发 TMA/ consumer warpgroupWGMMA——本章不实现见 §7。4. 决策表何时上 TMA、何时回退 B-07信号建议需要大 tile / 多维 / 卸掉 copy 指令压力试 TMA先看pipe2的sweep再考虑嵌算子bulk1d/tensor2d≤1但pipe21本机典型保留 prefetch不要只换 API 立刻 wait已 compute-bound本机约fma≥128停TMA 协议变纯开销B-07sweep已够、且无多维/指令带宽墙先留在 B-07对齐搞不定回退 sync / B-07勿硬上memcpy_async_txUBcluster 多播同一份 K/V概念上 TMA multicastB-01本章不做冲 GEMM/Attention SOLModule D / CUTLASS / FA3本 micro-bench 不复现库性能5. 实验怎么设计配套代码examples/02_memory_optim/08_tma_intro.cu主命令一条就够./bin/02_memory_optim_08_tma_intro--modesweepmode回答什么sync协作 sync load 整 tile公平基线bulk1d1D TMA 立刻 waittensor2d2D tensor-map 立刻 waitpipe22-stage TMA prefetchsweep扫fma-iters→ 加速比 vs intensity主证据Tile 固定1024 floats32×324 KiBn须整除 tile。sync是整 tile 协作 load不是 B-07 的 per-thread 单元素才能和 TMA bulk 公平对照。sweep输出 CSVfma_iters,sync_ms,bulk1d_ms,tensor2d_ms,pipe2_ms,speedup_*,...证据优先级① 裸跑sweep→docs/results/B-08_*.csv② 可选 NCU / SASS。不要在ncu附着时读程序自己打印的 ms。需要sm_90Blackwell 建议-DCMAKE_CUDA_ARCHITECTURES120。增删.cu后请重新cmake。5.1 预期形状本机已验证speedup (sync / mode) ▲ │ pipe2 ╭── 低 AI明显 1 │ ╱ │──────╱──────── 1.0 │ ╱ bulk1d ≈ 1tensor2d 常 1 │ ╱ 高 AI全体 → 1 └──────────────────────► fma-iters5.2 RTX 5090 实测裸跑GPURTX 5090sm_120sharedMemPerBlock48 KB载荷n4194304tiles/block64block2562D 视图2048×2048口径CUDA eventmedian明细docs/results/B-08_tma.mdIntensity sweepfma_iterssync_msbulk1dtensor2dpipe2sp_bulk1dsp_tensor2dsp_pipe210.03680.03500.04040.02181.052×0.913×1.688×20.03620.03730.04210.02380.970×0.859×1.521×40.04090.04290.04750.02880.952×0.860×1.419×80.05050.05090.05660.04020.992×0.893×1.258×160.06950.07100.07540.05880.979×0.922×1.183×320.10740.10950.11470.09640.981×0.937×1.114×640.18360.18520.19030.17320.991×0.964×1.060×1280.33630.33760.34320.32540.996×0.980×1.034×2560.64180.64330.64760.62970.998×0.991×1.019×数据源docs/results/B-08_sweep.csv重画python scripts/plot_b08_tma.py。固定 modefma_iters8取自 sweep 同行modemedian (ms)相对 sync一句话sync0.05051.00×协作 sync load 基线bulk1d0.05090.99×换引擎立刻 wait ≈ 打平tensor2d0.05660.89×2D 立刻 wait 更慢pipe20.04021.26×prefetch 才赚数据源docs/results/B-08_modes.csv。怎么读藏延迟靠 overlappipe2在fma1约1.69×fma8约1.26×。换引擎 ≠ 加速bulk1d/tensor2d立刻 wait 全程约0.861.05×。高 AI 回落fma256时pipe2≈1.02×——算力淹没协议开销与 B-07 同形。6. 工程边界6.1 对齐与 UB路径关键对齐1D bulkGMEM/SMEM 16Bsize ×162D tensorGMEM 16Bstride ×16SMEM 128Bmemcpy_async可不满足时回退不能当「一定是 TMA」的证据6.2 落地后仍要管 bank / swizzleTMA 只负责「把字节放到 shared」。描述符 swizzle32B/64B/128B是给 WGMMA 友好布局用的不是bank conflict 免死金牌见 B-02。6.3 Cluster multicast / DSMB-01 提过 TMA multicast。本章 MVP不做需要时进 Cluster / Module D。6.4 Blackwell 仍保留 cp.asyncB-07 在消费级新卡上仍然有用TMA 是另一条梯子不是「有 Blackwell 就必须重写一切」。7. 扩展阅读不抢 Module DColfax, Mastering the Hopper TMAColfax / CUTLASSwarp-specialized PipelineTmaAsyncShah et al.,FlashAttention-3NeurIPS’24 / arXiv:2407.08608ACTAGPGPU’25自动选 tile/queueYadav et al.,CypressPLDI’25 / arXiv:2504.07004PyTorch, Hopper TMA for FP8 GEMMsdescriptor 开销反例Luo et al., arXiv:2501.12084TMA 延迟/吞吐 microbench8. 工程 SOP 与常见误区建议流程确认 sm_90直接跑./bin/02_memory_optim_08_tma_intro --mode sweep看speedup_pipe2低 AI 1、高 AI →1 → 标题成立若只有bulk1d/tensor2d≤1、没有 overlap 计划 →不要为 TMA 而 TMA有收益再嵌真实算子需要布局 → B-09判停pipe2低 AI 明显 1本机约 1.31.7×→ 值得在 latency-bound 路径保留 prefetch全程只有立刻 wait ≤1 → 回退 B-07或加大 tile / 上真 overlap 后再测仅tensor2d异常慢 → 先查 encode / 坐标 / 128B 对齐高频误区把 HostcudaMemcpyAsync、B-07memcpy_async、TMA bulk 当成同一件事写了 TMA 却立刻 wait又期待「自动变快」本机已打脸用threadIdx.x0issue 却不 electG2S / S2G 完成模型混用不对齐仍调memcpy_async_txUB每次 launch 重建/重传巨大 descriptor以为上了 TMA 就不用管 bank / swizzle / 布局9. 小结与下一章三句话先证明你需要 overlap或指令带宽墙再上 TMA——本机pipe2极低 AI 约1.69×赚的是引擎搬运 ∥ compute不是指令名字立刻 wait 的bulk1d≈1.0×、tensor2d≈0.89×用sweep判停——高 AI 回落到 ~1 是正常物理结果。下一章B-09回到数据布局AoS/SoA/Transpose。TMA 再强也救不了从第一天就选错的布局。10. 参考文献官方文档CUDA Programming Guide — Asynchronous Data CopiesNVIDIA Hopper Architecture In-DepthCCCL — Tensor Memory Accelerator (TMA)NVIDIA Hopper Tuning Guide工程教程Colfax, Mastering the Hopper TMAMLC.ai, Pipelining GEMM with TMAPyTorch, Deep Dive on the Hopper TMA Unit for FP8 GEMMs实证 / 前沿Luo et al., IPDPS’24 / arXiv:2402.13499Luo et al., arXiv:2501.12084Shah et al., FlashAttention-3NeurIPS’24 / arXiv:2407.08608ACTAGPGPU’25 / DOI:10.1145/3725798.3725802Yadav et al., CypressPLDI’25 / arXiv:2504.07004