B-09 把「布局决定事务」钉成可复现数字。Module B 读完仍常卡在两件事我到底是哪类症状以及证据在哪个 binary本章不新开 micro-bench先给「认出症状」的分流再给「症状 → 证据 → 处方」总表§4回链 B-01B-09。TL;DR工程结论数字为 RTX 5090 /sm_120各章已落盘形状CUDA eventmedian完整表见docs/results/B-10_checklist.md。带宽章绝对 GB/s 可能含L2——主看加速比/相对形状。先认出症状层再选证据勿一上来换 API用 §2 的「时间花在哪 → SOL 分流 → 观察速查」定层层错了处方必废。每条处方挂证据入口examples/02_memory_optim/0N_*.cu --mode …若有docs/results/B-0N_*主证据 裸跑 event median禁止ncu 附着墙钟。三层 Async 分清B-06 Host CE ≠ B-07 SMcp.async/pipeline≠ B-08 TMA bulk。布局先于引擎sectors/request 偏高或 useful 形状异常 → 先 B-09/B-01本机 touch1SoA/AoS ≈13.6×。Checklist 是导航判停看「能否复现该章形状」细节回链各章正文同目录B-0N*.md。1. 问题有处方表仍对不上号各章都有决策表但排查时常卡在更前面一步——还不知道自己是表里哪一行。交付说明认出症状§2无 profiler 自问 → NSYS 时间线 → NCU SOL → 观察速查对上处方§4症状 → 证据 → 处方 → 回链防混层§4三层 Async 布局 vs 引擎索引§5 docs/results/B-10_checklist.md边界本章做本章不做B-01B-09汇总 回链本机形状不重讲机制、不重跑全家桶Module C一句钩子不写 warp / atomics / Graph左症状层。中证据binary --mode/docs/results。底线布局先于引擎。2. 如何认出自己的症状目标不是一次量对所有指标而是用最少步骤锁到 §4 的一行。顺序固定端到端墙钟 → 时间主要在 Host/UM 还是 Kernel │ ├─ Host / 拷贝 / UM 迁移占主导 → §2.12.2先别开 NCU 抠 kernel └─ Kernel 占主导 → §2.3 NCU SOL → §2.4 观察速查 → §4 对应行2.1 两分钟自问还没开 profiler按顺序回答第一个「是」就进该层不要并行猜 API。#自问「是」→ 先看症状层进 §4 行1热路径大量cudaMallocManaged/ 跨进程共享 managed且冷启动或尾延迟差UMUM 冷启动 / 迁移风暴2端到端里 H2D/D2H 很重或写了MemcpyAsync多 stream 仍像串行Host↔Device伪异步 / 伪 overlap3Kernel 慢且数据是宽 struct / 少字段扫描 / 行列转置布局布局 / 跨步写4已用 SMEM tiling但改访问下标或 pad 前后差很大SMEMbank 冲突5-Xptxas-v已见 spill或刚加__launch_bounds__后暴跌Regspill / Occupancy 误判6明确有可复用热点却在试 L2 persisting / 结果飘L2L2 驻留误用7低算术强度、GMEM→SMEM 搬运显眼已在试 async/TMA设备内 asyncB-07 / B-08先分清层8以上都弱但有效带宽差、访问跨步GMEM合并差不确定时先用秒表把「拷贝时间」和「kernel 时间」拆开两次 event 或一次 NSYS不要直接上 TMA。2.2 有 NSYS先看时间花在哪时间线主要看见症状层倾向下一步H2D/D2H 很长kernel 很短或 memcpy 与 compute 首尾相接无重叠Host↔Device§3 Host 行 → B-06Kernel 期间大量 page migrate / fault首轮远慢于稳态UM§3 UM 行 → B-05Kernel 占墙钟大头拷贝可忽略设备内§2.3 NCU勿在 Host 层打转Overlap / UM 是否成立以时间线为准不以函数名带不带 Async 为准。2.3 有 NCUSOL 分流后再对行只在kernel 已是主因时用。看 Speed of Light见 §10-2SOL 粗分含义Module B 里优先怀疑Memory ≫ Compute喂不饱或事务浪费布局 / 合并 / bank / 有效带宽B-09→B-01→B-02Compute ≫ Memory算力或短依赖墙少碰拷贝引擎spill/OccupancyB-03async 常不赚B-07/B-08双低latency 风险发不出足够 outstanding 工作低 AI 可试 pipelineB-07立刻 wait 的 async/TMA 常无效双高已近硬件墙换算法/融合留给后续 Module本表处方收益有限Memory 侧再一眼sectors/request相对理想偏高float 合并常看 ~4→先布局/合并再引擎。2.4 观察 → §3 行速查把「我看见的现象」映射到主表仍迷糊就从上到下试一次只验证一行。我看见…先对 §4 症状行最小验证pageable 缓冲 AsyncNSYS 无 overlapHost 伪异步 / 伪 overlap06_pinned_dmapageablevspinned/overlapmanagedfirst ≫ median或 migrate 风暴UM05_*faultvsprefetch少字段扫 particle/vertex或 NCU sec/req≈32布局09_* --mode sweep本机 touch1 ~13.6×SMEM tile 改 stride/pad 吞吐跳变SMEM bank02_*pad/swizzleptxas spill↑ 且热循环变慢Reg spill03_*-Xptxas-v开了 persisting 但 DRAM 未降 / 忘记 resetL204_*三件套低 AIasync1不比 sync 快设备内 copy07_*sync/async1/pipe2/sweep大 tile立刻 wait 的 TMA≈1×TMA 指令税08_* --mode sweepsm_90跨步明显但非 AoS 故事GMEM 合并01_* sectors锁到一行后才进入 §4 的证据与处方。3. Triage 闭环定层之后1. 用 §2 锁到症状行不要跳步换 API 2. 开该行 binary看 event median 是否同向 3. 需要时间线 → NSYS需要事务形状 → NCU sectors旁证 4. 按处方只改一处 → 同一入口回归 5. 形状对上 → 停对不上 → 回 §2 换层勿叠 API主证据 裸跑 CUDA eventncu/nsys/SASS 旁证。4. 症状 → 证据 → 处方主表正文回链同目录B-0N*.md勿依赖 URL 编码。结果链如下。症状层典型信号证据入口处方一句话章Host↔Device 伪异步 / 伪 overlappageable MemcpyAsync单流远低于 pinned06_pinned_dmapageable/pinned/serial/overlappinned 多非默认 stream成功≈贴单向 pinned~52 GB/sB-06UM 冷启动 / 迁移风暴first 与 steady 差大fault 时间线多 GPU ping-pong05_unified_memory_pffault/prefetch/advise先 prefetch同步advise 是 hint多 GPU 可写默认弃 UMB-05GMEM 合并差 / 有效带宽低跨步sectors/request 偏高未向量化01_global_mem_bandwidthNCU Memory / sectors合并 向量化大 tile 再谈 TMAB-01SMEM bank 冲突同 bank 多路tile/transpose 掉速02_shared_mem_bank_conflictpad/swizzlepadding / swizzletranspose 用TILE1B-02寄存器 spill / Occupancy 误判ptxas spill↑盲目追 Occupancy03_register_spill-Xptxas-v先看 spill 是否在 hot pathOccupancy 非 KPIB-03L2 热点未驻留 / 误用 persistingstreaming 却开 residency忘记 reset04_l2_residencytimeDRAML2 三件套有热点才 residency扫 set-aside×hitRatio必 resetB-04设备内 copy 延迟藏不住低 AI 仍 syncasync 立刻 wait07_cp_async_pipelinesync/async1/pipe2/sweep低 AI → thread-localpipe2~1.15–1.31×高 AI 停B-07大 tile / 多维搬运指令税pipeline 已够或不够立刻 wait 的 TMA08_tma_introsm_90sweep先证明 overlap 段1 再上 TMA立刻 wait 常不赚~0.86–1.05×B-08布局 / 跨步写自杀少字段扫 AoSnaive transpose09_layout_transformsweep/transpose_*少字段→SoAtranspose→SMEM tiletiled≈copy91%naive≈39%B-095. 防混层两张小表5.1 三层 Async层路径章成功长什么样Host CEH2D/D2H DMAB-06时间线真 overlap吞吐贴 pinned 上限SMcp.async/ pipelineGMEM→SMEMB-07低 AI 加速比1async1立刻 wait 可更慢TMA bulk专用引擎 GMEM↔SMEMB-08引擎∥computeprefetch立刻 wait 常 ≤15.2 布局 vs 引擎信号先做后做sectors/request 相对理想偏高float 合并常看 ~4B-09 / B-01B-07 / B-08useful GB/s 差但合并已好、低 AIB-07 intensitysweep仍不够再 B-08sm_90Host 侧伪异步B-06与设备内 async无关—本机锚点索引非本章新测B-09 touch1 SoA/AoS~13.6×NCU AoS32vs SoA4sec/reqB-07 高 AI →pipe2≤1B-08pipe2低 AI 才明显赚。6. 证据落点索引完整一页docs/results/B-10_checklist.md。章Binarybuild/bin/或等价建议主命令结果 / 图B-0102_memory_optim_01_global_mem_bandwidth章内 mode正文B-0202_memory_optim_02_shared_mem_bank_conflict章内 mode正文B-0302_memory_optim_03_register_spill章内 -Xptxas-v正文B-0402_memory_optim_04_l2_residency章内 mode正文B-0502_memory_optim_05_unified_memory_pffault/ profile 脚本docs/results/B-05_*B-0602_memory_optim_06_pinned_dmamodes/overlapdocs/results/B-06_*B-0702_memory_optim_07_cp_async_pipeline--mode sweepdocs/results/B-07_*B-0802_memory_optim_08_tma_intro--mode sweepsm_90docs/results/B-08_*B-0902_memory_optim_09_layout_transform--mode sweepdocs/results/B-09_*证据优先级event median →可选NSYS →可选NCU →可选SASS。7. 工程边界本章无新.cu、无新主结论 CSV数字全部回链已发布章。硬件总表不限架构B-08 行继承sm_90。口径禁止 ncu 附着程序自打印 ms/GB/s带宽看加速比形状。与 Module C访存收束于此并发原语另开模块。8. 扩展阅读不抢后续 ModuleGPU MODE Lecture 8 结构参考CUDA Performance Checklist数字以本仓库为准官方分流与 async 语义见 §10-2、§10-3下一站Module CWarp / CG / Atomics / Fusion / GraphModule-C.md——不重开 coalescing / TMA 教程各章机制细目回对应章 §10此处不重复抄表。9. SOP 与常见误区建议流程先 §2 认出症状行自问 → NSYS → SOL → 观察速查不要跳过定层直接换 API打开 §4 对应行的 binary只跑表内 mode对照docs/results或章内表看形状是否同向按处方改一处同一入口回归Host 问题勿用设备 async「补救」布局问题勿先上 TMA形状已对齐或加速比→1 →停判停证据入口复现不了该章形状 → 查口径payload / median / 是否 ncu 附着与问题规模换了 API 但层没对 → 回 §2 / §5只有转换成本、没有热路径收益 → 别为 SoA / pipeline / TMA 而上§2 多行同时像 →先拆墙钟Host vs Kernel再进 NCU高频误区还没定层就换 API最常见把 HostcudaMemcpyAsync、cuda::memcpy_async、TMA 当成同一层写了 async 却立刻 wait以为已经流水线化sectors/request 爆炸却先换拷贝引擎用 ncu 附着墙钟写进结论把 Occupancy / L2 hit rate / 绝对 GB/s 当唯一 KPI多 GPU 可写场景死撑 UMChecklist 当新理论不回链证据就改生产代码10. 小结与下一章三句话先认出症状§2再对处方§4API 名字不是诊断。三层 Async 分清布局先于引擎。Module B 到此收束细节在 B-01B-09证据在examples/docs/results/。下一模块Module C从 Warp 原语与并发控制起笔不再重开访存教程。11. 参考文献官方 / 工具CUDA C Best Practices GuidePinned / Coalescing / Shared / LocalNsight Compute Profiling GuideSpeed of Light / MemoryTriage 镜像若 404 以此为准CUDA Programming Guide — Asynchronous Data Copies · PipelinesNsight Systems User Guide工程 / 实证GPU MODE Lecture 8 笔记CUDA Performance Checklist结构参考数字以本仓库为准Svedin et al.,Benchmarking the Nvidia GPU Lineage…arXiv:2106.04979低 AI async 受益、高 AI 可 ≤1Li et al.,Performance Implications of Async Memcpy and UVM…IISWC’23DOI:10.1109/IISWC59245.2023.00024PDF扩展 TACO’25DOI:10.1145/3746231本仓库各章 §10机制级链接见 B-01B-09本章不重复整表。