Kernel Functions 與 Threading

重點總覽 (Overview)

本節說明「如何撰寫 kernel 本體」：用限定詞標記函式、理解 grid/block/thread 階層、用內建變數算出每個 thread 的全域索引，再以 if(i<n) 防止越界。啟動語法 <<<>>> 與編譯流程屬於 02-Heterogeneous-Data-Parallel-Computing/04-Calling-Kernels-Compilation-and-Summary。

Kernel Function 與 SPMD 模型 (Kernel Functions and the SPMD Model)

• Kernel function 定義「在平行階段中，每個 thread 要執行的程式碼」。
• 由於所有 thread 跑同一份程式碼,CUDA 屬於 SPMD 平行風格。

向量加法 kernel (書中 Fig. 2.10) —— 與序列版 (Fig. 2.4) 相比，for-loop 不見了：

// Compute vector sum C = A + B
// Each thread performs one pair-wise addition
__global__
void vecAddKernel(float* A, float* B, float* C, int n) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

三種函式限定詞 (Function Declaration Qualifiers)

• __global__ 前後各兩個底線,回傳型別必須是 void。
• 無任何限定詞 → 預設為 __host__。CUDA 程式多由 CPU-only 程式移植而來,讓原函式自動成為 host function 可省去大量改寫。

__device__               ─ device 執行,device 呼叫,不開新 thread
__global__  void K(...)   ─ device 執行,host 呼叫,launch 一個 grid
__host__                 ─ host 執行,host 呼叫 (default)

Thread 階層:Grid → Block → Thread (Thread Hierarchy)

呼叫 kernel 時,runtime 啟動一個組織成兩層階層的 grid:

            ┌───────────────────────── Grid (一次 kernel 呼叫的全部 thread) ─────────────────────────┐
            │                                                                                        │
        Block 0                         Block 1                         Block 2          ...    Block (gridDim-1)
   ┌──────────────────┐           ┌──────────────────┐           ┌──────────────────┐
   │ t0 t1 ... t255   │           │ t0 t1 ... t255   │           │ t0 t1 ... t255   │   每個 block 同樣大小
   └──────────────────┘           └──────────────────┘           └──────────────────┘   (本例 blockDim.x = 256)
   blockIdx.x = 0                  blockIdx.x = 1                  blockIdx.x = 2

• Grid = 一次 kernel 呼叫啟動的所有 thread。
• Grid = 一個由 block (thread block) 組成的陣列;同一 grid 內所有 block 大小相同。
• 每個 block 在目前硬體上最多 1024 threads (CUDA 3.0+;舊版僅 512)。

內建變數與全域索引 (Built-in Variables and the Global Index)

CUDA 提供一組唯讀內建變數,由 runtime 預先初始化,讓 thread 找到自己:

每個 thread 用以下公式算出整個 grid 內唯一的全域索引:

i = blockIdx.x * blockDim.x + threadIdx.x

以 blockDim.x = 256 為例,前三個 block 連續覆蓋 i = 0 .. 767:

 blockDim.x = 256

 Block 0 (blockIdx.x=0)      Block 1 (blockIdx.x=1)      Block 2 (blockIdx.x=2)
 threadIdx.x: 0 ... 255      threadIdx.x: 0 ... 255      threadIdx.x: 0 ... 255
        │                           │                           │
   i = 0*256+tIdx              i = 1*256+tIdx              i = 2*256+tIdx
        ▼                           ▼                           ▼
   i =   0 ... 255             i = 256 ... 511             i = 512 ... 767
   ───────────────────────────────────────────────────────────────────────►  連續覆蓋 0..767

• 每個 thread 用 i 去存取 A[i], B[i], C[i] —— 即 thread → data 的一對一映射。
• 第 b 個 block 覆蓋的索引區間 = [b*blockDim.x, b*blockDim.x + blockDim.x - 1]。
• 啟動 ≥ n 個 thread,即可處理長度 n 的向量。多維索引推廣見 03-Multidimensional-Grids-And-Data/02-Mapping-Threads-to-Multidimensional-Data。

邊界檢查 if(i < n) (Boundary Check)

並非所有向量長度都能被 block size 整除,因此 grid 通常會多啟動一些 thread:

• 例:n = 100,block size 取 32 → 需 ceil(100/32) = 4 個 block = 128 threads。
• 多出的 128 - 100 = 28 個 thread 必須被關閉,否則會存取超出陣列的記憶體。

 n = 100,   block size = 32,   launch 4 blocks = 128 threads

  thread i:  0 ............................... 99 | 100 ............ 127
             └──── i < n  → 執行 C[i]=A[i]+B[i] ──┘ └─ i >= n → 跳過 ─┘
                          (前 100 個)                    (後 28 個)

• 所有 thread 都跑同一份程式碼,因此都會測試 if (i < n):
  • i < n 的 thread → 執行加法;
  • i >= n 的 thread → 什麼都不做。
• 這讓同一個 kernel 能處理任意長度的向量。

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

Related Notes

• 02-Heterogeneous-Data-Parallel-Computing/01-Data-Parallelism-and-CUDA-Program-Structure
• 02-Heterogeneous-Data-Parallel-Computing/02-Vector-Addition-Host-Code-Device-Memory-and-Data-Transfer
• 02-Heterogeneous-Data-Parallel-Computing/04-Calling-Kernels-Compilation-and-Summary
• 03-Multidimensional-Grids-And-Data/02-Mapping-Threads-to-Multidimensional-Data
• 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
• 05-Memory-Architecture-And-Data-Locality/03-Boundary-Checks-and-Memory-Occupancy