Enter your search text in the box above

Select a result to preview

Vector Addition 主機端程式碼：Device Global Memory 與資料傳輸 (Host Code, Device Memory, and Data Transfer)

#cuda-programming #device-global-memory #vector-addition #heterogeneous-computing

重點總覽 (Overview)

本筆記聚焦 vector addition 的 host 端程式碼：先回顧序列版 vecAdd，再說明把工作「外包 (outsource)」到 device 所需的三件事——配置 device global memory、雙向搬資料、釋放記憶體。kernel 本身留待 02-Heterogeneous-Data-Parallel-Computing/03-Kernel-Functions-and-Threading。

項目 重點 備註
Sequential vecAdd for 迴圈逐一 C[i]=A[i]+B[i] 純 host code，全部變數用 _h 後綴
Stub / 外包模型 host 函式分 Part 1/2/3：配置+搬入 → 呼叫 kernel → 搬回+釋放 main 完全不需知道計算在 device 上跑
Device global memory device 卡上的 DRAM（例：V100 16/32GB） host 指標不可 dereference
cudaMalloc((void**)&p, size) 配置 device 記憶體 第一參數是「指標的位址」，size 以 bytes 計
cudaFree(p) 釋放 device 記憶體 只需指標值，不需位址
cudaMemcpy(dst, src, size, kind) host↔device 搬資料 dst 在前、src 在後（與直覺相反）
Error checking API 回傳 cudaError_t，比對 cudaSuccess 多數錯誤源自不當的參數值
效能警語 此例搬資料成本 > 計算量，GPU 版可能比 CPU 慢 arithmetic intensity 極低

序列版 vecAdd 與指標語意 (Sequential vecAdd & Pointer Semantics)

序列版（Fig. 2.4）是平行化的對照基準——data parallelism 的「Hello World」：

// Compute vector sum C_h = A_h + B_h
void vecAdd(float* A_h, float* B_h, float* C_h, int n) {
    for (int i = 0; i < n; ++i) {
        C_h[i] = A_h[i] + B_h[i];   // 每個 i 互相獨立 → 天生可平行
    }
}
int main() {
    // 配置 + I/O 讀入 A, B（各 N 個元素），略
    vecAdd(A, B, C, N);
}

Stub 與外包模型 (Stub & Outsourcing Model)

要平行化，最直接的做法是改寫 vecAdd，把計算搬到 device。改寫後的 host 函式扮演外包代理人 (outsourcing agent)，分成三部分：

        HOST (CPU)                                  DEVICE (GPU)
   ┌───────────────────┐                      ┌───────────────────────┐
   │  A_h, B_h, C_h     │                      │  device global memory │
   │  (host memory)     │                      │                       │
   └───────────────────┘                      └───────────────────────┘
            │                                            ▲
   Part 1   │  cudaMalloc(A_d,B_d,C_d)  ───────────────► │ 配置 A_d,B_d,C_d
            │  cudaMemcpy(H→D) A_h→A_d, B_h→B_d ───────► │ 搬入輸入
            │                                            │
   Part 2   │  vecAddKernel<<<...>>>(A_d,B_d,C_d,n) ───► │ 啟動 grid 做加法
            │                                            │
   Part 3   │  cudaMemcpy(D→H) C_d→C_h  ◄─────────────── │ 搬回結果
            │  cudaFree(A_d,B_d,C_d)    ───────────────► │ 釋放
            ▼
   結果在 C_h 可用
Important

此「transparent outsourcing」模型方便教學，但在實務上很沒效率——資料來回搬動成本高。真實程式會把大型資料留在 device 跨多個 kernel 呼叫重複使用，攤平搬運 overhead。

Device Global Memory 與配置/釋放 (cudaMalloc / cudaFree)

Device global memory（簡稱 global memory）：device 硬體卡上自帶的 DRAM（例：NVIDIA Volta V100 為 16GB 或 32GB）。叫「global」是為了與其他 device 記憶體類型區分（詳見 05-Memory-Architecture-And-Data-Locality/01-Memory-Access-Efficiency-and-CUDA-Memory-Types）。

float *A_d;
int size = n * sizeof(float);        // ← size 以「bytes」計算！
cudaMalloc((void**) &A_d, size);     // 配置；&A_d 轉成 generic pointer
...
cudaFree(A_d);                       // 釋放；只傳「值」即可
cudaMalloc C malloc
回傳 cudaError_t（錯誤碼） 指向配置物件的指標
參數數量 2：(指標變數位址, bytes) 1：(bytes)
寫入位址的方式 寫進你傳入位址的指標變數 由 return value 帶回

為何長這樣？

Warning

A_d、B_d、C_d 指向的是 device global memory，不可在 host code 裡 dereference（例如 A_d[0]）。這麼做會造成 exception 或 runtime error。它們只能用於 API 呼叫與 kernel 呼叫的參數。

Tip

