Unified Memory：平台差異與效能提示 (Platforms and Performance Hints)

重點總覽

延續 04-CUDA-Features/01-Unified-Memory-Full-Support，本筆記聚焦於「非 full support」平台（只有 managed memory 支援的裝置，以及 Windows/WSL/Tegra 上 concurrentManagedAccess = 0 的裝置）的行為限制，以及適用於所有 unified memory 的 效能提示 (performance hints)。

4.1.2 僅支援 Managed Memory 的裝置

針對 compute capability 6.x 或更高、但沒有 pageable memory access 的裝置（見 Overview of Unified Memory Paradigms）：

• CUDA managed memory 完整支援且 coherent。
• 但 GPU 不能存取 system-allocated memory（malloc/new/mmap 配置的記憶體）。
• programming model 與效能調校大致與 full support 相同，唯一例外：不能用 system allocator 配置 unified memory，必須用 cudaMallocManaged 等 CUDA API。

4.1.3 Windows / WSL / Tegra 上的 Unified Memory

4.1.3.1 Multi-GPU

• managed 配置透過 GPU 的 peer-to-peer (P2P) 能力，自動對系統中所有 GPU 可見。
• 行為類似 cudaMalloc()：當前 active device 是實體配置的 home，其他 GPU 經 PCIe 以降低的頻寬存取。

   active device = home (full bandwidth)
        ┌──────────┐
        │  GPU0    │ ← managed 實體配置在此
        └────┬─────┘
   PCIe (降頻) │   PCIe (降頻)
        ┌────┴─────┐  ┌──────────┐
        │  GPU1    │  │  GPU2    │  ← 經 P2P over PCIe 存取
        └──────────┘  └──────────┘

• Linux：只要所有「正在使用」的 GPU 都有 P2P，managed memory 就留在 GPU memory。一旦程式開始使用某個與其他持有 managed 配置的 GPU 沒有 P2P 的 GPU，driver 會把所有 managed 配置遷回 system memory，此時所有 GPU 都受 PCIe 頻寬限制。
• Windows：若 peer mapping 不可用（例如不同架構的 GPU 之間），系統自動 fallback 到 mapped memory，不論兩個 GPU 是否真的都被用到。若實際上只用一個 GPU，需在啟動前設定 CUDA_VISIBLE_DEVICES 限制可見 GPU，才能讓 managed memory 配置在 GPU memory。
• Windows 替代方案：設 CUDA_MANAGED_FORCE_DEVICE_ALLOC 為非零值，強制 driver 一律用 device memory 作為實體儲存。此時 process 中所有支援 managed memory 的裝置必須彼此 P2P 相容，否則回傳 cudaErrorInvalidDevice（即使已對該裝置呼叫 cudaDeviceReset 也一樣）。

4.1.3.2 Coherency and Concurrency

為確保 coherency，當 CPU 與 GPU 同時執行時，programming model 對資料存取加上約束：

• 任何 kernel 執行期間，GPU 獨佔所有 managed data，不論該 kernel 是否實際用到那塊資料。
• CPU 同時存取（即使是不同的 managed 配置）會造成 segmentation fault，因為該 page 被視為對 CPU 不可存取。

__device__ __managed__ int x, y = 2;
__global__ void kernel() { x = 10; }
int main() {
    kernel<<<1, 1>>>();
    y = 20;                    // 在不支援 concurrent access 的 GPU 上出錯
    cudaDeviceSynchronize();
}

上例在 CC 6.x（有 GPU page faulting）可成功，但在 pre-6.x 與 Windows 上失敗，因 CPU 觸碰 y 時 kernel 仍在執行。正確寫法是先同步再存取：

int main() {
    kernel<<<1, 1>>>();
    cudaDeviceSynchronize();   // 確保 GPU 完工
    y = 20;                    // 成功
}

• 任何「邏輯上保證 GPU 完工」的呼叫都有效（見 Explicit Synchronization）。
• 若在 GPU 仍 active 時用 cudaMallocManaged()/cuMemAllocManaged() 動態配置，其行為在「再啟動工作或同步 GPU」之前未定義，期間 CPU 存取可能（也可能不）segfault。例外：用 cudaMemAttachHost / CU_MEM_ATTACH_HOST flag 配置者不受此限。

