CPU/GPU 同步與完整流程 (Synchronization and Full Workflow)

重點總覽

Kernel Launch 非同步與為何需要同步 (2.1.4)

Kernel launch 相對於呼叫它的 CPU thread 是非同步的：launch 語句送出工作後立刻回傳，host 的控制流會在 kernel 完成執行之前繼續往下跑，甚至可能在 kernel 真正啟動之前就往下執行。因此若要在 host code 中保證 kernel 已完成才繼續（例如讀取結果），就必須加入同步機制。

最簡單的方式是 cudaDeviceSynchronize：它阻塞呼叫它的 host thread，直到 GPU 上所有先前發出的工作都完成為止。

vecAdd<<<blocks, threads>>>(A, B, C, vectorLength); // 非同步，立即回傳
cudaDeviceSynchronize();                            // 阻塞 host，等 GPU 全部做完
// 此後才能安全地讀取 C 的結果

• 非同步：launch 不等 kernel 跑完即回傳，讓 CPU 能與 GPU 重疊工作。
• cudaDeviceSynchronize 等待的是「所有先前發出的工作」，不限單一 kernel。
• 本章範例因為 GPU 上只有單一操作，用 cudaDeviceSynchronize 已足夠。

   CPU thread                 GPU
   ───────────                ───
   vecAdd<<<>>>()  ──launch──▶ [ kernel 開始排程/執行 ]
   (立即回傳)  │                        │
   ...可做其他事 │                        │  (CPU 與 GPU 並行)
   cudaDeviceSynchronize() ──┐          │
        (阻塞等待) ◀──────────┘  完成 ◀──┘
   讀取 C[] ← 此時結果才保證就緒

完整 vecAdd 端到端流程 (2.1.5)

本節把整個向量相加範例串起來，含 kernel、host code 與驗證工具函式。範例預設向量長度 1024，也可由命令列參數指定不同長度。共有兩個版本：Unified Memory 與 Explicit Memory Management，兩者用的是同一個 kernel。

共用的 kernel 與驗證工具

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength) {
    int workIndex = threadIdx.x + blockIdx.x * blockDim.x;
    if (workIndex < vectorLength)          // 邊界檢查，避免越界
        C[workIndex] = A[workIndex] + B[workIndex];
}

• 每個 thread 由 threadIdx.x + blockIdx.x*blockDim.x 算出唯一全域索引，各做一筆獨立加法。
• if (workIndex < vectorLength) 防止當 blocks*threads 大於資料量時越界存取。
• 驗證流程：CPU 用序列版 serialVecAdd 算同一題，再用 vectorApproximatelyEqual（帶 epsilon=0.00001 容差）比對，浮點不可用 == 直接比。

Unified Memory 版本（最精簡）

float *A, *B, *C;
cudaMallocManaged(&A, vectorLength*sizeof(float));   // 統一記憶體，CPU/GPU 皆可存取
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); // 向上取整算 grid 大小
vecAdd<<<blocks, threads>>>(A, B, C, vectorLength);
cudaDeviceSynchronize();                             // 等 kernel 完成

serialVecAdd(A, B, comparisonResult, vectorLength); // CPU 對照
// ...vectorApproximatelyEqual(C, comparisonResult, ...) 驗證...

cudaFree(A); cudaFree(B); cudaFree(C);               // unified memory 用 cudaFree 釋放

• 用 cudaMallocManaged 配置統一記憶體，不需手動 cudaMemcpy：unified memory 會確保 A、B、C 對 GPU 可存取。
• 流程精簡為：配置 → host 初始化 → launch → sync → 驗證 → cudaFree。

Explicit Memory 版本（手動搬移）

// host 指標 A,B,C；device 指標 devA,devB,devC
cudaMallocHost(&A, n*sizeof(float)); // page-locked host 記憶體，做 copy 的最佳實踐
// ... B, C 同 ...
initArray(A, n); initArray(B, n);

cudaMalloc(&devA, n*sizeof(float));  // 配置 device 記憶體
cudaMalloc(&devB, n*sizeof(float));
cudaMalloc(&devC, n*sizeof(float));

cudaMemcpy(devA, A, n*sizeof(float), cudaMemcpyDefault); // H2D
cudaMemcpy(devB, B, n*sizeof(float), cudaMemcpyDefault);
cudaMemset(devC, 0, n*sizeof(float));

vecAdd<<<blocks, threads>>>(devA, devB, devC, n);
cudaDeviceSynchronize();

cudaMemcpy(C, devC, n*sizeof(float), cudaMemcpyDefault); // D2H 取回結果

cudaFree(devA); cudaFree(devB); cudaFree(devC);          // device 用 cudaFree
cudaFreeHost(A); cudaFreeHost(B); cudaFreeHost(C);       // page-locked 用 cudaFreeHost
free(comparisonResult);                                  // 一般 malloc 用 free

• host 緩衝區用 cudaMallocHost（page-locked / pinned memory），對要與 GPU 互傳的緩衝區是最佳實踐，可提升 copy 效能、且為非同步傳輸所必需。
• cudaMemcpyDefault 讓 CUDA 依來源/目的指標自動判斷拷貝方向。
• 對應釋放 API 要配對：cudaMalloc↔cudaFree、cudaMallocHost↔cudaFreeHost、malloc↔free。

   端到端流程 (Explicit Memory)
   ┌──────────┐   ┌───────────┐   ┌──────────┐   ┌──────┐   ┌──────────┐   ┌──────┐
   │ 配置記憶體 │──▶│ H2D copy  │──▶│ launch   │──▶│ sync │──▶│ D2H copy │──▶│ 驗證/ │
   │cudaMalloc│   │cudaMemcpy │   │ kernel   │   │等GPU │   │cudaMemcpy│   │ free │
   └──────────┘   └───────────┘   └──────────┘   └──────┘   └──────────┘   └──────┘
   (Unified Memory 省略 H2D / D2H 兩步，由系統按需搬移)

建置與執行

nvcc vecAdd_unifiedMemory.cu -o vecAdd_unifiedMemory
./vecAdd_unifiedMemory          # 預設長度 1024
./vecAdd_unifiedMemory 4096     # 命令列參數指定長度
# 輸出: Unified Memory: CPU and GPU answers match

Runtime Initialization (2.1.6)

CUDA runtime 會為系統中每個 device 建立一個 CUDA context，稱為該 device 的 primary context。此 context 在「第一個需要該 device 上有 active context 的 runtime 函式」被呼叫時初始化，並由應用程式中所有 host thread 共享。

• 初始化是隱式且透明的：作為 context 建立的一部分，device code 會在必要時JIT (just-in-time) 編譯並載入 device memory，整個過程對使用者透明。
• 若在 cudaInitDevice / cudaSetDevice 之前就發出 runtime API 請求，runtime 會隱式使用 device 0 並自我初始化以處理請求。
• primary context 可從 driver API 存取，以做 runtime 與 driver API 的互通。
• cudaDeviceReset 會銷毀當前 device 的 primary context；之後再呼叫 runtime API 會重新建立一個新的 primary context。

  第一個需要 context 的 runtime 呼叫
            │  (隱式觸發)
            ▼
   建立 primary context ──▶ device code JIT 編譯 ──▶ 載入 device memory
            │
            ▼   (此初始化會花費可觀的一次性開銷)
        runtime 就緒

考試/測驗重點

Related Notes

• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/02-CUDA-Cpp-Memory-Management
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 02-Programming-GPUs/05-CUDA-Python
• 02-Programming-GPUs/06-SIMT-Basics-and-Thread-Hierarchy
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 02-Programming-GPUs/17-NVCC-Compiler
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps