」到「處方」)
B-09 把「布局決定事務」釘成可復現(xiàn)數(shù)字。Module B 讀完仍常卡在兩件事我到底是哪類癥狀以及證據(jù)在哪個 binary本章不新開 micro-bench先給「認出癥狀」的分流再給「癥狀 → 證據(jù) → 處方」總表§4回鏈 B-01B-09。TL;DR工程結論數(shù)字為 RTX 5090 /sm_120各章已落盤形狀CUDA eventmedian完整表見docs/results/B-10_checklist.md。帶寬章絕對 GB/s 可能含L2——主看加速比/相對形狀。先認出癥狀層再選證據(jù)勿一上來換 API用 §2 的「時間花在哪 → SOL 分流 → 觀察速查」定層層錯了處方必廢。每條處方掛證據(jù)入口examples/02_memory_optim/0N_*.cu --mode …若有docs/results/B-0N_*主證據(jù) 裸跑 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 是導航判?？础改芊駨同F(xiàn)該章形狀」細節(jié)回鏈各章正文同目錄B-0N*.md。1. 問題有處方表仍對不上號各章都有決策表但排查時?？ㄔ诟懊嬉徊健€不知道自己是表里哪一行。交付說明認出癥狀§2無 profiler 自問 → NSYS 時間線 → NCU SOL → 觀察速查對上處方§4癥狀 → 證據(jù) → 處方 → 回鏈防混層§4三層 Async 布局 vs 引擎索引§5 docs/results/B-10_checklist.md邊界本章做本章不做B-01B-09匯總 回鏈本機形狀不重講機制、不重跑全家桶Module C一句鉤子不寫 warp / atomics / Graph左癥狀層。中證據(jù)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 慢且數(shù)據(jù)是寬 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首輪遠慢于穩(wěn)態(tài)UM§3 UM 行 → B-05Kernel 占墻鐘大頭拷貝可忽略設備內§2.3 NCU勿在 Host 層打轉Overlap / UM 是否成立以時間線為準不以函數(shù)名帶不帶 Async 為準。2.3 有 NCUSOL 分流后再對行只在kernel 已是主因時用。看 Speed of Light見 §10-2SOL 粗分含義Module B 里優(yōu)先懷疑Memory ? Compute喂不飽或事務浪費布局 / 合并 / bank / 有效帶寬B-09→B-01→B-02Compute ? Memory算力或短依賴墻少碰拷貝引擎spill/OccupancyB-03async 常不賺B-07/B-08雙低latency 風險發(fā)不出足夠 outstanding 工作低 AI 可試 pipelineB-07立刻 wait 的 async/TMA 常無效雙高已近硬件墻換算法/融合留給后續(xù) Module本表處方收益有限Memory 側再一眼sectors/request相對理想偏高float 合并常看 ~4→先布局/合并再引擎。2.4 觀察 → §3 行速查把「我看見的現(xiàn)象」映射到主表仍迷糊就從上到下試一次只驗證一行。我看見…先對 §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↑ 且熱循環(huán)變慢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 的證據(jù)與處方。3. Triage 閉環(huán)定層之后1. 用 §2 鎖到癥狀行不要跳步換 API 2. 開該行 binary看 event median 是否同向 3. 需要時間線 → NSYS需要事務形狀 → NCU sectors旁證 4. 按處方只改一處 → 同一入口回歸 5. 形狀對上 → 停對不上 → 回 §2 換層勿疊 API主證據(jù) 裸跑 CUDA eventncu/nsys/SASS 旁證。4. 癥狀 → 證據(jù) → 處方主表正文回鏈同目錄B-0N*.md勿依賴 URL 編碼。結果鏈如下。癥狀層典型信號證據(jù)入口處方一句話章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. 證據(jù)落點索引完整一頁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_*證據(jù)優(yōu)先級event median →可選NSYS →可選NCU →可選SASS。7. 工程邊界本章無新.cu、無新主結論 CSV數(shù)字全部回鏈已發(fā)布章。硬件總表不限架構B-08 行繼承sm_90?？趶浇?ncu 附著程序自打印 ms/GB/s帶寬看加速比形狀。與 Module C訪存收束于此并發(fā)原語另開模塊。8. 擴展閱讀不搶后續(xù) ModuleGPU MODE Lecture 8 結構參考CUDA Performance Checklist數(shù)字以本倉庫為準官方分流與 async 語義見 §10-2、§10-3下一站Module CWarp / CG / Atomics / Fusion / GraphModule-C.md——不重開 coalescing / TMA 教程各章機制細目回對應章 §10此處不重復抄表。9. SOP 與常見誤區(qū)建議流程先 §2 認出癥狀行自問 → NSYS → SOL → 觀察速查不要跳過定層直接換 API打開 §4 對應行的 binary只跑表內 mode對照docs/results或章內表看形狀是否同向按處方改一處同一入口回歸Host 問題勿用設備 async「補救」布局問題勿先上 TMA形狀已對齊或加速比→1 →停判停證據(jù)入口復現(xiàn)不了該章形狀 → 查口徑payload / median / 是否 ncu 附著與問題規(guī)模換了 API 但層沒對 → 回 §2 / §5只有轉換成本、沒有熱路徑收益 → 別為 SoA / pipeline / TMA 而上§2 多行同時像 →先拆墻鐘Host vs Kernel再進 NCU高頻誤區(qū)還沒定層就換 API最常見把 HostcudaMemcpyAsync、cuda::memcpy_async、TMA 當成同一層寫了 async 卻立刻 wait以為已經流水線化sectors/request 爆炸卻先換拷貝引擎用 ncu 附著墻鐘寫進結論把 Occupancy / L2 hit rate / 絕對 GB/s 當唯一 KPI多 GPU 可寫場景死撐 UMChecklist 當新理論不回鏈證據(jù)就改生產代碼10. 小結與下一章三句話先認出癥狀§2再對處方§4API 名字不是診斷。三層 Async 分清布局先于引擎。Module B 到此收束細節(jié)在 B-01B-09證據(jù)在examples/docs/results/。下一模塊Module C從 Warp 原語與并發(fā)控制起筆不再重開訪存教程。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結構參考數(shù)字以本倉庫為準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本章不重復整表。 B-10. Module B Checklist：從「癥狀」到「證據(jù)」到「處方」
B-09 把「布局決定事務」釘成可復現(xiàn)數(shù)字。Module B 讀完仍常卡在兩件事我到底是哪類癥狀以及證據(jù)在哪個 binary本章不新開 micro-bench先給「認出癥狀」的分流再給「癥狀 → 證據(jù) → 處方」總表§4回鏈 B-01B-09。TL;DR工程結論數(shù)字為 RTX 5090 /sm_120各章已落盤形狀CUDA eventmedian完整表見docs/results/B-10_checklist.md。帶寬章絕對 GB/s 可能含L2——主看加速比/相對形狀。先認出癥狀層再選證據(jù)勿一上來換 API用 §2 的「時間花在哪 → SOL 分流 → 觀察速查」定層層錯了處方必廢。每條處方掛證據(jù)入口examples/02_memory_optim/0N_*.cu --mode …若有docs/results/B-0N_*主證據(jù) 裸跑 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 是導航判?？础改芊駨同F(xiàn)該章形狀」細節(jié)回鏈各章正文同目錄B-0N*.md。1. 問題有處方表仍對不上號各章都有決策表但排查時?？ㄔ诟懊嬉徊健€不知道自己是表里哪一行。交付說明認出癥狀§2無 profiler 自問 → NSYS 時間線 → NCU SOL → 觀察速查對上處方§4癥狀 → 證據(jù) → 處方 → 回鏈防混層§4三層 Async 布局 vs 引擎索引§5 docs/results/B-10_checklist.md邊界本章做本章不做B-01B-09匯總 回鏈本機形狀不重講機制、不重跑全家桶Module C一句鉤子不寫 warp / atomics / Graph左癥狀層。中證據(jù)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 慢且數(shù)據(jù)是寬 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首輪遠慢于穩(wěn)態(tài)UM§3 UM 行 → B-05Kernel 占墻鐘大頭拷貝可忽略設備內§2.3 NCU勿在 Host 層打轉Overlap / UM 是否成立以時間線為準不以函數(shù)名帶不帶 Async 為準。2.3 有 NCUSOL 分流后再對行只在kernel 已是主因時用。看 Speed of Light見 §10-2SOL 粗分含義Module B 里優(yōu)先懷疑Memory ? Compute喂不飽或事務浪費布局 / 合并 / bank / 有效帶寬B-09→B-01→B-02Compute ? Memory算力或短依賴墻少碰拷貝引擎spill/OccupancyB-03async 常不賺B-07/B-08雙低latency 風險發(fā)不出足夠 outstanding 工作低 AI 可試 pipelineB-07立刻 wait 的 async/TMA 常無效雙高已近硬件墻換算法/融合留給后續(xù) Module本表處方收益有限Memory 側再一眼sectors/request相對理想偏高float 合并常看 ~4→先布局/合并再引擎。2.4 觀察 → §3 行速查把「我看見的現(xiàn)象」映射到主表仍迷糊就從上到下試一次只驗證一行。我看見…先對 §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↑ 且熱循環(huán)變慢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 的證據(jù)與處方。3. Triage 閉環(huán)定層之后1. 用 §2 鎖到癥狀行不要跳步換 API 2. 開該行 binary看 event median 是否同向 3. 需要時間線 → NSYS需要事務形狀 → NCU sectors旁證 4. 按處方只改一處 → 同一入口回歸 5. 形狀對上 → 停對不上 → 回 §2 換層勿疊 API主證據(jù) 裸跑 CUDA eventncu/nsys/SASS 旁證。4. 癥狀 → 證據(jù) → 處方主表正文回鏈同目錄B-0N*.md勿依賴 URL 編碼。結果鏈如下。癥狀層典型信號證據(jù)入口處方一句話章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. 證據(jù)落點索引完整一頁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_*證據(jù)優(yōu)先級event median →可選NSYS →可選NCU →可選SASS。7. 工程邊界本章無新.cu、無新主結論 CSV數(shù)字全部回鏈已發(fā)布章。硬件總表不限架構B-08 行繼承sm_90?？趶浇?ncu 附著程序自打印 ms/GB/s帶寬看加速比形狀。與 Module C訪存收束于此并發(fā)原語另開模塊。8. 擴展閱讀不搶后續(xù) ModuleGPU MODE Lecture 8 結構參考CUDA Performance Checklist數(shù)字以本倉庫為準官方分流與 async 語義見 §10-2、§10-3下一站Module CWarp / CG / Atomics / Fusion / GraphModule-C.md——不重開 coalescing / TMA 教程各章機制細目回對應章 §10此處不重復抄表。9. SOP 與常見誤區(qū)建議流程先 §2 認出癥狀行自問 → NSYS → SOL → 觀察速查不要跳過定層直接換 API打開 §4 對應行的 binary只跑表內 mode對照docs/results或章內表看形狀是否同向按處方改一處同一入口回歸Host 問題勿用設備 async「補救」布局問題勿先上 TMA形狀已對齊或加速比→1 →停判停證據(jù)入口復現(xiàn)不了該章形狀 → 查口徑payload / median / 是否 ncu 附著與問題規(guī)模換了 API 但層沒對 → 回 §2 / §5只有轉換成本、沒有熱路徑收益 → 別為 SoA / pipeline / TMA 而上§2 多行同時像 →先拆墻鐘Host vs Kernel再進 NCU高頻誤區(qū)還沒定層就換 API最常見把 HostcudaMemcpyAsync、cuda::memcpy_async、TMA 當成同一層寫了 async 卻立刻 wait以為已經流水線化sectors/request 爆炸卻先換拷貝引擎用 ncu 附著墻鐘寫進結論把 Occupancy / L2 hit rate / 絕對 GB/s 當唯一 KPI多 GPU 可寫場景死撐 UMChecklist 當新理論不回鏈證據(jù)就改生產代碼10. 小結與下一章三句話先認出癥狀§2再對處方§4API 名字不是診斷。三層 Async 分清布局先于引擎。Module B 到此收束細節(jié)在 B-01B-09證據(jù)在examples/docs/results/。下一模塊Module C從 Warp 原語與并發(fā)控制起筆不再重開訪存教程。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結構參考數(shù)字以本倉庫為準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本章不重復整表。