」到「處方」)
B-09 把「布局決定事務(wù)」釘成可復(fù)現(xiàn)數(shù)字。Module B 讀完仍?？ㄔ趦杉挛业降资悄念?lèi)癥狀以及證據(jù)在哪個(gè) binary本章不新開(kāi) micro-bench先給「認(rèn)出癥狀」的分流再給「癥狀 → 證據(jù) → 處方」總表§4回鏈 B-01B-09。TL;DR工程結(jié)論數(shù)字為 RTX 5090 /sm_120各章已落盤(pán)形狀CUDA eventmedian完整表見(jiàn)docs/results/B-10_checklist.md。帶寬章絕對(duì) GB/s 可能含L2——主看加速比/相對(duì)形狀。先認(rèn)出癥狀層再選證據(jù)勿一上來(lái)?yè)Q API用 §2 的「時(shí)間花在哪 → SOL 分流 → 觀察速查」定層層錯(cuò)了處方必廢。每條處方掛證據(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本機(jī) touch1SoA/AoS ≈13.6×。Checklist 是導(dǎo)航判?？础改芊駨?fù)現(xiàn)該章形狀」細(xì)節(jié)回鏈各章正文同目錄B-0N*.md。1. 問(wèn)題有處方表仍對(duì)不上號(hào)各章都有決策表但排查時(shí)?？ㄔ诟懊嬉徊健€不知道自己是表里哪一行。交付說(shuō)明認(rèn)出癥狀§2無(wú) profiler 自問(wèn) → NSYS 時(shí)間線(xiàn) → NCU SOL → 觀察速查對(duì)上處方§4癥狀 → 證據(jù) → 處方 → 回鏈防混層§4三層 Async 布局 vs 引擎索引§5 docs/results/B-10_checklist.md邊界本章做本章不做B-01B-09匯總 回鏈本機(jī)形狀不重講機(jī)制、不重跑全家桶Module C一句鉤子不寫(xiě) warp / atomics / Graph左癥狀層。中證據(jù)binary --mode/docs/results。底線(xiàn)布局先于引擎。2. 如何認(rèn)出自己的癥狀目標(biāo)不是一次量對(duì)所有指標(biāo)而是用最少步驟鎖到 §4 的一行。順序固定端到端墻鐘 → 時(shí)間主要在 Host/UM 還是 Kernel │ ├─ Host / 拷貝 / UM 遷移占主導(dǎo) → §2.12.2先別開(kāi) NCU 摳 kernel └─ Kernel 占主導(dǎo) → §2.3 NCU SOL → §2.4 觀察速查 → §4 對(duì)應(yīng)行2.1 兩分鐘自問(wèn)還沒(méi)開(kāi) profiler按順序回答第一個(gè)「是」就進(jìn)該層不要并行猜 API。#自問(wèn)「是」→ 先看癥狀層進(jìn) §4 行1熱路徑大量cudaMallocManaged/ 跨進(jìn)程共享 managed且冷啟動(dòng)或尾延遲差UMUM 冷啟動(dòng) / 遷移風(fēng)暴2端到端里 H2D/D2H 很重或?qū)懥薓emcpyAsync多 stream 仍像串行Host?Device偽異步 / 偽 overlap3Kernel 慢且數(shù)據(jù)是寬 struct / 少字段掃描 / 行列轉(zhuǎn)置布局布局 / 跨步寫(xiě)4已用 SMEM tiling但改訪(fǎng)問(wèn)下標(biāo)或 pad 前后差很大SMEMbank 沖突5-Xptxas-v已見(jiàn) spill或剛加__launch_bounds__后暴跌Regspill / Occupancy 誤判6明確有可復(fù)用熱點(diǎn)卻在試 L2 persisting / 結(jié)果飄L2L2 駐留誤用7低算術(shù)強(qiáng)度、GMEM→SMEM 搬運(yùn)顯眼已在試 async/TMA設(shè)備內(nèi) asyncB-07 / B-08先分清層8以上都弱但有效帶寬差、訪(fǎng)問(wèn)跨步GMEM合并差不確定時(shí)先用秒表把「拷貝時(shí)間」和「kernel 時(shí)間」拆開(kāi)兩次 event 或一次 NSYS不要直接上 TMA。2.2 有 NSYS先看時(shí)間花在哪時(shí)間線(xiàn)主要看見(jiàn)癥狀層傾向下一步H2D/D2H 很長(zhǎng)kernel 很短或 memcpy 與 compute 首尾相接無(wú)重疊Host?Device§3 Host 行 → B-06Kernel 期間大量 page migrate / fault首輪遠(yuǎn)慢于穩(wěn)態(tài)UM§3 UM 行 → B-05Kernel 占?jí)︾姶箢^拷貝可忽略設(shè)備內(nèi)§2.3 NCU勿在 Host 層打轉(zhuǎn)Overlap / UM 是否成立以時(shí)間線(xiàn)為準(zhǔn)不以函數(shù)名帶不帶 Async 為準(zhǔn)。2.3 有 NCUSOL 分流后再對(duì)行只在kernel 已是主因時(shí)用?？?Speed of Light見(jiàn) §10-2SOL 粗分含義Module B 里優(yōu)先懷疑Memory ? Compute喂不飽或事務(wù)浪費(fèi)布局 / 合并 / bank / 有效帶寬B-09→B-01→B-02Compute ? Memory算力或短依賴(lài)墻少碰拷貝引擎spill/OccupancyB-03async 常不賺B-07/B-08雙低latency 風(fēng)險(xiǎn)發(fā)不出足夠 outstanding 工作低 AI 可試 pipelineB-07立刻 wait 的 async/TMA 常無(wú)效雙高已近硬件墻換算法/融合留給后續(xù) Module本表處方收益有限Memory 側(cè)再一眼sectors/request相對(duì)理想偏高float 合并?？?~4→先布局/合并再引擎。2.4 觀察 → §3 行速查把「我看見(jiàn)的現(xiàn)象」映射到主表仍迷糊就從上到下試一次只驗(yàn)證一行。我看見(jiàn)…先對(duì) §4 癥狀行最小驗(yàn)證pageable 緩沖 AsyncNSYS 無(wú) overlapHost 偽異步 / 偽 overlap06_pinned_dmapageablevspinned/overlapmanagedfirst ? median或 migrate 風(fēng)暴UM05_*faultvsprefetch少字段掃 particle/vertex或 NCU sec/req≈32布局09_* --mode sweep本機(jī) touch1 ~13.6×SMEM tile 改 stride/pad 吞吐跳變SMEM bank02_*pad/swizzleptxas spill↑ 且熱循環(huán)變慢Reg spill03_*-Xptxas-v開(kāi)了 persisting 但 DRAM 未降 / 忘記 resetL204_*三件套低 AIasync1不比 sync 快設(shè)備內(nèi) copy07_*sync/async1/pipe2/sweep大 tile立刻 wait 的 TMA≈1×TMA 指令稅08_* --mode sweepsm_90跨步明顯但非 AoS 故事GMEM 合并01_* sectors鎖到一行后才進(jìn)入 §4 的證據(jù)與處方。3. Triage 閉環(huán)定層之后1. 用 §2 鎖到癥狀行不要跳步換 API 2. 開(kāi)該行 binary看 event median 是否同向 3. 需要時(shí)間線(xiàn) → NSYS需要事務(wù)形狀 → NCU sectors旁證 4. 按處方只改一處 → 同一入口回歸 5. 形狀對(duì)上 → 停對(duì)不上 → 回 §2 換層勿疊 API主證據(jù) 裸跑 CUDA eventncu/nsys/SASS 旁證。4. 癥狀 → 證據(jù) → 處方主表正文回鏈同目錄B-0N*.md勿依賴(lài) URL 編碼。結(jié)果鏈如下。癥狀層典型信號(hào)證據(jù)入口處方一句話(huà)章Host?Device 偽異步 / 偽 overlappageable MemcpyAsync單流遠(yuǎn)低于 pinned06_pinned_dmapageable/pinned/serial/overlappinned 多非默認(rèn) stream成功≈貼單向 pinned~52 GB/sB-06UM 冷啟動(dòng) / 遷移風(fēng)暴first 與 steady 差大fault 時(shí)間線(xiàn)多 GPU ping-pong05_unified_memory_pffault/prefetch/advise先 prefetch同步advise 是 hint多 GPU 可寫(xiě)默認(rèn)棄 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 熱點(diǎn)未駐留 / 誤用 persistingstreaming 卻開(kāi) residency忘記 reset04_l2_residencytimeDRAML2 三件套有熱點(diǎn)才 residency掃 set-aside×hitRatio必 resetB-04設(shè)備內(nèi) copy 延遲藏不住低 AI 仍 syncasync 立刻 wait07_cp_async_pipelinesync/async1/pipe2/sweep低 AI → thread-localpipe2~1.15–1.31×高 AI 停B-07大 tile / 多維搬運(yùn)指令稅pipeline 已夠或不夠立刻 wait 的 TMA08_tma_introsm_90sweep先證明 overlap 段1 再上 TMA立刻 wait 常不賺~0.86–1.05×B-08布局 / 跨步寫(xiě)自殺少字段掃 AoSnaive transpose09_layout_transformsweep/transpose_*少字段→SoAtranspose→SMEM tiletiled≈copy91%naive≈39%B-095. 防混層兩張小表5.1 三層 Async層路徑章成功長(zhǎng)什么樣Host CEH2D/D2H DMAB-06時(shí)間線(xiàn)真 overlap吞吐貼 pinned 上限SMcp.async/ pipelineGMEM→SMEMB-07低 AI 加速比1async1立刻 wait 可更慢TMA bulk專(zhuān)用引擎 GMEM?SMEMB-08引擎∥computeprefetch立刻 wait 常 ≤15.2 布局 vs 引擎信號(hào)先做后做sectors/request 相對(duì)理想偏高float 合并?？?~4B-09 / B-01B-07 / B-08useful GB/s 差但合并已好、低 AIB-07 intensitysweep仍不夠再 B-08sm_90Host 側(cè)偽異步B-06與設(shè)備內(nèi) async無(wú)關(guān)—本機(jī)錨點(diǎn)索引非本章新測(cè)B-09 touch1 SoA/AoS~13.6×NCU AoS32vs SoA4sec/reqB-07 高 AI →pipe2≤1B-08pipe2低 AI 才明顯賺。6. 證據(jù)落點(diǎn)索引完整一頁(yè)docs/results/B-10_checklist.md。章Binarybuild/bin/或等價(jià)建議主命令結(jié)果 / 圖B-0102_memory_optim_01_global_mem_bandwidth章內(nèi) mode正文B-0202_memory_optim_02_shared_mem_bank_conflict章內(nèi) mode正文B-0302_memory_optim_03_register_spill章內(nèi) -Xptxas-v正文B-0402_memory_optim_04_l2_residency章內(nèi) 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)先級(jí)event median →可選NSYS →可選NCU →可選SASS。7. 工程邊界本章無(wú)新.cu、無(wú)新主結(jié)論 CSV數(shù)字全部回鏈已發(fā)布章。硬件總表不限架構(gòu)B-08 行繼承sm_90?？趶浇?ncu 附著程序自打印 ms/GB/s帶寬看加速比形狀。與 Module C訪(fǎng)存收束于此并發(fā)原語(yǔ)另開(kāi)模塊。8. 擴(kuò)展閱讀不搶后續(xù) ModuleGPU MODE Lecture 8 結(jié)構(gòu)參考CUDA Performance Checklist數(shù)字以本倉(cāng)庫(kù)為準(zhǔn)官方分流與 async 語(yǔ)義見(jiàn) §10-2、§10-3下一站Module CWarp / CG / Atomics / Fusion / GraphModule-C.md——不重開(kāi) coalescing / TMA 教程各章機(jī)制細(xì)目回對(duì)應(yīng)章 §10此處不重復(fù)抄表。9. SOP 與常見(jiàn)誤區(qū)建議流程先 §2 認(rèn)出癥狀行自問(wèn) → NSYS → SOL → 觀察速查不要跳過(guò)定層直接換 API打開(kāi) §4 對(duì)應(yīng)行的 binary只跑表內(nèi) mode對(duì)照docs/results或章內(nèi)表看形狀是否同向按處方改一處同一入口回歸Host 問(wèn)題勿用設(shè)備 async「補(bǔ)救」布局問(wèn)題勿先上 TMA形狀已對(duì)齊或加速比→1 →停判停證據(jù)入口復(fù)現(xiàn)不了該章形狀 → 查口徑payload / median / 是否 ncu 附著與問(wèn)題規(guī)模換了 API 但層沒(méi)對(duì) → 回 §2 / §5只有轉(zhuǎn)換成本、沒(méi)有熱路徑收益 → 別為 SoA / pipeline / TMA 而上§2 多行同時(shí)像 →先拆墻鐘Host vs Kernel再進(jìn) NCU高頻誤區(qū)還沒(méi)定層就換 API最常見(jiàn)把 HostcudaMemcpyAsync、cuda::memcpy_async、TMA 當(dāng)成同一層寫(xiě)了 async 卻立刻 wait以為已經(jīng)流水線(xiàn)化sectors/request 爆炸卻先換拷貝引擎用 ncu 附著墻鐘寫(xiě)進(jìn)結(jié)論把 Occupancy / L2 hit rate / 絕對(duì) GB/s 當(dāng)唯一 KPI多 GPU 可寫(xiě)場(chǎng)景死撐 UMChecklist 當(dāng)新理論不回鏈證據(jù)就改生產(chǎn)代碼10. 小結(jié)與下一章三句話(huà)先認(rèn)出癥狀§2再對(duì)處方§4API 名字不是診斷。三層 Async 分清布局先于引擎。Module B 到此收束細(xì)節(jié)在 B-01B-09證據(jù)在examples/docs/results/。下一模塊Module C從 Warp 原語(yǔ)與并發(fā)控制起筆不再重開(kāi)訪(fǎng)存教程。11. 參考文獻(xiàn)官方 / 工具CUDA C Best Practices GuidePinned / Coalescing / Shared / LocalNsight Compute Profiling GuideSpeed of Light / MemoryTriage 鏡像若 404 以此為準(zhǔn)CUDA Programming Guide — Asynchronous Data Copies · PipelinesNsight Systems User Guide工程 / 實(shí)證GPU MODE Lecture 8 筆記CUDA Performance Checklist結(jié)構(gòu)參考數(shù)字以本倉(cāng)庫(kù)為準(zhǔn)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擴(kuò)展 TACO’25DOI:10.1145/3746231本倉(cāng)庫(kù)各章 §10機(jī)制級(jí)鏈接見(jiàn) B-01B-09本章不重復(fù)整表。 B-10. Module B Checklist：從「癥狀」到「證據(jù)」到「處方」
B-09 把「布局決定事務(wù)」釘成可復(fù)現(xiàn)數(shù)字。Module B 讀完仍?？ㄔ趦杉挛业降资悄念?lèi)癥狀以及證據(jù)在哪個(gè) binary本章不新開(kāi) micro-bench先給「認(rèn)出癥狀」的分流再給「癥狀 → 證據(jù) → 處方」總表§4回鏈 B-01B-09。TL;DR工程結(jié)論數(shù)字為 RTX 5090 /sm_120各章已落盤(pán)形狀CUDA eventmedian完整表見(jiàn)docs/results/B-10_checklist.md。帶寬章絕對(duì) GB/s 可能含L2——主看加速比/相對(duì)形狀。先認(rèn)出癥狀層再選證據(jù)勿一上來(lái)?yè)Q API用 §2 的「時(shí)間花在哪 → SOL 分流 → 觀察速查」定層層錯(cuò)了處方必廢。每條處方掛證據(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本機(jī) touch1SoA/AoS ≈13.6×。Checklist 是導(dǎo)航判?？础改芊駨?fù)現(xiàn)該章形狀」細(xì)節(jié)回鏈各章正文同目錄B-0N*.md。1. 問(wèn)題有處方表仍對(duì)不上號(hào)各章都有決策表但排查時(shí)?？ㄔ诟懊嬉徊健€不知道自己是表里哪一行。交付說(shuō)明認(rèn)出癥狀§2無(wú) profiler 自問(wèn) → NSYS 時(shí)間線(xiàn) → NCU SOL → 觀察速查對(duì)上處方§4癥狀 → 證據(jù) → 處方 → 回鏈防混層§4三層 Async 布局 vs 引擎索引§5 docs/results/B-10_checklist.md邊界本章做本章不做B-01B-09匯總 回鏈本機(jī)形狀不重講機(jī)制、不重跑全家桶Module C一句鉤子不寫(xiě) warp / atomics / Graph左癥狀層。中證據(jù)binary --mode/docs/results。底線(xiàn)布局先于引擎。2. 如何認(rèn)出自己的癥狀目標(biāo)不是一次量對(duì)所有指標(biāo)而是用最少步驟鎖到 §4 的一行。順序固定端到端墻鐘 → 時(shí)間主要在 Host/UM 還是 Kernel │ ├─ Host / 拷貝 / UM 遷移占主導(dǎo) → §2.12.2先別開(kāi) NCU 摳 kernel └─ Kernel 占主導(dǎo) → §2.3 NCU SOL → §2.4 觀察速查 → §4 對(duì)應(yīng)行2.1 兩分鐘自問(wèn)還沒(méi)開(kāi) profiler按順序回答第一個(gè)「是」就進(jìn)該層不要并行猜 API。#自問(wèn)「是」→ 先看癥狀層進(jìn) §4 行1熱路徑大量cudaMallocManaged/ 跨進(jìn)程共享 managed且冷啟動(dòng)或尾延遲差UMUM 冷啟動(dòng) / 遷移風(fēng)暴2端到端里 H2D/D2H 很重或?qū)懥薓emcpyAsync多 stream 仍像串行Host?Device偽異步 / 偽 overlap3Kernel 慢且數(shù)據(jù)是寬 struct / 少字段掃描 / 行列轉(zhuǎn)置布局布局 / 跨步寫(xiě)4已用 SMEM tiling但改訪(fǎng)問(wèn)下標(biāo)或 pad 前后差很大SMEMbank 沖突5-Xptxas-v已見(jiàn) spill或剛加__launch_bounds__后暴跌Regspill / Occupancy 誤判6明確有可復(fù)用熱點(diǎn)卻在試 L2 persisting / 結(jié)果飄L2L2 駐留誤用7低算術(shù)強(qiáng)度、GMEM→SMEM 搬運(yùn)顯眼已在試 async/TMA設(shè)備內(nèi) asyncB-07 / B-08先分清層8以上都弱但有效帶寬差、訪(fǎng)問(wèn)跨步GMEM合并差不確定時(shí)先用秒表把「拷貝時(shí)間」和「kernel 時(shí)間」拆開(kāi)兩次 event 或一次 NSYS不要直接上 TMA。2.2 有 NSYS先看時(shí)間花在哪時(shí)間線(xiàn)主要看見(jiàn)癥狀層傾向下一步H2D/D2H 很長(zhǎng)kernel 很短或 memcpy 與 compute 首尾相接無(wú)重疊Host?Device§3 Host 行 → B-06Kernel 期間大量 page migrate / fault首輪遠(yuǎn)慢于穩(wěn)態(tài)UM§3 UM 行 → B-05Kernel 占?jí)︾姶箢^拷貝可忽略設(shè)備內(nèi)§2.3 NCU勿在 Host 層打轉(zhuǎn)Overlap / UM 是否成立以時(shí)間線(xiàn)為準(zhǔn)不以函數(shù)名帶不帶 Async 為準(zhǔn)。2.3 有 NCUSOL 分流后再對(duì)行只在kernel 已是主因時(shí)用?？?Speed of Light見(jiàn) §10-2SOL 粗分含義Module B 里優(yōu)先懷疑Memory ? Compute喂不飽或事務(wù)浪費(fèi)布局 / 合并 / bank / 有效帶寬B-09→B-01→B-02Compute ? Memory算力或短依賴(lài)墻少碰拷貝引擎spill/OccupancyB-03async 常不賺B-07/B-08雙低latency 風(fēng)險(xiǎn)發(fā)不出足夠 outstanding 工作低 AI 可試 pipelineB-07立刻 wait 的 async/TMA 常無(wú)效雙高已近硬件墻換算法/融合留給后續(xù) Module本表處方收益有限Memory 側(cè)再一眼sectors/request相對(duì)理想偏高float 合并?？?~4→先布局/合并再引擎。2.4 觀察 → §3 行速查把「我看見(jiàn)的現(xiàn)象」映射到主表仍迷糊就從上到下試一次只驗(yàn)證一行。我看見(jiàn)…先對(duì) §4 癥狀行最小驗(yàn)證pageable 緩沖 AsyncNSYS 無(wú) overlapHost 偽異步 / 偽 overlap06_pinned_dmapageablevspinned/overlapmanagedfirst ? median或 migrate 風(fēng)暴UM05_*faultvsprefetch少字段掃 particle/vertex或 NCU sec/req≈32布局09_* --mode sweep本機(jī) touch1 ~13.6×SMEM tile 改 stride/pad 吞吐跳變SMEM bank02_*pad/swizzleptxas spill↑ 且熱循環(huán)變慢Reg spill03_*-Xptxas-v開(kāi)了 persisting 但 DRAM 未降 / 忘記 resetL204_*三件套低 AIasync1不比 sync 快設(shè)備內(nèi) copy07_*sync/async1/pipe2/sweep大 tile立刻 wait 的 TMA≈1×TMA 指令稅08_* --mode sweepsm_90跨步明顯但非 AoS 故事GMEM 合并01_* sectors鎖到一行后才進(jìn)入 §4 的證據(jù)與處方。3. Triage 閉環(huán)定層之后1. 用 §2 鎖到癥狀行不要跳步換 API 2. 開(kāi)該行 binary看 event median 是否同向 3. 需要時(shí)間線(xiàn) → NSYS需要事務(wù)形狀 → NCU sectors旁證 4. 按處方只改一處 → 同一入口回歸 5. 形狀對(duì)上 → 停對(duì)不上 → 回 §2 換層勿疊 API主證據(jù) 裸跑 CUDA eventncu/nsys/SASS 旁證。4. 癥狀 → 證據(jù) → 處方主表正文回鏈同目錄B-0N*.md勿依賴(lài) URL 編碼。結(jié)果鏈如下。癥狀層典型信號(hào)證據(jù)入口處方一句話(huà)章Host?Device 偽異步 / 偽 overlappageable MemcpyAsync單流遠(yuǎn)低于 pinned06_pinned_dmapageable/pinned/serial/overlappinned 多非默認(rèn) stream成功≈貼單向 pinned~52 GB/sB-06UM 冷啟動(dòng) / 遷移風(fēng)暴first 與 steady 差大fault 時(shí)間線(xiàn)多 GPU ping-pong05_unified_memory_pffault/prefetch/advise先 prefetch同步advise 是 hint多 GPU 可寫(xiě)默認(rèn)棄 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 熱點(diǎn)未駐留 / 誤用 persistingstreaming 卻開(kāi) residency忘記 reset04_l2_residencytimeDRAML2 三件套有熱點(diǎn)才 residency掃 set-aside×hitRatio必 resetB-04設(shè)備內(nèi) copy 延遲藏不住低 AI 仍 syncasync 立刻 wait07_cp_async_pipelinesync/async1/pipe2/sweep低 AI → thread-localpipe2~1.15–1.31×高 AI 停B-07大 tile / 多維搬運(yùn)指令稅pipeline 已夠或不夠立刻 wait 的 TMA08_tma_introsm_90sweep先證明 overlap 段1 再上 TMA立刻 wait 常不賺~0.86–1.05×B-08布局 / 跨步寫(xiě)自殺少字段掃 AoSnaive transpose09_layout_transformsweep/transpose_*少字段→SoAtranspose→SMEM tiletiled≈copy91%naive≈39%B-095. 防混層兩張小表5.1 三層 Async層路徑章成功長(zhǎng)什么樣Host CEH2D/D2H DMAB-06時(shí)間線(xiàn)真 overlap吞吐貼 pinned 上限SMcp.async/ pipelineGMEM→SMEMB-07低 AI 加速比1async1立刻 wait 可更慢TMA bulk專(zhuān)用引擎 GMEM?SMEMB-08引擎∥computeprefetch立刻 wait 常 ≤15.2 布局 vs 引擎信號(hào)先做后做sectors/request 相對(duì)理想偏高float 合并?？?~4B-09 / B-01B-07 / B-08useful GB/s 差但合并已好、低 AIB-07 intensitysweep仍不夠再 B-08sm_90Host 側(cè)偽異步B-06與設(shè)備內(nèi) async無(wú)關(guān)—本機(jī)錨點(diǎn)索引非本章新測(cè)B-09 touch1 SoA/AoS~13.6×NCU AoS32vs SoA4sec/reqB-07 高 AI →pipe2≤1B-08pipe2低 AI 才明顯賺。6. 證據(jù)落點(diǎn)索引完整一頁(yè)docs/results/B-10_checklist.md。章Binarybuild/bin/或等價(jià)建議主命令結(jié)果 / 圖B-0102_memory_optim_01_global_mem_bandwidth章內(nèi) mode正文B-0202_memory_optim_02_shared_mem_bank_conflict章內(nèi) mode正文B-0302_memory_optim_03_register_spill章內(nèi) -Xptxas-v正文B-0402_memory_optim_04_l2_residency章內(nèi) 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)先級(jí)event median →可選NSYS →可選NCU →可選SASS。7. 工程邊界本章無(wú)新.cu、無(wú)新主結(jié)論 CSV數(shù)字全部回鏈已發(fā)布章。硬件總表不限架構(gòu)B-08 行繼承sm_90?？趶浇?ncu 附著程序自打印 ms/GB/s帶寬看加速比形狀。與 Module C訪(fǎng)存收束于此并發(fā)原語(yǔ)另開(kāi)模塊。8. 擴(kuò)展閱讀不搶后續(xù) ModuleGPU MODE Lecture 8 結(jié)構(gòu)參考CUDA Performance Checklist數(shù)字以本倉(cāng)庫(kù)為準(zhǔn)官方分流與 async 語(yǔ)義見(jiàn) §10-2、§10-3下一站Module CWarp / CG / Atomics / Fusion / GraphModule-C.md——不重開(kāi) coalescing / TMA 教程各章機(jī)制細(xì)目回對(duì)應(yīng)章 §10此處不重復(fù)抄表。9. SOP 與常見(jiàn)誤區(qū)建議流程先 §2 認(rèn)出癥狀行自問(wèn) → NSYS → SOL → 觀察速查不要跳過(guò)定層直接換 API打開(kāi) §4 對(duì)應(yīng)行的 binary只跑表內(nèi) mode對(duì)照docs/results或章內(nèi)表看形狀是否同向按處方改一處同一入口回歸Host 問(wèn)題勿用設(shè)備 async「補(bǔ)救」布局問(wèn)題勿先上 TMA形狀已對(duì)齊或加速比→1 →停判停證據(jù)入口復(fù)現(xiàn)不了該章形狀 → 查口徑payload / median / 是否 ncu 附著與問(wèn)題規(guī)模換了 API 但層沒(méi)對(duì) → 回 §2 / §5只有轉(zhuǎn)換成本、沒(méi)有熱路徑收益 → 別為 SoA / pipeline / TMA 而上§2 多行同時(shí)像 →先拆墻鐘Host vs Kernel再進(jìn) NCU高頻誤區(qū)還沒(méi)定層就換 API最常見(jiàn)把 HostcudaMemcpyAsync、cuda::memcpy_async、TMA 當(dāng)成同一層寫(xiě)了 async 卻立刻 wait以為已經(jīng)流水線(xiàn)化sectors/request 爆炸卻先換拷貝引擎用 ncu 附著墻鐘寫(xiě)進(jìn)結(jié)論把 Occupancy / L2 hit rate / 絕對(duì) GB/s 當(dāng)唯一 KPI多 GPU 可寫(xiě)場(chǎng)景死撐 UMChecklist 當(dāng)新理論不回鏈證據(jù)就改生產(chǎn)代碼10. 小結(jié)與下一章三句話(huà)先認(rèn)出癥狀§2再對(duì)處方§4API 名字不是診斷。三層 Async 分清布局先于引擎。Module B 到此收束細(xì)節(jié)在 B-01B-09證據(jù)在examples/docs/results/。下一模塊Module C從 Warp 原語(yǔ)與并發(fā)控制起筆不再重開(kāi)訪(fǎng)存教程。11. 參考文獻(xiàn)官方 / 工具CUDA C Best Practices GuidePinned / Coalescing / Shared / LocalNsight Compute Profiling GuideSpeed of Light / MemoryTriage 鏡像若 404 以此為準(zhǔn)CUDA Programming Guide — Asynchronous Data Copies · PipelinesNsight Systems User Guide工程 / 實(shí)證GPU MODE Lecture 8 筆記CUDA Performance Checklist結(jié)構(gòu)參考數(shù)字以本倉(cāng)庫(kù)為準(zhǔn)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擴(kuò)展 TACO’25DOI:10.1145/3746231本倉(cāng)庫(kù)各章 §10機(jī)制級(jí)鏈接見(jiàn) B-01B-09本章不重復(fù)整表。