4.1.3.3 Stream-Associated Unified Memory

stream 表達 kernel launch 間的依賴與獨立關係：同一 stream 內保證連續執行，不同 stream 間允許並行（見 02-Programming-GPUs/14-Async-Streams-and-Events）。

Stream callbacks：在 stream callback 內 CPU 可存取 managed data，前提是沒有其他可能存取 managed data 的 stream 仍 active。後面不接 device work 的 callback 還可當同步點（例如在 callback 內 signal condition variable）。

四個重點：

1. CPU 隨時可存取非 managed 的 mapped memory，即使 GPU active。
2. 只要有任何 kernel 在跑，GPU 即視為 active，即使該 kernel 沒用到 managed data；只要 kernel 可能用到，存取就被禁止。
3. 對 managed memory 的並行 inter-GPU 存取，限制與非 managed 的 multi-GPU 存取相同。
4. 並行 GPU kernel 存取 managed data 沒有額外約束——也就是允許 kernel 間 race（與非 managed memory 現況相同）。從 GPU 角度看，managed 與非 managed memory 行為一致。

綁定 stream 以取得更細粒度控制：

cudaError_t cudaStreamAttachMemAsync(cudaStream_t stream, void *ptr,
                                     size_t length = 0, unsigned int flags = 0);

• 把 ptr 起算 length bytes 綁到 stream：只要該 stream 的所有操作完成，CPU 就可存取這塊記憶體，不論其他 stream 是否 active。
• 等於把「active GPU 對 managed 區的獨佔」從 whole-GPU 縮小為 per-stream。
• 未綁定 stream 的配置對所有 stream 的 kernel 可見（cudaMallocManaged 與 __managed__ 的預設）——這正是「任何 kernel 在跑時 CPU 不可碰」規則的由來。

預設 (global visibility)            綁定 stream 後
┌─────────────────────────┐        ┌─────────────────────────┐
│ 任一 kernel active       │        │ 只有 stream1 active 才   │
│  → CPU 全面禁止存取      │        │   禁止；其他 stream 跑   │
│                         │        │   時 CPU 仍可存取        │
└─────────────────────────┘        └─────────────────────────┘

多執行緒 host 程式：每個 CPU thread 建立自己的 stream（用 NULL stream 會在 thread 間造成依賴），並用 cudaStreamAttachMemAsync 把該 thread 的 managed 配置綁到自己的 stream，通常一生不變。

void run_task(int *in, int *out, int length) {
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    int *data;
    cudaMallocManaged((void**)&data, length, cudaMemAttachHost); // 初始 device 不可見
    cudaStreamAttachMemAsync(stream, data);                      // 綁到私有 stream
    cudaStreamSynchronize(stream);
    // ... host 與 device 交替使用 data ...
}

• cudaMemAttachHost flag 讓配置初始對 device-side 執行不可見，避免「配置到取得」之間被別的 thread 的 kernel 意外視為 in-use。
• 替代方案是配置後跨所有 thread 做 process-wide barrier；但 stream 銷毀會讓配置還原為預設可見性，故銷毀前還需第二道 barrier。cudaMemAttachHost 正是為簡化此流程而存在。

Stream-associated UM 的資料搬移（concurrentManagedAccess = 0 時）：

當 memcpy/memset 從 device 存取資料時，該 stream 視為 GPU active；期間 CPU 存取「綁該 stream」或「global visibility」的資料會 segfault（若 concurrentManagedAccess = 0）。

4.1.4 Performance Hints（效能提示）

效能提示給 CUDA 更多 unified memory 用法資訊，讓 driver 更有效率地管理 managed memory。

4.1.4.1 Data Prefetching（資料預取）

cudaMemPrefetchAsync 是 async、stream-ordered API，可把資料遷到指定 processor 附近：

cudaError_t cudaMemPrefetchAsync(const void *devPtr, size_t count,
                                 struct cudaMemLocation location,
                                 unsigned int flags, cudaStream_t stream = 0);

