The Modern CUDA Toolbox in Practice: A Step-by-Step Optimization Walkthrough

現代 CUDA 工具箱實務：一步步優化示範

NVIDIA CUDA 仍是 GPU 加速運算的根基，從科學模擬到大規模 AI 訓練皆仰賴它。然而，要寫出正確、易維護且效能佳的 CUDA 程式並非易事：記憶體錯誤往往隱蔽、效能瓶頸若未配合適當的量測工具很難被發現，而自行實作的 GPU 演算法也很少能匹敵經過高度優化的函式庫。所幸，近年 CUDA 工具鏈已相當成熟，許多常見問題都有直接的解法。

1. 問題的起點與原始程式概況

範例程式從 CPU 讀入三張 RGB 圖片，將資料傳至 GPU，先執行 RGB→灰階 轉換，接著在每個 32×32 的區塊內以排序方式求取 中位數，最後把每個區塊的中位數結果拷回 CPU。程式使用 OpenMP 於 CPU 端平行化三張圖的處理流程，並在 GPU 端分別呼叫兩個 kernel（轉換與中位數計算）。

2. 第一步：利用 Compute Sanitizer 捕捉非法存取

執行時若出現「illegal memory access」訊息，第一個動作應是啟動 Compute Sanitizer（CUDA 的功能正確性檢測套件）。此工具即時指出在 0_base_error_example.cu 第 55 行發生 共享記憶體寫入越界，原因是程式使用全域索引寫入 block 內的 shared memory。

3. 第二步：改用 CCCL 的新索引 API

為避免此類索引錯誤，NVIDIA 在 CCCL（CUDA C++ 標準函式庫） 中加入了區分全域與 block 級索引的 API。配合 cuda::launch 取代傳統的 <<<grid, block>>> 呼叫方式，可在 kernel 內使用 cuda::grid::thread_idx、cuda::grid::block_idx 等安全索引，確保共享記憶體的存取僅限於 block 內的範圍。

cuda::launch(kernel, grid, block, args...);

在 kernel 內則改寫為：

auto idx = cuda::grid::thread_idx.x + cuda::grid::block_idx.x * blockDim.x;

如此即可根除越界寫入的根本原因。

4. 第三步：引入 cuda::std::span 提升可讀性與安全性

若不使用 cuda::std::span（類似 C++20 std::span 的 GPU 版）或其 n 維變形 cuda::std::mdspan，開發者必須自行管理指標與長度，極易產生錯誤。將原始的原始指標改為 cuda::std::span<float>，編譯器即可在編譯期檢查範圍，降低記憶體錯誤的機會。

5. 第四步：效能基準與逐步優化

在每一次小幅度的程式修改後，使用 nvprof 或 Nsight Compute 進行基準測試，觀察記憶體吞吐量、指令佔用率與執行時間。透過這些工具可以快速定位 記憶體共用衝突（shared memory bank conflict）或 指令流水線瓶頸，進一步調整 block 大小或使用向量化載入。

6. 結語：從安全到效能的完整循環

本次示範展示了從 錯誤偵測 → 安全索引 → 語意化容器 → 效能量測 的完整優化流程。隨著 CUDA 工具箱的持續演進，開發者只要善用 Compute Sanitizer、CCCL 新 API 以及 Nsight 系列分析器，就能在不大幅改寫原始演算法的前提下，將程式變得更安全、易維護且效能更佳。未來，隨著 CUDA Graph 與 Unified Memory 的進一步成熟，類似的工作流程將更加自動化，值得每位 GPU 程式設計師持續關注與實踐。