記得換算單位：要配置 n 個 single-precision float，size = n * sizeof(float) = n * 4 bytes。忘了乘 sizeof 是最常見的 bug。

資料傳輸 cudaMemcpy (Data Transfer)

cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice);  // 搬入 A
cudaMemcpy(B_d, B_h, size, cudaMemcpyHostToDevice);  // 搬入 B
...
cudaMemcpy(C_h, C_d, size, cudaMemcpyDeviceToHost);  // 搬回 C

四個參數（順序很重要）：

位置 參數 說明
1 dst 目的地指標（先寫目的地！）
2 src 來源指標
3 size 要複製的 bytes 數
4 kind 傳輸方向（predefined constant）

方向常數 (transfer kind)：

常數 方向
cudaMemcpyHostToHost host → host
cudaMemcpyHostToDevice host → device（搬入輸入）
cudaMemcpyDeviceToHost device → host（搬回結果）
cudaMemcpyDeviceToDevice device global memory 內部互搬
Warning

cudaMemcpy 是 (dst, src, ...)，與英文「copy A to B」的直覺相反。寫成 (src, dst, ...) 會把資料搬反方向，是經典考點與 bug。同一個函式靠調換 dst/src 與選對 kind 即可雙向傳輸。

完整 host code 與錯誤檢查 (Complete Host Code & Error Checking)

把 Part 1/3 補齊後的完整 vecAdd（Fig. 2.8 / 2.13），Part 2 的 kernel 呼叫留待後續章節：

void vecAdd(float* A_h, float* B_h, float* C_h, int n) {
    int size = n * sizeof(float);
    float *A_d, *B_d, *C_d;

    // Part 1: 配置 device memory + 搬入
    cudaMalloc((void**) &A_d, size);
    cudaMalloc((void**) &B_d, size);
    cudaMalloc((void**) &C_d, size);
    cudaMemcpy(A_d, A_h, size, cudaMemcpyHostToDevice);
    cudaMemcpy(B_d, B_h, size, cudaMemcpyHostToDevice);

    // Part 2: 呼叫 kernel（execution configuration，後述）
    // vecAddKernel<<<ceil(n/256.0), 256>>>(A_d, B_d, C_d, n);

    // Part 3: 搬回結果 + 釋放
    cudaMemcpy(C_h, C_d, size, cudaMemcpyDeviceToHost);
    cudaFree(A_d);
    cudaFree(B_d);
    cudaFree(C_d);
}

Error checking：每個 CUDA API 都回傳 cudaError_t flag，務必檢查（多數錯誤來自不當的參數值，例如記憶體不足）：

cudaError_t err = cudaMalloc((void**) &A_d, size);
if (err != cudaSuccess) {
    printf("%s in %s at line %d\n",
           cudaGetErrorString(err), __FILE__, __LINE__);
    exit(EXIT_FAILURE);
}
Tip

書中範例為求精簡通常省略 error check，但實務上應以一個 C macro 包住每個 API 呼叫，能省下大量除錯時間。

效能警語：為何此例 GPU 可能比 CPU 慢

Warning

vector addition 只是「簡單」範例。配置 device memory、H→D 搬入、D→H 搬回、釋放等 overhead，往往使 GPU 版比序列版還慢。

原因是 arithmetic intensity 極低：對 2 個 float 輸入 + 1 個 float 輸出（共 12 bytes 流量）只做 1 次加法。

Arithmetic Intensity=1 FLOP12 bytes≈0.083 FLOP/byte (memory-bound)

真實應用的 kernel 計算量相對資料量大得多，且會把資料留在 device 跨多次 kernel 呼叫攤平 overhead，才划得來。延伸見 06-Performance-Considerations/04-Optimization-Checklist-and-Bottlenecks。

考試/面試重點 (Exam / Test Patterns)

情境 / 關鍵字 答案 / 技巧
cudaMalloc 配置 v 個 int 的 第二參數 v * sizeof(int)（bytes，不是元素數）
cudaMalloc 讓 A_d 指向配置的 第一參數 (void**) &A_d（指標的位址＋轉型）
為何 cudaMalloc 有 2 個參數（C malloc 只有 1） return value 空出來回報錯誤；位址寫進指標變數
cudaFree 傳什麼 傳指標值 cudaFree(A_d)，不傳位址
把 3000 bytes 從 A_h 複製到 A_d cudaMemcpy(A_d, A_h, 3000, cudaMemcpyHostToDevice)（dst 先）
cudaMemcpy 參數順序 (dst, src, size, kind)——目的地在前
宣告接收 API 回傳值的變數 cudaError_t err;
搬回結果用哪個 kind cudaMemcpyDeviceToHost
host code 可否 A_d[i] = ... 不可，dereference device 指標→runtime error
為何 GPU vecAdd 可能比 CPU 慢 搬資料 overhead > 計算量；arithmetic intensity 太低
stub 是什麼 只負責配置/搬資料/呼叫 kernel 的 host 包裝函式