目錄
CUDA優化
在GPU平行運算中，CUDA優化是提升效能的關鍵。本指南詳細解析 CUDA 快取機制與記憶體存取流程，協助開發者降低延遲、提升資料吞吐量。透過正確的記憶體架構配置與快取策略，您能有效優化CUDA程式，發揮 GPU 最大效能。本文由 TAKI Cloud 專業團隊透過專注於AI伺服器與高效能運算的實務經驗，帶您全面掌握 CUDA優化技巧。
專業實務經驗說明
TAKI Cloud技術團隊擁有多年實際協助客戶進行 CUDA 程式優化的經驗，服務對象涵蓋 AI 影像辨識、深度學習模型訓練、科學運算等高效能運算領域。我們協助國內外多家AI新創與研究機構將程式在 RTX A6000、RTX4090 等 GPU 伺服器上執行效能提升最高可達 5倍，具體作法包括：
重新設計記憶體存取模式，避免全域記憶體隨機訪問造成瓶頸。
善用共享記憶體與 L1 快取，減少讀寫延遲。
精細調整 Grid 與 Block 配置，確保每個執行緒群能充分利用 GPU 資源。
對 CUDA 核函數進行 Profile 分析，鎖定效能熱點進行逐步優化。
透過以上最佳化經驗，我們不僅幫助客戶縮短 AI 模型訓練時間，還協助他們在有限預算內達成雲端資源最高使用效益。
L1/L2 Cache
這篇文章針對 ncu 與 kernel 優化：快取 & 訪存流程，主要是來了解 GPU 中的各級快取以及訪存的主要流程，後面會簡單介紹 Memory Chart 和 Memory Chart 中的各個指標及其意義，主要參考的是 NVIDIA 官方文件：https://docute.nvidiaCompute/bvidia
這文章我們不會有相關程式碼調優分析，更多的是合併訪存的一個補充。
我們先看下 NVIDIA 官方文件對於 Caches 的描述：
當 SM（Streaming Multiprocessor）需要存取資料時，會先從快取中查詢：如果 L1 快取中已經有資料，就能立即取得，存取速度非常快；如果L1沒命中，系統會再到 L2 快取查找；若 L2 也找不到，最後才會去全域記憶體（DRAM）讀取，這步驟會花最多時間。啟用L1快取後，SM 每次讀取全域記憶體資料都會先經過 L1 和 L2 快取，能大幅提升存取速度，減少等待時間。
上圖給出了 NVIDIA GPU 中 L1TEX 級別的緩存其內部各單元如何協同完成請求分發、命中/未命中判斷、與 L2 二級快取交互以及最終把數據送回給 SM（Streaming Multiprocessor）的流水線模型。
下面我們分段來講下其中各項資料/請求的流向：(from ChatGPT)
SM 如何發起請求
Warp Scheduler → MIO
在 GPU 中，SM（Streaming Multiprocessor）裡的多個 warp（即同時執行的執行緒組）會不斷向記憶體提出讀取或寫入的需求，像是讀取全域記憶體（global）、區域記憶體（local）、共享記憶體（shared），或是使用紋理/表面（texture/surface）資料。這些需求會集中由一個稱作 MIO（Memory I/O crossbar）的模組來負責統一收發，MIO 就像是「記憶體請求的交通樞紐」，確保每個 warp 的記憶體需求能高效被處理並送到對應的記憶體單元。
MIO → L1TEX
- global 和 local 記憶體訪問，連同 shared（訪問，一起透過 MIO → LSUIN。
- texture 和 Surface 記憶體存取則透過 MIO → TEXIN。
L1TEX 內部關鍵階段
Tag Stage
- 只針對 global/local 的 load/store 以及所有 texture/surface 請求做 Tag 查詢。
- 命中（hit）：把對應 cache line（sector）送到 Data Stage。
- 未命中（miss）：把缺少的 sector 請求送到 Miss Stage 去向下游（L2）獲取。
Miss Stage
- 總結所有未命中扇區（sector），把它們打包後發給 L2。
- 當 L2 回傳資料時，再把磁區交給 Data Stage。
Data Stage
- 接收來自 Tag Stage 或 Miss Stage 的扇形資料。
- 根據請求類型，分發到 LSU Data 或 TEX Data。
LSU Data
把資料送回給 SM 的 Warp Scheduler，用於普通的 GPU load/store。
TEX Data + TEX Filter
- 先把扇區原始位元組送入 TEX Data，再透過 TEX Filter 做插值/過濾（linear，point，shadow sampling 等）。
- 最終返回給 SM，用作 texture 採樣的顏色或深度值。
Shared 記憶體請求
直接從 LSUIN 繞過 Tag/Miss 階段，走專門的資料通路到 LSU Data，因為 shared memory 不走 L1 cache 的常規標籤配置。
與 L2 Cache 的交互
- Miss Stage → L2：所有 L1 未命中的扇區統一發給 L2。
- L2 → L1（Miss Stage/Data Stage）：L2 將磁區資料傳回 Miss Stage，再流入 Data Stage 分發給 LSU 或 TEX。
總和整體流程：
- 雙重輸入：Load/Store 請求（LSUIN）和 Texture 請求（TEXIN）同時競爭 L1TEX 的標籤和資料通路。
- 三級流水：Tag Stage 判定 → Miss Stage（若未命中）→ Data Stage 匯聚。
- 雙輸出：LSU Data（給普通 Load/Store）和 TEX Data+Filter（給紋理取樣）兩條迴路。
- Shared 記憶體：特殊繞過路徑，效率更高。
- 多層快取協作：L1TEX 負責快速命中，Miss 時下探到 L2，再回填 L1，以提高整體頻寬與吞吐。
L2 Cache：
上圖呈現了 GPU 中二級快取（L2 Cache）的工作邏輯，以及資料和控制訊息如何在不同層級之間流動。所有上層的快取單元（例如 L1TEX、L1V 等）和其他記憶體用戶端，都會透過 L2 快取來和主記憶體（DRAM）溝通。當有資料要讀取或寫入時，L2 快取會負責判斷是否命中、未命中，或處理原子操作與記憶體一致性的請求。
值得特別一提的是，NVIDIA GPU 的 L2 快取並不是一個單一、巨大的快取區，而是被拆分成多個「切片」（slice），每個切片負責特定地址範圍的資料。這些切片分散在晶片上各處，能同時處理不同資料區塊的存取請求，實現平行化加速。
在 NVIDIA 官方文件中，整體的 L2 快取稱為 ltc（L2 cache），而單個切片則叫作 lts（L2 cache slice）。每個 lts 切片內部又會分成多個階段，包括標籤比對階段（Tag Stage，lts_t）、未命中處理階段（Miss Stage，lts_m）以及資料讀寫階段（Data Stage，lts_d）。
圖中綠色的大框標示的是「LTS」，指的是單個 L2 快取切片內部的資料流處理，而不是整個 L2 系統。雖然 L2 是所有SM（Streaming Multiprocessor）共享的記憶體資源，但硬體上透過多個 lts 切片來實作，可同時處理多個存取請求，提升記憶體存取效能，並幫助分析性能瓶頸時能細緻地觀察每個切片的行為。
L2快取的定位與功能說明
L2 Cache（L2快取）位於 GPU 晶片內，扮演各 SM（流式多處理器）和物理記憶體之間的資料中繼站。當 SM 需要讀取資料卻在 L1 快取找不到時，就會向L2請求；此外像 DMA、Texture 資料寫回或畫面緩衝區的讀寫，也都會經過 L2 統一管理。
L2 以物理地址作為索引，能夠管理各SM之間，以及 CPU/主機與 GPU 之間的資料共享與快取一致性，確保多個單元同時讀寫同一區域時不會出錯。
除了基本的快取功能，L2 還內建了：
壓縮/解壓縮硬體：可以把資料壓縮後再寫入記憶體，減少對外部記憶體頻寬的壓力，加速資料傳輸。
全域原子操作單元：支援像 atomicAdd、CAS 等 GPU 級原子運算，方便多個執行緒同時安全地修改共享變數。
請求的入口與標籤判斷
來自 L1 的未命中階段（Miss Stage）
當程式需要存取資料時，首先會在 L1 快取（包括資料快取 L1D、指令快取 L1C、紋理快取 L1TEX 等）中查找。如果在 L1 沒找到資料（這叫「未命中」），這些資料請求就會被彙整起來，一起交給下一層 —— L2 快取，進行後續處理。
標籤檢查階段（Tag Stage）
L2 快取會檢查這些請求資料在自己的記憶體裡是否已經存在，這個步驟就像對照資料的地址標籤：
命中（hit）：找到資料，直接回傳給程式使用，不需要再往更慢的主記憶體找。
未命中（miss）：沒找到資料，就會記錄這個「miss」，並通知下一階段（Data Stage）去主記憶體中把資料取回來。
數據階段與主存交互
資料階段（Data Stage）
在這個階段，L2快取接收來自「Tag階段」的兩種情況：
命中（hit）扇區：資料已存在於L2快取，直接從L2的資料記憶體（Data RAM）讀出回應給請求方。
未命中（miss）扇區：資料不在L2快取，需要將請求「打包」後送到對應的記憶體控制器（Memory Controller），去外部記憶體（DRAM）讀取資料。
當 DRAM 把資料取回後，資料會先送回到 Data Stage，再寫入（回填）到 L2 快取裡，同時把資料回應給發出請求的 L1 快取。
快取一致性（Coherence）
L2 快取透過晶片內部的高速互連（L2 fabric），負責處理跨多個 SM（Streaming Multiprocessor）之間的快取一致性問題。
舉例來說：
如果其他 SM 對相同扇區進行寫入，或有失效通知（invalidate），需要確保所有 SM 的快取狀態一致。
這些一致性相關的訊息也會作為「扇區級請求」進入 Data Stage，由硬體自動完成一致性處理，無需軟體介入。
全局原子操作（Atomic）
當CUDA程式中需要全局原子操作(像是 atomicAdd、atomicCAS)，這些指令會把資料從 L1 快取送到 L2 Cache 的「原子單元」。
壓縮 / 解壓縮（可選特性）
當大量資料同時訪問外部記憶體時，會造成記憶體頻寬成為效能瓶頸。
總結：L1/L2 Cache 訪存流程
- 所有 L1 Miss、DMA、幀緩衝區訪問、texture 寫回等操作，最終都會集中到 L2 Cache 作統一處理。
- L2 Cache 的”標籤單元”會快速判斷資料是否命中。
- 如果命中，直接從 L2 內部 RAM 回應；未命中則將資料打包送往主存，資料返回後會同時更新 L2 快取並回傳給 L1。
- L2 Cache 透過高速交換架構(fabric)，確保多個 SM 的快取資料能保持一致性。
- 當有全域原子操作時，會在 L2 的 Atomic 單元完成操作，保證多 SM 之間的原子性。
- 如果啟用硬體壓縮，資料傳輸時可大幅降低對外部記憶體頻寬的壓力，進一步提升效能。
想了解更多細節，可參考 NVIDIA 官方文件：Nsight Compute 硬體模型說明
Memory Chart
Memory Chart（記憶體圖）以圖形化邏輯方式顯示 GPU 上和 GPU 外記憶體子單元的效能數據，包括傳輸大小、命中率、指令或請求數等。
上篇文章我們簡單分析了下合併訪存程式碼的記憶體圖，下面我們來跟隨官方文件一起來看看 Memory Chart 的詳細分析：
邏輯單元（綠色）
邏輯單元指的是 CUDA 程式執行時會用到的各種”虛擬”記憶體類型，它們在示意圖中以綠色表示：
- Kernel：在 GPU 的串流多處理器 (SM) 上執行的 CUDA 核心程式。
- Global Memory：CUDA 的”全域記憶體”，是所有執行緒都能讀取和寫入的主要記憶體空間。
- Local Memory：CUDA 的”局部記憶體”，只供單一執行緒使用，保存該執行緒的私有資料。
- Texture Memory：CUDA 的”紋理記憶體”，適合做特殊資料存取(例如：影像處理中的取樣）。
- Surface Memory：CUDA 的”表面記憶體”，支援 2D/3D 資料結構的高效存取。
物理單元（藍色）
物理單元指的是 GPU 上真正存在的硬體模組，圖中以藍色表示：
- L1/TEX Cache：一級快取和紋理緩存，是 GPU 用來加速記憶體存取的重要元件。
- Shared Memory：CUDA 程式設計師可以自行管理的”共享記憶體”，位於每個 SM 上，速度比全域記憶體快。
- L2 Cache：二級快取，用來快取大量數據，減少對慢速系統記憶體的存取。
- L2 Compression：二級快取的”記憶體壓縮單元”，可透過資料壓縮減少記憶體頻寬佔用。
- System Memory：GPU 之外的系統內存，也就是 CPU 的記憶體，存取速度較慢。
- Device Memory：GPU 晶片本身的顯存，供執行 CUDA 程式使用。
- Peer Memory：當系統有多張 GPU 時，其他 GPU 的顯存可以當作”對等記憶體”來存取。
Links 連線
在 CUDA 效能分析圖中，Kernel（運算核心）和其他單元之間的連線，代表這些單元之間執行的指令數量。例如，Kernel 和 Global Memory 之間的連線，表示從全域記憶體讀取或寫入資料的指令。如果使用 NVIDIA A100 GPU，還會額外顯示「Load Global」或「Store Shared」等指令，因為它們對記憶體或快取的訪問行為，和一般全域存取或共享記憶體操作不同。
當你看到綠色（邏輯單元）和藍色（物理單元）之間的連線，這代表相關指令所產生的記憶體請求數。例如，從 L1/TEX Cache 到 Global Memory 的連線，會顯示全域記憶體載入指令所發出的請求數量。
每條連線的顏色表示該連線的使用率──顏色越接近紅色，代表接近該路徑的峰值性能；如果是灰色，代表這條連線當前沒有活躍傳輸。右側的顏色圖例會告訴你 0% 到 100% 的利用率分布，圖例左側的三角形標記，對應每條線在圖中的位置，幫助你更準確估算性能使用情況。
Ports 端口
每個單元通常都有一個共用的資料端口，負責所有資料的輸入與輸出。即使連線本身看起來沒到達各自的峰值性能，但當多條連線同時共享一個端口時，端口本身還是可能達到最大利用率。
在分析圖中，單元內的彩色小方塊表示端口的利用率：
顏色同樣採用 0% 到 100% 的顏色漸變，對應右側圖例。
如果端口不活躍，會顯示為灰色。
圖例左側的標記也能幫助你判斷各端口的使用情況。
透過觀察端口利用率，你可以快速發現是否因端口瓶頸，導致整個運算過程的效能受限。
Metrics 指標
可使用 –set full、–section MemoryWorkloadAnalysis_Chart 或 –metrics group:memory__chart 命令列收集該圖表的各項指標。下圖展示了記憶體表（Memory Tables）中報告的峰值與記憶體圖表（Memory Chart）中連接 Port 之間的對應：
Memory Tables（內存表）
在上一節中，我們簡單看了內存圖（Memory Chart），接下來來認識 內存表（Memory Tables）。
內存表會詳細列出各種GPU內部的記憶體單元，包括共享記憶體、不同層級的快取（L1/L2）以及顯示卡的裝置記憶體等，幫助你清楚了解每個環節的效能數據。
在表格中：
大多數欄位可以把滑鼠移到格子上，就會跳出指標的名稱與詳細解釋。
有些欄位不是直接統計得來，而是由其他指標計算出來的「衍生指標」，這些會直接顯示計算結果，不會有名稱。
如果某個指標在該次分析中沒被使用到，會在提示中看到「UNUSED」。
你也可以將滑鼠移到表格的行標題或列標題上，查看該區塊的說明，幫助你了解每部分代表什麼。
以上資訊皆參考自 NVIDIA 官方文件：NVIDIA Nsight Compute Profiling Guide – Memory Tables。
Shared Memory
Instructions
每個 warp 實際執行的共享記憶體指令總數，代表程式中需要多少次共享記憶體讀取或寫入操作。
Requests
共享記憶體的總請求次數。在Volta（SM 7.0）及以後的架構中，每條共享記憶體指令都會產生一次請求。
Wavefronts
執行完一個 warp 所有共享記憶體請求所需要的時間週期數。因為共享記憶體的請求需要序列化處理（不能一次做完），所以 Wavefronts 表示完成一組32條執行緒（warp）共享記憶體讀寫所花的時鐘週期。
簡單來說：Wavefronts 可以想像成「共享記憶體完成一次warp訪問需要幾個處理波段」，數值越高代表需要越多時鐘週期來完成這次warp的共享記憶體操作。
% Peak
共享記憶體使用率相對於硬體理論峰值的百分比。越接近100%，說明單元利用率高，代表程式可能在共享記憶體上形成瓶頸，但不一定代表執行效率高，因為也可能是有大量衝突或不良訪問模式。
Bank Conflicts
當warp中多個執行緒同時訪問共享記憶體，若它們的地址對應到相同的memory bank但偏移不同，就會形成「bank衝突」。此時硬體需將衝突的訪問拆分成多次非衝突的請求，導致共享記憶體有效帶寬下降，延遲增加。
關於 Wavefronts 與 Instructions/Requests 的關係
提到表格中的 Wavefronts 值看起來是「Instructions / Requests 的32倍」，其實這不是巧合，而是因為每個 warp 包含 32 條執行緒：
每個 request 是一次 32 執行緒的統一請求
Instructions 記錄每條執行緒各自的指令數
Wavefronts 表示完成一個 warp 所有請求需要經過多少序列化處理的時鐘週期
因此如果共享記憶體訪問沒有衝突且充分並行化，Wavefronts 應該接近 Instructions/32；但若存在bank衝突或訪問不良，Wavefronts 就會拉長，因為需要額外週期來完成被拆分的訪問。
一句話總結
Wavefronts 表示 GPU 要花多少時間序列化處理一個 warp 對共享記憶體的訪問，數值大小能反映訪問效率是否良好。
什麼是 Wavefront？
在 NVIDIA GPU 中，一個 warp 是最小的平行執行單位，包含 32 條執行緒。編譯器會把 warp 視為一起排程、同時存取記憶體的單位。在 NCU（NVIDIA Compute Profiler）裡，這個 warp 的行為被稱作 wavefront，意思是「同時被GPU調度、一起執行的32個執行緒」。
Requests 與 Wavefronts 有什麼差別？
Requests：你的 kernel 在程式中，理論上發出了多少次「warp級別」的共享記憶體（SMEM）訪問請求。
Wavefronts：GPU實際上處理這些請求時，總共啟動了多少次 warp 來完成訪問。
如果共享記憶體訪問完全沒有 bank 衝突，每個 warp 的一次訪問只要一個 cycle（時脈週期）就能完成，這時 Wavefronts 與 Requests 數量會相等。
為什麼表中出現 Wavefronts = Requests × 32？
當你在 kernel 中的存取模式非常不理想，出現了最嚴重的 bank 衝突：
SMEM 有 32 個 bank，理想情況是 32 條執行緒各自訪問不同的 bank，所有執行緒一次完成。
但如果 32 條執行緒全都同時去搶同一個 bank，這時 GPU 只能一次處理一條執行緒的訪問，其他必須依序排隊。
這會造成：
- 每個 warp 本該 1 個 cycle 完成的請求，要拉長到 32 個 cycle才能處理完全部 32 條執行緒。
- NCU 會顯示公式：Wavefronts = Requests + Bank Conflicts
- 如果完全衝突，就會測到 Wavefronts = Requests × 32，因為每個 warp 內的32條執行緒都需要各自分開的 cycle。
L1/TEX Cache
CUDA 記憶體存取分析指標簡明解說
Instructions（指令總數）：每個 warp 執行的所有匯編指令總數，可用來衡量計算工作量。
Requests（L1 請求數）：針對每條指令產生的 L1 請求總數，例如每個全域記憶體加載指令都會發出一次請求。如果請求數量遠高於指令數，可能表示記憶體訪問合併效率不佳。
Wavefronts（記憶體響應批次）：為完成所有請求所需的「wavefronts」數量。wavefronts 是GPU處理這些記憶體操作的執行批次，會被依序分批處理。
Wavefronts % Peak（處理單元利用率）：顯示處理wavefronts時使用單元的峰值利用率百分比。若數值接近100%，代表硬體幾乎已滿載，可能成為效能瓶頸。
Sectors（L1扇區訪問數）：所有 L1 請求訪問到的 L1 扇區總數。每次加載或儲存都會至少訪問1個扇區，原子或歸約操作通常會額外穿透到 L2。
Sectors/Req（每次請求訪問的平均扇區數）：平均每次 L1 請求需要訪問的扇區數。這個指標能反映訪問合併（Coalescing）程度：比率越低表示記憶體訪問更集中、更有效率；比率高則代表每個執行緒訪問的位置分散，導致效率低下。舉例：
- 32-bit 每次請求的理想值是 4
- 64-bit理想值是 8
- 128-bit理想值是 16
若比率遠高於這些值，代表訪問是「uncoalesced」（不合併），造成 L1 記憶體流量增加，影響效能。
Hit Rate（L1命中率）：記憶體請求直接命中L1快取的比例。命中率高代表大多數請求可快速在L1完成，降低延遲；反之，未命中的請求必須訪問L2，增加延遲。
Bytes（總請求位元組數）：請求從 L1 中讀取的資料總大小。每個 L1 扇區固定 32 bytes，因此 Bytes = Sectors × 32。
Sector Misses to L2（L1未命中數）：未命中 L1 後需要進一步訪問 L2 的扇區總數。高未命中率代表需要從更慢的 L2 取得資料，影響效能。
% Peak to L2（L1到L2的峰值利用率）：代表 L1 到 L2 資料通道的繁忙程度。若數值高，可能是大量寫入、原子操作或歸約操作導致L2成為效能瓶頸。
Returns to SM（返回到SM的資料包數）：從L1回傳到SM（Streaming Multiprocessor）的資料包數量。若資料請求範圍大，返回的包數也會增加。
% Peak to SM（L1回傳到SM的峰值利用率）：顯示 XBAR 到 L1 返回路徑的使用率。高比例代表主要瓶頸在讀取操作上，需考慮提高合併度或提升 L1 命中率來降低壅塞。
補充範例：
在一個測試中，如果全域加載平均每次請求32個sectors，代表32個執行緒各自訪問不同的sector，完全沒有合併，導致低效能。理想情況是同一warp中的多個執行緒訪問同一cache line，減少L1流量、提升效能。
如何判讀 Rows（資料列）
Access Types（訪問類型）：依記憶體操作區分，例如「全域讀取」「紋理訪問」「surface歸約」等。
Loads（所有讀取訪問的總和）：顯示該指標中所有讀取操作的合計。
Stores（所有寫入訪問的總和）：顯示該指標中所有寫入操作的合計。
Total（讀取與寫入的合計）：將 Loads 與 Stores 相加後的總數。
L2 Cache Eviction Policies
Columns（欄位說明）
First（evict_first 策略）
表示使用 evict_first 策略時，二級快取（L2 Cache）中被存取的區塊（sector）數量。這種策略適合「一次用完就丟」的流式資料，因為資料會排在快取清單的最前面，需要清理快取時最先被清掉。First Hit Rate
使用 evict_first 策略時的快取命中率，代表你在二級快取中再次讀取同一塊資料時的成功率。命中率高表示多次重用，低表示資料是一次性使用。Last（evict_last 策略）
表示使用 evict_last 策略時，在二級快取中被存取的區塊數量。這種策略把資料排在清單最後，表示這些資料最好保留得久一點，通常用於需要長時間留在快取裡的關鍵資料。Last Hit Rate
使用 evict_last 策略時的快取命中率，能顯示需要長期保留的資料在快取中的使用效益。Normal（evict_normal 策略）
表示使用預設 evict_normal 策略時，在二級快取中被存取的區塊數。這是大部分資料使用的標準策略，會根據資料訪問模式決定何時驅逐資料。Normal Hit Rate
使用 evict_normal 策略時的快取命中率。Normal Demote（evict_normal_demote 策略）
表示使用 evict_normal_demote 策略時，在二級快取中被存取的區塊數量。這種策略適合先留在快取中，但當空間緊張時會比 evict_normal 更容易被降級或驅逐。Normal Demote Hit Rate
使用 evict_normal_demote 策略時的快取命中率。
Rows（列說明）
Access Types（訪問類型）
指各種資料訪問來源，例如來自 L1 快取的加載（Load）或運算（Reduce）操作。L1/TEX Total
所有從 L1 Cache 或 Texture Cache 發出的操作總數，反映 L1 階段的存取頻率。L2 Fabric Total
如果你的 GPU 有多個 L2 分區，L2 Fabric 就是這些分區彼此間的高速連接；這個欄位表示 L2 Fabric 上所有的存取總數。只會在支援 L2 Fabric 的 CUDA 設備、且執行含 Fabric 通訊的 kernel 時出現。GPU Total
所有二級快取（L2 Cache）用戶端（像是 L1 Cache、Texture 單元、Tensor Core 等）的總操作次數，加總出 GPU 整體對 L2 的訪問量。
Device Memory
Columns：
Sectors：每種存取方式向設備記憶體發出的「區塊請求」總數。數字越大，表示該類操作對記憶體的使用頻繁。
% Peak：該存取單元的「使用峰值百分比」。數值越高，表示該部分的使用量接近設備所能承受的最高負載，可能代表潛在的效能瓶頸，但高數值並不一定意味效率好。
Bytes：記錄二級快取與設備記憶體之間傳輸的資料總量，單位是位元組（Bytes），數字越大，表示資料交換頻繁。
Throughput：代表設備記憶體的資料吞吐量（每秒能傳輸多少位元組），數值越高表示該單元在資料讀寫上的活動更密集。
Rows：
Access Types：各種記憶體讀取或寫入的操作類型，例如資料載入（load）、儲存（store）等。
Total：將同一欄位裡所有操作類型的數據加總，方便你總覽整體存取情況。
add1 核函數快取 & 訪存分析
前面我們跟隨 NVIDIA 官方文件簡單學習了 L1/L2 Cache 以及 NCU 分析後提供的 Memory Chart 和 Memory Tables
下面我們就簡單看下我們自己寫的 add1 核函數經過 NCU 分析後的記憶體圖表，程式碼如下：
void __global__ add1(float* x, float* y, float* z){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
z[idx] = x[idx] + y[idx];
}
Memory Chart 如下：
全域記憶體存取次數
在執行這個加法 kernel 時，總共發出了大約 78.6 萬次全域記憶體指令，其中52.4萬次是讀取（load），26.2萬次是寫入（store）。
這表示程式對全域記憶體的讀寫非常頻繁，是影響效能的第一個重點。
L1/TEX Cache 完全沒命中
L1/TEX 快取的命中率是 0%，代表這些讀寫指令完全沒有透過L1快取就取得資料，也就是每一次記憶體讀寫都需要直接「穿透」到下一層的L2 Cache。
換句話說，L1 Cache在這次運算中沒有發揮任何加速效果。
從 L1 轉到 L2 的請求量
僅就讀取（load）而言，52.4萬次讀取指令，因為每次會以128位元組（128B）的單位進行訪問，總共向L2 Cache發送了大約67.1 MB的資料請求量。
這數量說明在L1失效後，L2 Cache承擔了龐大的資料存取負擔。
L2 Cache 部分命中
L2快取在這67.1MB的資料請求中，大約有33.4%被命中（成功在L2取得資料），剩下66.6%的請求沒命中，只能往下送到DRAM。
這個命中率表示L2雖有幫助，但還不足以大幅減少存取延遲。
未命中部分送到DRAM
L2快取沒命中的那些請求，最後會被轉發到片外的DRAM（顯示卡的記憶體），這是最慢的存取路徑，也容易成為效能瓶頸來源。
簡單來說：
因為 L1 快取完全沒幫上忙，大部分資料流直接從 L2 或 DRAM 取得，導致每次存取都要花比較長的時間完成，這會大幅降低 kernel 效能。想要優化 CUDA 程式，就需要讓資料能更好地命中快取，減少每次存取的延遲。
接著來看下 Memory Tables，先看 L1 Cache：
透過這張 L1/TEX Cache 表，我們可以清楚了解這段 kernel（add1）在訪問顯存時的行為：
指令與請求數量
Global Load（全域讀取）：
執行了 52 萬次 warp 級別的 load 指令
發出了 52 萬次合併後的記憶體請求
Global Store（全域寫入）：
執行了 26 萬次 store 指令，對應同樣 26 萬次請求
沒有使用到其他 Local、Texture、Shared 等記憶體類型，代表這個 kernel 純粹在做全域記憶體的讀寫。
Wavefronts 與請求對應
每次 warp 發出的請求與 wavefront 數量完全相符（Load 和 Store 的 Wavefronts = Requests），代表訪問非常「整齊」，完全沒有發生記憶體 bank 衝突，每個 warp 都能一次性同時完成訪問。
合併訪問（Coalescing）效果 – 每個請求跨4個 sector
表中顯示：
Load 讀取了 209 萬個 sectors
Store 寫入了 104 萬個 sectors
由於每個請求跨了 4 個 sectors（Sectors/Req = 4），說明 warp 訪問的記憶體範圍非常寬，但又剛好與 warp 自身 128 bytes (32條4B float) 對齊，達成「完美合併訪問」。
具體資料量：
Load 共讀取了 67MB
Store 共寫入了 34MB
每個 sector 大小是 32 bytes，所以4個 sectors = 128 bytes，完美符合一個 warp 訪問需求。
L1 完全未命中
表中 Hit Rate = 0%，代表所有 Load 和 Store 的請求在 L1 完全 miss，全部直接轉給 L2 cache。
Sector Misses to L2 的數量與 sectors 完全相符（Load + Store 共314萬次 miss），證明沒有任何請求在 L1 命中。
L1/TEX 帶寬利用率極低
Load 只使用了 L1 峰值帶寬的 4.6%，Store 使用了 2.3%，加起來不到 7%，代表L1沒有被大量使用或造成瓶頸。
L2 負載情況
Load 佔用了 L2 最大帶寬的18.5%，Store也用了9.3%，總計快要28%。
回傳給SM的 sectors 數量為 78萬次，代表大量資料最終還是要回到SM進行運算。
簡單總結：
這個 kernel 只做單純的全域讀寫，沒有用到L1/TEX的快取命中，所有請求都直接跑到L2。
記憶體訪問非常理想，warp 訪問完全合併（每次請求剛好跨4個sector、128 bytes對齊），且完全沒有 bank 衝突。
L1 利用率很低（不到7%），大部分流量都落到L2，後續部分又必須從DRAM補足，顯示記憶體訪問雖然整齊，但若能提高L1命中率，效能還能再進一步優化。
透過這張表，我們可以非常細緻地看到每次 warp 訪問在 L1 發生了什麼、拿了多少 sector、衝不衝突、用了多少頻寬，以及 miss 後打到 L2 的流量情況。
最後我們來看下 L2 Cache 和 Device Memory：
從上面這張 L2 Cache + Device Memory 的表裡，我們可以看出 add1 在二級快取和片外記憶體上的整個資料流特徵：
L2 Cache 層級
類別 | 請求(warp) | 扇區數 | 扇區/Req | Hit Rate | Misses→Device | % Peak |
|---|---|---|---|---|---|---|
L1/TEX Load | 524 288 | 2 097 152 | 4 | 0 % | 2 097 152 | 30.37 % |
L1/TEX Store | 262 144 | 1 048 576 | 4 | 100 % | 0 | 15.18 % |
總和 | 786 432 | 3 145 728 | 4 | 33.33 % | 2 097 152 | 45.55 % |
在這次 CUDA 運算中，L2 Cache 負責快取記憶體的讀寫，以下是觀察到的幾個關鍵重點：
讀取（Load）全部未命中
一共有 52 萬次 warp 發出的讀取請求（共 2,097,152 個資料扇區，約67MB），這些資料完全沒命中L2 Cache，直接從DRAM（顯示卡的主記憶體）中讀取。
這代表讀取效率不佳，會導致較高的延遲。
寫入（Store）全部命中
有 26 萬次 warp 發出的寫入請求（共 1,048,576 個資料扇區，約33MB），完全在L2 Cache就命中，沒發生Miss。
不過，這些資料在之後還是會被「正常驅逐」（Eviction）機制從L2 Cache寫回到DRAM。
整體命中率
綜合 Load 和 Store，總請求數是78萬次，其中 Store佔了33.33%，恰好就是整體的命中率（因為Store全命中而Load全沒命中）。
L2 頻寬使用情況
整體 L2 頻寬利用率約 45.6%：其中30.4%用於處理 Load，15.2%用於 Store。
表示 GPU 的 L2 頻寬還有超過一半的餘裕，未被完全利用。
驅逐策略
無論是讀取 Miss 還是寫入 Hit，L2 都採用「正常驅逐」(Normal Eviction)來將資料移到 DRAM，沒有使用特殊的逐出策略。
Device Memory 層級
類型 | 扇區數 | 位元組量 | % Peak | 吞吐 (B/s) |
|---|---|---|---|---|
Load | 2 097 160 | 67 109 120 B | 62.58 % | 216 536 912 751.68 |
Store | 1 013 704 | 32 438 528 B | 30.25 % | 104 667 423 851.32 |
總和 | 3 110 864 | 99 547 648 B | 92.83 % | 321 204 336 602.99 |
DRAM 訪存流量與帶寬概況
1️⃣ 在這個 add1 kernel 中：
DRAM 讀取流量大約是 67 MB（2,097,160 個扇區），
寫入流量大約是 32 MB（1,013,704 個扇區），
總共加起來大約是 99.5 MB。
2️⃣ 頻寬利用率非常高，達到 92.8%，代表整個 kernel 幾乎把片外記憶體（DRAM）的頻寬用到極限。
3️⃣ 實際的數據傳輸速度也很驚人：
DRAM 讀取頻寬大約 216 GB/s，
寫入頻寬大約 105 GB/s，
總頻寬大約 321 GB/s。
問題核心簡單說
add1 中的資料讀取幾乎完全 跳過 L2 快取（L2 命中率是 0%），資料直接從 DRAM 讀取，造成主要的讀取流量。
資料寫入雖然在 L2 快取有命中，但最後還是因為快取驅逐（被擠出）寫回到 DRAM，因此也產生了約 32 MB 的寫入流量。
頻寬使用情況
二級快取（L2）的頻寬只用了大約 45%，
而 DRAM（片外記憶體）幾乎被用到 93% 的極限，
這顯示此 kernel 的主要瓶頸就是 DRAM 頻寬不足。
如何改善效能？
如果能透過 資料重用 或 分塊（tiling）技術，讓 L2 對讀取資料的命中率提升，就能大幅減少DRAM的負擔、降低頻繁訪問外部記憶體，進而有效提升運算效能。
結論
add1 這個 kernel 最大的效能瓶頸就是 DRAM 帶寬快要被用滿了，若能想辦法提升 L2 快取的命中率，就能讓效能有顯著進步。
參考資料
CUDA優化常見問題
透過優化記憶體配置與採用頁面鎖定技術，能顯著減少 Host 到 Device 的資料交換延遲。
CUDA 流 是執行緒的抽象化，可實現多指令同時發送到 GPU，提升運算效率。
關於 TAKI Cloud
TAKI Cloud 提供專業的 AI 伺服器租用、GPU 雲端運算並專注於AI伺服器與GPU雲端服務，提供符合國際標準的中華電信 IDC 機房與 24 小時技術支援，協助國內外 AI 團隊快速布署最新大模型，具備多年 A100 / H100 環境管理與支援經驗。
官方網站：www.taki.com.tw
本文作者：張亮
TAKI Cloud 資深系統架構師，專精於 AI 運算優化、CUDA 調校與高效能 GPU 叢集架構，擁有 10 年以上經驗，曾協助多家 AI 新創完成模型加速專案。
RTX - 3090 GPU 主機
視頻渲染、科學模擬和機器學習
支援 DeepSeek-R1 32B
實例
8卡 NVIDIA RTX-4090 24G
數量
1
全台唯一提供高階 AI / GPU 主機租用
價格殺很大 / 量大可談
RTX - 4090 GPU 主機
視頻渲染、科學模擬和機器學習
支援 DeepSeek-R1 70B
實例
8卡 NVIDIA RTX-4090 24G
數量 庫存緊張，欲租從速
1
全台唯一提供高階 AI / GPU 主機租用
價格殺很大 / 量大可談
HGX H100 GPU 主機
原價 499,999元/月 特惠價 450,000元/月
支援 DeepSeek-R1 671B 滿血版
實例
8顆 NVIDIA HGX H100 80G
數量
1
全台唯一提供高階 AI / GPU 主機租用
價格殺很大 / 量大可談
