進階啟動與 Clusters (Advanced Launch and Clusters)

重點總覽

進階 CUDA API 與功能的定位

本章涵蓋較進階的 CUDA API 與功能。這些技巧多半不需要修改 kernel 程式碼,而是從 host 端影響應用層行為,包含 GPU 工作的執行方式、效能,以及 CPU 端的效能。

• 通常以 host 端 API 呼叫或啟動屬性來達成。
• 既影響 GPU work execution,也影響 CPU/GPU 兩側效能表現。

3.1.1 cudaLaunchKernelEx

最初版本引入 triple chevron <<<>>> 記法時,kernel 的執行配置只有四個可程式化參數:

• thread block 維度
• grid 維度
• dynamic shared memory(選填,未指定為 0)
• stream(未指定則用 default stream)

部分 CUDA 功能可受惠於啟動時提供的額外屬性與提示。cudaLaunchKernelEx 讓程式能以 cudaLaunchConfig_t 結構設定上述執行配置參數,並可再傳入零或多個 cudaLaunchAttribute 來控制或建議其他啟動參數。

cudaLaunchConfig_t config = {0};
config.gridDim   = numBlocks;        // 等同 <<<>>> 第一參數
config.blockDim  = threadsPerBlock;  // 等同 <<<>>> 第二參數
config.attrs     = attribute;        // cudaLaunchAttribute 陣列
config.numAttrs  = 1;                // 屬性數量
cudaLaunchKernelEx(&config, kernel, args...);

cudaLaunchConfig_t 內含 gridDim、blockDim、dynamicSmemBytes、stream,再加上 attrs(屬性陣列指標)與 numAttrs(陣列長度)。

• 屬性陣列每個元素是 cudaLaunchAttribute,以 id 標示屬性種類,val 為對應的值(union)。
• 範例屬性:cudaLaunchAttributePreferredSharedMemoryCarveout(L1/Shared Memory 平衡,見 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config)、cudaLaunchAttributeClusterDimension(指定 cluster 大小)。
• 完整支援屬性清單見 CUDA Runtime API Reference Documentation。

triple chevron  <<< grid, block, [smem], [stream] >>>
                 └─ 只有 4 個固定參數

cudaLaunchKernelEx(&config, kernel, args)
   config: gridDim / blockDim / dynamicSmemBytes / stream
         + attrs[0..numAttrs-1]   ← 可擴充的屬性陣列
              ├ cudaLaunchAttributeClusterDimension
              ├ cudaLaunchAttributePreferredSharedMemoryCarveout
              └ ...（更多屬性）

3.1.2 Launching Clusters

Thread block clusters 是 compute capability 9.0 以上才有的可選 thread block 組織層級。它保證一個 cluster 內的所有 thread blocks 同時在單一 GPC(GPU Processing Cluster)上執行,使得超過單一 SM 容量的更大一群 threads 能彼此交換資料與同步。

Grid
 └── Cluster (同時駐留在單一 GPC)
       ├── Thread Block 0   ┐
       ├── Thread Block 1   ├─ 可跨 block 交換資料 / 同步
       └── Thread Block ... ┘
             └── Threads

• cluster 讓多個 SM 上的 blocks能協作,突破單一 SM/單一 block 的規模限制。
• 用 <<<>>> 啟動 cluster 時,cluster 大小是隱式決定的(由 __cluster_dims__ 指定,見 §3.1.2.2)。
• 與 cudaLaunchKernelEx 搭配時,則可**逐次啟動(per-launch)**地配置 cluster 大小。

3.1.2.1 用 cudaLaunchKernelEx 啟動 Clusters

與 <<<>>> 不同,使用 cudaLaunchKernelEx 時 cluster 大小可逐次啟動配置。透過 cudaLaunchAttributeClusterDimension 屬性指定 cluster 維度。

// Kernel 定義：無編譯期屬性附加在 kernel 上
__global__ void cluster_kernel(float *input, float* output) {}

int main() {
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    cudaLaunchConfig_t config = {0};
    // grid 維度不受 cluster 影響，仍以 block 數枚舉；且須為 cluster size 的倍數
    config.gridDim  = numBlocks;
    config.blockDim = threadsPerBlock;

    cudaLaunchAttribute attribute[1];
    attribute[0].id = cudaLaunchAttributeClusterDimension;
    attribute[0].val.clusterDim.x = 2;  // X 維度的 cluster 大小
    attribute[0].val.clusterDim.y = 1;
    attribute[0].val.clusterDim.z = 1;
    config.attrs    = attribute;
    config.numAttrs = 1;

    cudaLaunchKernelEx(&config, cluster_kernel, input, output);
}

此處 kernel 無任何編譯期 cluster 屬性,cluster 維度完全由 runtime 的 attribute 決定。

與 thread block clusters 相關的兩種 attribute:

• cudaLaunchAttributeClusterDimension 的效果類似在 kernel 定義上加 __cluster_dims__,差別在於可在 runtime 對同一 kernel 的不同啟動改變。
• 使用 preferred 維度時,必須同時指定最小 cluster 維度(由 __cluster_dims__ 或 cudaLaunchAttributeClusterDimension 提供)。
• 所有 thread blocks 都會以「至少最小 cluster 維度」執行;在可能時採用 preferred 維度,但不保證所有 cluster 都以 preferred 維度執行。

3.1.2.2 Blocks as Clusters

當 kernel 以 __cluster_dims__ 標註時,grid 中的 cluster 數量是隱式的,由 grid 大小除以指定 cluster 大小算出。

__cluster_dims__((2, 2, 2)) __global__ void foo();
// 每個 cluster 含 2x2x2 個 thread blocks → 8x8x8 個 clusters
foo<<<dim3(16, 16, 16), dim3(1024, 1, 1)>>>();

此例 grid 為 16x16x16 個 thread blocks,因每 cluster 是 2x2x2 blocks,故得 8x8x8 個 clusters。<<<>>> 第一參數仍是 block 數。

另一種方式是用 __block_size__ 標註,在 kernel 定義時同時指定 block size 與 cluster size。啟用後,<<<>>> 的 grid 維度改以 cluster 為單位而非 thread block。

// 每 block 多少 threads、每 cluster 多少 blocks，皆作為 kernel 的屬性處理
__block_size__((1024, 1, 1), (2, 2, 2)) __global__ void foo();
// 8x8x8 clusters
foo<<<dim3(8, 8, 8)>>>();

__block_size__ 需兩個欄位,各為 3 元素 tuple:第一個是 block 維度,第二個是 cluster 大小;第二個 tuple 若未傳入則預設為 (1,1,1)。

__block_size__((1024,1,1),(2,2,2))  →  啟用 Blocks as Clusters
   foo<<< dim3(8,8,8) >>>();              // 第一參數 = cluster 數
   foo<<< dim3(8,8,8), 1, stream >>>();   // 帶 smem/stream 時，第二參數必須是 1
                              ▲
                              └─ 佔位符；其他值 = UB

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/02-Advanced-Streams-and-Dependent-Launch
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 03-Advanced-CUDA/05-Thread-Scopes-and-Scoped-Atomics
• 03-Advanced-CUDA/06-Asynchronous-Barriers-and-Pipelines
• 03-Advanced-CUDA/07-Async-Data-Copies-and-L1-Config
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/04-CUDA-Cpp-Errors-and-Specifiers
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps