CUDA 程式設計入門：異構模型、記憶體管理、線程組織與 Kernel 啟動全流程詳解

前言

CUDA 讓 C/C++ 開發者以熟悉語法釋放 GPU 大規模並行算力。本文系統介紹異構運算模型、Host/Device 記憶體操作、線程層級組織、Kernel 啟動語法、函數類型限制，並附帶完整向量加法範例代碼。

異構運算模型

• Host：CPU 及主機記憶體（前綴 h_）。
• Device：GPU 及顯存（前綴 d_）。
• PCIe 總線連接，資料需顯式搬移。
• Kernel：標記 __global__ 的 Device 函數，從 Host 啟動，異步執行。

記憶體操作 API

cudaMemcpyKind：cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice。

錯誤檢查模式：

#define CHECK(call) do { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "%s:%d %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
        exit(1); \
    } \
} while(0)

線程組織：Grid → Block → Thread

• Grid：Kernel 啟動的所有線程集合，1D/2D/3D。
• Block：線程塊，共享 Shared Memory、可同步，1D/2D/3D。
• Thread：基本執行單元，擁有唯一 ID。

內建座標變數

threadIdx.x/y/z   // Block 內索引
blockIdx.x/y/z    // Grid 內 Block 索引
blockDim.x/y/z    // Block 尺寸
gridDim.x/y/z     // Grid 尺寸

全域線性索引計算（1D）

int idx = blockIdx.x * blockDim.x + threadIdx.x;

2D 矩陣映射

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * N + col;

Kernel 啟動語法

kernel_name<<>>(args...);
// grid, block 為 dim3 型別
dim3 block(256);
dim3 grid((N + block.x - 1) / block.x);

異步特性：啟動後控制權立即返回 Host，需 cudaDeviceSynchronize() 等待完成。

函數類型修飾符

Kernel 限制：僅存取 Device 記憶體、無可變參數、無靜態變數、無函數指標、異步啟動。

完整範例：向量加法

#include 
#include 
#include 
#include 
#include 

#define CHECK(call) ...  // 如上

void initialData(float *ip, int size) {
    srand((unsigned)time(NULL));
    for (int i = 0; i < size; ++i) ip[i] = (rand() & 0xFF) / 10.0f;
}

void sumArraysOnHost(float *A, float *B, float *C, int N) {
    for (int i = 0; i < N; ++i) C[i] = A[i] + B[i];
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

bool checkResult(float *hostRef, float *gpuRef, int N) {
    const double eps = 1e-8;
    for (int i = 0; i < N; ++i)
        if (fabs(hostRef[i] - gpuRef[i]) > eps) return false;
    return true;
}

int main(int argc, char **argv) {
    printf("%s Starting...\n", argv[0]);
    CHECK(cudaSetDevice(0));

    int nElem = 32;
    size_t nBytes = nElem * sizeof(float);

    float *h_A = (float*)malloc(nBytes);
    float *h_B = (float*)malloc(nBytes);
    float *hostRef = (float*)malloc(nBytes);
    float *gpuRef = (float*)malloc(nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);

    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((void**)&d_A, nBytes));
    CHECK(cudaMalloc((void**)&d_B, nBytes));
    CHECK(cudaMalloc((void**)&d_C, nBytes));

    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));

    dim3 block(nElem);
    dim3 grid(1);
    sumArraysOnGPU<<>>(d_A, d_B, d_C);
    printf("Execution config <<<%d, %d>>>\n", grid.x, block.x);

    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    if (checkResult(hostRef, gpuRef, nElem))
        printf("Arrays match.\n\n");
    else
        printf("Arrays mismatch!\n");

    CHECK(cudaFree(d_A)); CHECK(cudaFree(d_B)); CHECK(cudaFree(d_C));
    free(h_A); free(h_B); free(hostRef); free(gpuRef);
    CHECK(cudaDeviceReset());
    return 0;
}

編譯與運行

nvcc -o vecAdd vecAdd.cu
./vecAdd
# 輸出：
# vecAdd Starting...
# Vector size 32
# Execution config <<<1, 32>>>
# Arrays match.

進階學習路徑

1. 2D/3D 線程組織與矩陣運算（sumMatrixOnGPU2D）。
2. Device 查詢與多 GPU 選擇策略。
3. Shared Memory、Constant Memory、Texture Memory 進階用法。
4. 非同步流與事件重疊計算傳輸。
5. 效能分析：nvprof、Nsight Compute、Nsight Systems。

總結

掌握 Host/Device 記憶體邊界、線程索引映射、Kernel 異步啟動模型，是 CUDA 程式設計的「三大支柱」。建議從向量加法起步，逐步遷移至矩陣乘法、卷積、減運算等實際算法，並養成「先在 Host 寫參考實現 → 再移植 Device → 對數驗證」的工程習慣。