CUDA C++ 記憶體管理 (CUDA C++ Memory Management)
要在 GPU 上執行 vecAdd 這類 kernel，輸入/輸出陣列（A、B、C）必須位於 GPU 可存取的記憶體中。CUDA 提供多種做法，本筆記聚焦兩種最基本的方式：Unified Memory（由 driver 自動搬移）與 Explicit Memory Management（程式設計師手動配置與複製），並比較兩者在效能上的取捨。
重點總覽
| 項目 | 重點 |
|---|---|
| 問題本質 | kernel 要存取的資料必須在 GPU 可存取的 memory；兩種主要做法：unified vs explicit |
| Unified Memory | cudaMallocManaged（或 __managed__ specifier）配置；單一指標 CPU/GPU 皆可存取，driver 自動搬移 |
cudaFree |
unified 與 explicit 的 device buffer 都用同一個 cudaFree 釋放 |
| Explicit Management | cudaMalloc 配置 device memory，cudaMemcpy 手動搬移；host pointer 與 device pointer 分離 |
cudaMemcpy |
參數：dest、src、size(bytes)、cudaMemcpyKind_t；同步（複製完才返回） |
| Copy 方向 | cudaMemcpyHostToDevice / DeviceToHost / DeviceToDevice / Default（依指標自動判斷） |
| Page-locked memory | cudaMallocHost 配置 page-locked host buffer，提升複製效能、async 傳輸的必要條件 |
| 效能取捨 | explicit 較冗長但可精確控制搬移時機/位置，並與運算重疊；unified 較簡潔，可用 advise/prefetch 補回部分效能 |
2.1.3 Memory in GPU Computing（總覽）
要使用前面定義的 vecAdd kernel，陣列 A、B、C 必須放在 GPU 可存取的 memory 中。本章示範兩種做法，其餘（如更進階的 unified memory 機制）留待後續章節。
「資料要讓 GPU 看得到」的兩條路徑
┌─────────────────────────────┬─────────────────────────────┐
│ Unified Memory │ Explicit Management │
│ cudaMallocManaged │ cudaMalloc + cudaMemcpy │
│ → 單一指標, driver 自動搬移 │ → host/device 指標分離,手動搬 │
│ 簡潔、好寫 │ 冗長,但可精確控制與重疊 │
└─────────────────────────────┴─────────────────────────────┘
兩種方式並非互斥；可先用 unified memory 快速寫出正確版本，再針對熱點改用 explicit management（或加上 prefetch/advise）來調效能。
2.1.3.1 Unified Memory
Unified Memory 是 CUDA runtime 的功能，讓 NVIDIA Driver 負責管理 host 與 device(s) 之間的資料搬移。記憶體用 cudaMallocManaged API 配置，或用 __managed__ specifier 宣告變數。Driver 會確保：當 CPU 或 GPU 任一方嘗試存取該記憶體時，記憶體都是可存取的。
關鍵事實：
cudaMallocManaged配置的 buffer 可從 CPU 或 GPU 存取；使用 單一指標（不需要另外維護 host/device 兩份指標）。- 這些 buffer 用
cudaFree釋放（與 explicit 的 device memory 相同 API）。 - host 端可直接
initArray(A, ...)初始化，kernel 直接收同一個指標A，kernel 完成後 CPU 也可直接讀C，中間不需要手寫任何cudaMemcpy。 - 支援所有 CUDA 支援的 OS 與 GPU，但底層機制與效能會因系統架構而異。
void unifiedMemExample(int vectorLength)
{
float *A = nullptr, *B = nullptr, *C = nullptr;
// 用 unified memory 配置：CPU/GPU 皆可存取
cudaMallocManaged(&A, vectorLength * sizeof(float));
cudaMallocManaged(&B, vectorLength * sizeof(float));
cudaMallocManaged(&C, vectorLength * sizeof(float));
initArray(A, vectorLength); // host 直接初始化
initArray(B, vectorLength);
int threads = 256;
int blocks = cuda::ceil_div(vectorLength, threads);
vecAdd<<<blocks, threads>>>(A, B, C, vectorLength); // 同一指標傳給 GPU
cudaDeviceSynchronize(); // 等 kernel 完成
// CPU 直接讀 C 做比對，無需 cudaMemcpy
cudaFree(A); cudaFree(B); cudaFree(C);
}
重點：同一個指標 A/B/C 既給 host 初始化也給 kernel 使用，driver 自動保證資料在被存取時就在正確的位置。
例外：在某些 Linux 系統上（例如具備 address translation services 或 heterogeneous memory management 者），所有 system memory 自動就是 unified memory，此時無需使用 cudaMallocManaged 或 __managed__ specifier。
2.1.3.2 Explicit Memory Management
Explicit memory management 由程式設計師明確控制記憶體配置與 host/device 間的資料搬移。優點是可改善效能，缺點是程式碼更冗長。device memory 用 cudaMalloc 配置、用 cudaFree 釋放（與 unified 相同）。
關鍵事實：
- host pointer 與 device pointer 分離：
A/B/C指向 host buffer，devA/devB/devC指向 device buffer，兩者是不同位址空間的指標，不可互相直接 dereference。 - 用
cudaMallocHost配置 host buffer：這會配置 page-locked memory，是「buffer 將用於 CPU↔GPU 複製」時的 best practice；對應釋放 API 為cudaFreeHost。 - 用
cudaMemcpy在 host/device 間搬資料；cudaMemset可將 device memory 初始化為某值（例：cudaMemset(devC, 0, ...)）。 - kernel 收的是 device 指標
devA/devB/devC；算完後要再cudaMemcpy把結果devC拷回 host 的C才能在 CPU 讀。
// host buffer 用 page-locked memory（複製效能較佳）
cudaMallocHost(&A, vectorLength * sizeof(float)); // ... B, C 同理
initArray(A, vectorLength); initArray(B, vectorLength);
// device buffer
cudaMalloc(&devA, vectorLength * sizeof(float)); // ... devB, devC 同理
// host → device
cudaMemcpy(devA, A, vectorLength * sizeof(float), cudaMemcpyDefault);
cudaMemcpy(devB, B, vectorLength * sizeof(float), cudaMemcpyDefault);
cudaMemset(devC, 0, vectorLength * sizeof(float));
vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
cudaDeviceSynchronize();
cudaMemcpy(C, devC, vectorLength * sizeof(float), cudaMemcpyDefault); // device → host
cudaFree(devA); /* ... */ cudaFreeHost(A); /* ... */
重點：相比 unified 版本，多了 cudaMalloc + 兩段 cudaMemcpy（上傳輸入、下載結果），這正是「冗長但可控」的代價。
cudaMemcpy 與複製方向
cudaMemcpy 的參數依序為：destination pointer、source pointer、size in bytes、cudaMemcpyKind_t。cudaMemcpyKind_t 常見值：
| Kind | 方向 |
|---|---|
cudaMemcpyHostToDevice |
CPU → GPU |
cudaMemcpyDeviceToHost |
GPU → CPU |
cudaMemcpyDeviceToDevice |
GPU 內部，或 GPU↔GPU |
cudaMemcpyDefault |
由 source/dest 指標位址自動判斷複製類型 |
host buffer (pageable/pinned) device buffer (global memory)
A ──cudaMemcpyHostToDevice──▶ devA
B ──cudaMemcpyHostToDevice──▶ devB
C ◀──cudaMemcpyDeviceToHost── devC (kernel 寫完結果後回拷)
cudaMemcpy 是 同步（synchronous） 的：在複製完成前不會返回，會阻塞呼叫它的 host thread。非同步複製（async copy）另見 streams 章節 02-Programming-GPUs/14-Async-Streams-and-Events。
cudaMemcpyDefault 因為靠 unified virtual address space 由指標判斷方向，使用上最省心；但寫明確方向（HostToDevice 等）能讓意圖更清楚、也便於除錯。
Page-locked (pinned) host memory
cudaMallocHost配置的是 page-locked memory，可提升複製效能，且是 非同步記憶體傳輸的必要條件。- 一般 best practice：凡是會用於與 GPU 之間傳資料的 CPU buffer，都用 page-locked memory。
例外/陷阱：若過多 host memory 被 page-lock，某些系統上效能會反而下降。最佳做法是只對真正用於 GPU 收/送資料的 buffer 做 page-lock，而非全部 host memory。
2.1.3.3 Memory Management and Application Performance
如範例所見，explicit memory management 較冗長，需要程式設計師自行指定 host/device 間的複製。這同時是它的優點與缺點：
| 面向 | Explicit | Unified |
|---|---|---|
| 程式碼量 | 冗長（手寫 malloc/memcpy） | 簡潔（單一指標） |
| 搬移時機控制 | 由程式設計師精確控制 | 由 driver 自動決定 |
| 資料 residency / 配置位置 | 完全掌握「資料在哪、配置在哪」 | driver 管理 |
| 與運算重疊 | 可控制傳輸並與其他運算重疊 | 預設不易 |
| 調效手段 | 直接控制 transfer | 用 Memory Advise / Prefetch 提示 driver |
- explicit 的核心價值：對「何時在 host/device 間複製、記憶體駐留在哪、什麼配置在哪」有完整控制，因而能透過控制傳輸、並讓傳輸與其他運算重疊，創造效能機會。
- 使用 unified memory 時，CUDA 也提供一組 API（Memory Advise 與 Prefetch）對管理記憶體的 NVIDIA driver 給出提示（hints），讓 unified memory 也能拿回部分 explicit management 的效能優勢。
共通的效能原則：避免不必要的 host↔device 搬移。資料一旦在 GPU 上，盡量留在 GPU 連續做多個 kernel，再一次性拷回；無論 explicit 或 unified（搭配 prefetch）都遵循此原則。
反例（搬移過多） 正例（資料常駐 GPU）
H→D → kernel1 → D→H H→D → kernel1 → kernel2 → kernel3 → D→H
H→D → kernel2 → D→H （中間不來回搬）
H→D → kernel3 → D→H
每步都來回 PCIe/NVLINK 只進/出一次,傳輸可與運算重疊
考試/測驗重點
| 情境/關鍵字 | 答案 |
|---|---|
| 配置 CPU/GPU 皆可存取的單一指標 buffer | cudaMallocManaged（或 __managed__ specifier） |
| unified buffer 用什麼釋放？ | cudaFree（與 device memory 相同） |
| explicit：配置 device memory / 釋放 | cudaMalloc / cudaFree |
| explicit：配置 page-locked host memory / 釋放 | cudaMallocHost / cudaFreeHost |
cudaMemcpy 的最後一個參數型別 |
cudaMemcpyKind_t |
cudaMemcpy 參數順序 |
dest, src, size(bytes), kind |
| 讓 CUDA 依指標自動判斷複製方向 | cudaMemcpyDefault |
| GPU→CPU 的方向常數 | cudaMemcpyDeviceToHost |
| CPU→GPU 的方向常數 | cudaMemcpyHostToDevice |
| GPU 內或 GPU↔GPU | cudaMemcpyDeviceToDevice |
cudaMemcpy 是同步還是非同步？ |
同步；複製完才返回（會阻塞 host thread） |
| 非同步傳輸的前提 | host buffer 須為 page-locked（cudaMallocHost） |
| page-lock 太多會怎樣？ | 某些系統效能反而下降；只 page-lock 用於傳輸的 buffer |
unified memory 一定要 cudaMallocManaged 嗎？ |
例外：某些 Linux（ATS/HMM）所有 system memory 自動是 unified，無需呼叫 |
| unified memory 的調效 API | Memory Advise / Prefetch（給 driver hints） |
| explicit vs unified 的根本差別 | explicit 手動控制搬移/residency 但冗長；unified 由 driver 自動但較難精細調效 |
| host pointer 可否直接傳給 kernel（explicit）？ | 否；kernel 收的是 device pointer（devA 等），兩者位址空間分離 |
| 共通效能準則 | 避免不必要的 host↔device 搬移，讓資料常駐 GPU、傳輸與運算重疊 |
Related Notes
- 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
- 02-Programming-GPUs/03-CUDA-Cpp-Sync-and-Workflow
- 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
- 02-Programming-GPUs/05-CUDA-Python
- 02-Programming-GPUs/14-Async-Streams-and-Events
- 02-Programming-GPUs/16-Unified-and-System-Memory
- 02-Programming-GPUs/07-SIMT-Device-Memory-Spaces
- 01-Introduction-to-CUDA/04-GPU-Memory-Hierarchy
- 02-Programming-GPUs/Practice-Programming-GPUs
- 00-Dashboard/Exam-Traps