Enter your search text in the box above

Select a result to preview

多維 Grid 組織 (Multidimensional Grid Organization)

#cuda-programming #multidimensional-grids #dim3 #thread-hierarchy #thread-indexing

重點總覽 (Overview)

項目 內容 重點
階層 (hierarchy) grid → block → thread，兩層 grid 是 block 的 3D 陣列；block 是 thread 的 3D 陣列
執行設定 (execution config) <<<dimGrid, dimBlock>>> 第 1 參數 = grid 維度 (以 block 計)；第 2 參數 = block 維度 (以 thread 計)
dim3 型別 三元素整數向量 (x, y, z) 未用維度設為 1 即可降維
座標變數 (built-in) blockIdx, threadIdx thread 自身座標，用來定位資料
維度變數 (built-in) gridDim, blockDim grid/block 的尺寸，由執行設定預先初始化
block 上限 總 thread 數 ≤ 1024 x*y*z ≤ 1024，分配方式任意
gridDim 上限 x: 1..2³¹-1；y,z: 1..65535 各維度範圍不同
Important

grid 中所有 thread 執行同一個 kernel function，靠 blockIdx/threadIdx 座標彼此區分並認領各自負責的資料區段。grid 一旦啟動，grid/block 維度在整個 grid 執行完畢前保持不變。

兩層執行緒階層 (Two-Level Thread Hierarchy)

內建變數 意義 由誰決定 在 kernel 內可否改名
gridDim grid 各維 block 數 第 1 個執行設定參數 否 (CUDA C 規範)
blockDim block 各維 thread 數 第 2 個執行設定參數 否
blockIdx 本 block 在 grid 的座標 排程時指定 否
threadIdx 本 thread 在 block 的座標 排程時指定 否
Tip

記法：*Dim 是「尺寸」(總共幾個)，*Idx 是「索引」(我是第幾個)。Dim 來自 host 的執行設定；Idx 由硬體在排程時賦予。

執行設定參數與 dim3 (Execution Configuration & dim3)

每個執行設定參數型別為 dim3：三個元素 x, y, z 的整數向量。用不到的維度設為 1 即可。

// 1D grid：32 個 block，每個 block 128 threads -> 共 128*32 = 4096 threads
dim3 dimGrid(32, 1, 1);
dim3 dimBlock(128, 1, 1);
vecAddKernel<<<dimGrid, dimBlock>>>(...);

// host 變數名稱任意 (只要型別是 dim3)，下列效果完全相同
dim3 dog(32, 1, 1);
dim3 cat(128, 1, 1);
vecAddKernel<<<dog, cat>>>(...);

維度可由其他變數在 runtime 計算，讓 block 數隨資料量變化：

// block 固定 256；grid 大小隨 n 變動
dim3 dimGrid(ceil(n/256.0), 1, 1);
dim3 dimBlock(256, 1, 1);
vecAddKernel<<<dimGrid, dimBlock>>>(...);
// n=1000 -> 4 blocks；n=4000 -> 16 blocks (皆足以覆蓋所有元素)

1D 捷徑 (shorthand)：直接傳算術運算式，編譯器把它當 x 維、y/z 自動為 1：

vecAddKernel<<<ceil(n/256.0), 256>>>(...);  // 等同上面的 dim3 寫法
Tip

此捷徑利用 C++ 建構子的預設參數：dim3 的三個參數預設值皆為 1，傳單一值時只填第一個參數 (x)，y、z 取預設 1。

Warning

kernel 內讀到的是維度值，例如 n=4000 時 gridDim.x = 16、blockDim.x = 256。這些名稱不可更改 (與 host 端可自取的 dimGrid/dog 不同)。

各維度尺寸限制 (Per-Dimension Size Limits)

變數 / 限制 範圍
gridDim.x 1 .. 2³¹ − 1
gridDim.y 1 .. 2¹⁶ − 1 (65,535)
gridDim.z 1 .. 2¹⁶ − 1 (65,535)
blockIdx.x 0 .. gridDim.x − 1
blockIdx.y 0 .. gridDim.y − 1
blockIdx.z 0 .. gridDim.z − 1
block 總 thread 數 blockDim.x * blockDim.y * blockDim.z ≤ 1024

block 形狀範例 (上限 1024)：

(x, y, z) 總數 合法?
(512, 1, 1) 512 ✅
(8, 16, 4) 512 ✅
(32, 16, 2) 1024 ✅
(32, 32, 2) 2048 ❌ 超過 1024
Warning

邊界情況：compute capability < 3.0 的舊裝置，gridDim.x (即 x 維 block 數) 上限只到 2¹⁶ − 1，而非 2³¹ − 1。寫可攜程式時別假設第一維可無限大。

維度標記順序 (Dimension Labeling Order)

圖 3.1 玩具範例：gridDim = (2, 2, 1)、blockDim = (4, 2, 2)。

dim3 dimGrid(2, 2, 1);    // C 寫法：低維 (x) 在前
dim3 dimBlock(4, 2, 2);
KernelFunction<<<dimGrid, dimBlock>>>(...);

Grid  gridDim=(x=2, y=2, z=1)        block 標記 = (blockIdx.y, blockIdx.x)
        bIdx.x=0     bIdx.x=1
      +-----------+-----------+
y=0   | block(0,0)| block(0,1)|
      +-----------+-----------+
y=1   | block(1,0)| block(1,1)|   block(1,0): blockIdx.y=1, blockIdx.x=0
      +-----------+-----------+

每個 block  blockDim=(x=4, y=2, z=2) -> 4*2*2 = 16 threads
   thread 標記 = (threadIdx.z, threadIdx.y, threadIdx.x)
   例: thread(1,0,2)  ->  z=1, y=0, x=2
   全 grid 共 4 blocks * 16 = 64 threads
Important

標記順序與 C 設定順序相反：圖示中的 block/thread 標記最高維在前 (z, y, x)；但在 dim3(...) 與執行設定裡是最低維在前 (x, y, z)。此「反序」是為了之後把 thread 座標映射到多維資料索引時更直覺 (見 03-Multidimensional-Grids-And-Data/02-Mapping-Threads-to-Multidimensional-Data)。

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

情境 / 關鍵字 答案 / 技巧
<<<A, B>>> 兩參數意義 A = grid 維度 (block 數)，B = block 維度 (thread 數)
算 grid 中總 thread 數 (1D) gridDim.x * blockDim.x；2D/3D 再乘上各維
dimGrid 名稱能改嗎? host 端可 (型別需 dim3)；kernel 內 gridDim/blockDim 不可
block 最多幾 threads? 1024，且 x*y*z ≤ 1024
(32, 32, 2) 合法嗎? 否，2048 > 1024
gridDim.y / .z 上限 65,535 (2¹⁶−1)；gridDim.x 才是 2³¹−1
blockIdx.x 範圍 0 .. gridDim.x − 1 (索引從 0 起)
<<<ceil(n/256.0), 256>>> 意義 算術式當 x 維、y/z=1；block 固定 256，grid 隨 n 變
grid 與 block 必須同維? 否，可不同維度
圖中 block(1,0) 代表 blockIdx.y=1, blockIdx.x=0 (高維在前，與 C 設定反序)
啟動後維度可變嗎? 否，整個 grid 執行完前維度固定