• [devPtr, devPtr + count) 在 stream 中執行 prefetch task 時遷到：location.type == cudaMemLocationTypeDevice 時遷到 location.id 指定的 GPU；cudaMemLocationTypeHost 時遷到 CPU。
• 遷移在 stream 中所有前序操作完成後才開始，並在後續操作前完成；prefetch 進行中仍可存取資料。
• 用法：kernel 使用前先 prefetch 到 GPU，使用後再 prefetch 回 CPU。system allocator（malloc）與 managed（cudaMallocManaged）配置寫法一致。

4.1.4.2 Data Usage Hints（cudaMemAdvise）

當多個 processor 同時存取同一資料時，用 cudaMemAdvise 提示存取模式：

cudaError_t cudaMemAdvise(const void *devPtr, size_t count,
                          enum cudaMemoryAdvise advice,
                          struct cudaMemLocation location);

• 三者各有對應的 Unset 版本：cudaMemAdviseUnsetReadMostly、cudaMemAdviseUnsetPreferredLocation、cudaMemAdviseUnsetAccessedBy。

cudaMemLocation loc = {.type = cudaMemLocationTypeDevice, .id = myGpuId};
cudaMemAdvise(dataPtr, dataSize, cudaMemAdviseSetReadMostly, loc);
// 外層: CPU 每圈改寫資料 -> 對各 GPU prefetch (read duplication)
// 內層: kernel 只讀 dataPtr

4.1.4.3 Memory Discarding（記憶體丟棄）

cudaMemDiscardBatchAsync 告知 runtime 指定範圍的內容已無用。UM driver 會因 fault-based migration 或 oversubscription 的 eviction 自動搬移資料，這些搬移有時是多餘的、嚴重拖慢效能；標記為 discard 後，driver 在 prefetch 或 page eviction 時不需再搬移這塊資料。

cudaError_t cudaMemDiscardBatchAsync(void **dptrs, size_t *sizes, size_t count,
                                     unsigned long long flags, cudaStream_t stream);

• 對 dptrs/sizes 指定的多個範圍批次丟棄；兩陣列長度皆為 count。每個範圍必須是 cudaMallocManaged 或 __managed__ 的 managed memory。
• cudaMemDiscardAndPrefetchBatchAsync 合併 discard 與 prefetch，語意等同「cudaMemDiscardBatchAsync 後接 cudaMemPrefetchBatchAsync」但更省。用於「需要記憶體在目標位置、但不需要其內容」時。其 prefetchLocs 指定各 prefetch 目的地、prefetchLocIdxs 指出每個 prefetch location 套用到哪些操作（例：10 個操作前 6 個到 A、後 4 個到 B → numPrefetchLocs=2、prefetchLocIdxs={0, 6}）。

4.1.4.4 查詢資料使用屬性

cudaMemRangeGetAttribute 查詢經 cudaMemAdvise/cudaMemPrefetchAsync 設於 managed memory 的屬性（範圍須為 cudaMallocManaged 或 __managed__）：

cudaMemRangeGetAttribute(void *data, size_t dataSize,
                         enum cudaMemRangeAttribute attribute,
                         const void *devPtr, size_t count);

• 可用對應的 cudaMemRangeGetAttributes（複數）一次查多個屬性。

4.1.4.5 GPU Memory Oversubscription（超額配置）

unified memory 讓應用程式超額配置 (oversubscribe) 任一 processor 的記憶體：可配置並共享大於系統中任一 processor 記憶體容量的陣列，在不顯著增加程式複雜度下，支援放不進單一 GPU 的資料集做 out-of-core 處理。

考試/測驗重點

Related Notes

• 04-CUDA-Features/01-Unified-Memory-Full-Support
• 04-CUDA-Features/06-Stream-Ordered-Memory-Allocator
• 04-CUDA-Features/20-Virtual-Memory-Management
• 04-CUDA-Features/21-Extended-GPU-Memory
• 04-CUDA-Features/19-Interprocess-Communication
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/15-Async-Callbacks-Ordering-Graphs
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps