SIMT 基礎與 Thread 階層 (SIMT Basics and Thread Hierarchy)

重點總覽

SIMT 基礎 (Basics of SIMT)

CUDA kernel 大致可以用寫傳統 CPU 程式的方式來撰寫，但理解 GPU 上 thread 如何被排程、如何存取記憶體、如何推進執行，能幫你寫出最大化資源利用率的 kernel。

• 從開發者角度，CUDA thread 是平行的基本單位 (fundamental unit of parallelism)。
• SIMT 模型讓每個 thread 維護自己的 state 與 control flow：功能上每個 thread 可以執行完全不同的 code path。
• 但若同一 warp 內的 thread 走 divergent code path（分支發散），會造成效能損失；好的 kernel 應盡量減少 warp 內 divergence。

Thread 階層 (Thread Hierarchy)

Thread 被組織成 thread block，block 再組織成 grid。grid 與 block 都可以是 1、2 或 3 維。kernel 內用以下 built-in 變數查詢自己的位置與維度。

這些變數的主要用途，是組出每個 thread 唯一的 global thread index，讓每個 thread 去讀寫 global memory 的特定資料、執行需要的 code path：

// 1D grid + 1D block 的標準慣用法
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < vectorLength) {
    C[idx] = A[idx] + B[idx];
}

# Python 提供 cuda.grid(1) 直接回傳一維全域索引
work_index = cuda.grid(1)
C[work_index] = A[work_index] + B[work_index]

C++ 需手動 blockIdx.x * blockDim.x + threadIdx.x；Python 的 cuda.grid(1) 是同一公式的便捷封裝。加上 if (idx < length) 的邊界檢查，是因為 grid 通常會比資料量略大。

多維與 linearization

多維 block/grid 只是為了撰寫方便，不影響效能。block 內的 thread 以可預測的方式被線性化：

• threadIdx.x 變化最快（stride 1）
• threadIdx.y 的 stride = blockDim.x
• threadIdx.z 的 stride = blockDim.x * blockDim.y

linear = threadIdx.x
       + threadIdx.y * blockDim.x
       + threadIdx.z * blockDim.x * blockDim.y

2D grid (gridDim = 3 x 2) 含 1D thread block (blockDim.x = 4)
─────────────────────────────────────────────────────────────
 blockIdx (y)
   1 │ [B(0,1)] [B(1,1)] [B(2,1)]
   0 │ [B(0,0)] [B(1,0)] [B(2,0)]
     └──────────────────────────── blockIdx (x)
          0        1        2

每個 block 內 (blockDim.x = 4)，threadIdx.x 連續、x 最快：
   B(1,0):  t0   t1   t2   t3
            ↑ threadIdx.x = 0,1,2,3  （連續 thread → 連續 warp lane）

Thread Block 同步 (__syncthreads())

在前面的 vecAdd 等範例中，每個 thread 做獨立工作，不需要互相同步。一旦 thread 之間要合作、或存取同一塊記憶體位址（尤其是 shared memory），就必須同步，以避免 race condition 與 memory hazard。

block 層級最基本的同步機制是 __syncthreads()（Python：cuda.syncthreads()），它是一個 barrier：

• block 內所有 thread 都到達該呼叫之前，沒有任何 thread 被允許往下執行。
• barrier 也建立記憶體順序保證：barrier 之前對 shared memory 的所有寫入，會「排序」在任何 thread 通過 barrier 之後的讀取之前（先寫後讀的可見性）。
• __syncthreads() 被設計成 lightweight（搭配近核心、低延遲的 shared memory）。

// 假設 blockDim.x == 128
__global__ void example_syncthreads(int* input_data, int* output_data) {
    __shared__ int shared_data[128];
    shared_data[threadIdx.x] = input_data[blockDim.x * blockIdx.x + threadIdx.x];

    // barrier：保證所有對 shared_data 的寫入都完成且有序，
    // 之後才有任何 thread 被解除阻塞：
    __syncthreads();

    // 只讓單一 thread 安全地讀取整塊 shared_data：
    if (threadIdx.x == 0) {
        float sum = 0;
        for (int i = 0; i < blockDim.x; ++i)
            sum += shared_data[i];
        output_data[blockIdx.x] = sum;
    }
}

典型樣式：「各 thread 寫 shared memory → __syncthreads() → 再讀其他 thread 寫的資料」。少了中間的 barrier，thread 0 可能在其他 warp 還沒寫完 shared_data 時就開始加總，導致讀到未完成的值。

時間軸（block 內）：
  warp0  ──寫 shared_data──┐
  warp1  ──寫 shared_data──┤  (barrier) ===> 全部到齊才放行 ==> 讀取/加總
  warp2  ──寫 shared_data──┤        ▲
  warp3  ──寫 shared_data──┘   沒有 barrier 時，快的 warp 會提早讀到舊值

考試/測驗重點

Related Notes

• 02-Programming-GPUs/05-CUDA-Python
• 02-Programming-GPUs/07-SIMT-Device-Memory-Spaces
• 02-Programming-GPUs/08-SIMT-Memory-Performance
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 01-Introduction-to-CUDA/02-Execution-Model-and-SIMT
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps