CUDA Grid 和 Block

這篇備忘錄記錄了 CUDA Grid 和 Block。

資訊

在 CUDA 中，並行執行的程式碼稱為核函數 (Kernel)。當你啟動一個核函數時，你必須指定執行該核函數的線程數量以及它們如何組織。這就是 Grid (網格) 和 Block (塊) 的概念。

1. 線程層次結構

CUDA 的線程模型是一個分層結構：

• Grid (網格)：由一個或多個 Block 組成。Grid 可以是一維、二維或三維的。
• Block (塊)：由多個線程組成。Block 也可以是一維、二維或三維的。
• Thread (線程)：Grid 中執行的最小單位。每個線程都執行相同的核函數程式碼。

這種層次結構允許你將問題分解為獨立的子任務，並將這些子任務映射到 Grid 和 Block 的維度上。

2. 定義 Grid 和 Block 維度

在啟動核函數時，你使用 <<<Dg, Db>>> 語法來指定 Grid 和 Block 的維度：

kernel_function<<<Dg, Db>>>(arguments);

• Dg (Dimension Grid)：指定 Grid 的維度。這是一個 dim3 類型的變量，可以是一維 (Dg.x)、二維 (Dg.x, Dg.y) 或三維 (Dg.x, Dg.y, Dg.z)。
• Db (Dimension Block)：指定每個 Block 的維度。這也是一個 dim3 類型的變量，可以是一維 (Db.x)、二維 (Db.y) 或三維 (Db.z)。

dim3 結構體

dim3 是一個結構體，通常用於表示 Grid 或 Block 的大小。

struct dim3
{
    unsigned int x, y, z;
};

如果你只指定一個值，它會被賦給 .x，其他為 1。例如 dim3(10) 等同於 dim3(10, 1, 1)。

3. 線程索引

在核函數內部，每個線程都可以通過幾個內置變量來識別自己的唯一索引：

• blockIdx.x, blockIdx.y, blockIdx.z：當前線程所在的 Block 在 Grid 中的索引。
• blockDim.x, blockDim.y, blockDim.z：每個 Block 的維度（線程數）。
• threadIdx.x, threadIdx.y, threadIdx.z：當前線程在 Block 中的索引。
• gridDim.x, gridDim.y, gridDim.z：Grid 的維度（Block 數）。

通過這些變量，可以計算出每個線程在整個 Grid 中的全局唯一索引。

一維索引計算

對於一維 Grid 和 Block，全局索引的計算方式為：

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

二維索引計算

對於二維 Grid 和 Block，全局索引的計算方式為：

int x_index = blockIdx.x * blockDim.x + threadIdx.x;
int y_index = blockIdx.y * blockDim.y + threadIdx.y;

// 如果將二維數據存儲在一維數組中，則需要將 (x_index, y_index) 轉換為一維索引
// 假設矩陣寬度為 width
int global_index = y_index * width + x_index;

範例：向量加法

這是一個簡單的向量加法核函數，演示了 Grid 和 Block 的使用。

#include <iostream>
#include <vector>
#include <cuda_runtime.h>

// 核函數：在設備上執行向量加法
__global__ void add_vectors(int* a, int* b, int* c, int N) {
    // 計算當前線程在 Grid 中的全局索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 確保索引在有效範圍內
    if (idx < N) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int N = 100000; // 向量大小
    size_t size = N * sizeof(int);

    // 主機 (CPU) 數據
    std::vector<int> h_a(N), h_b(N), h_c(N);

    // 初始化主機數據
    for (int i = 0; i < N; ++i) {
        h_a[i] = i;
        h_b[i] = i * 2;
    }

    // 設備 (GPU) 數據指針
    int *d_a, *d_b, *d_c;

    // 在設備上分配內存
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);

    // 將數據從主機複製到設備
    cudaMemcpy(d_a, h_a.data(), size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b.data(), size, cudaMemcpyHostToDevice);

    // 配置 Grid 和 Block 維度
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock; // 確保覆蓋所有元素

    // 啟動核函數
    add_vectors<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);

    // 將結果從設備複製回主機
    cudaMemcpy(h_c.data(), d_c, size, cudaMemcpyDeviceToHost);

    // 驗證結果 (可選)
    std::cout << "h_c[0] = " << h_c[0] << std::endl;      // 0 + 0 = 0
    std::cout << "h_c[1] = " << h_c[1] << std::endl;      // 1 + 2 = 3
    std::cout << "h_c[N-1] = " << h_c[N-1] << std::endl; // (N-1) + 2*(N-1) = 3*(N-1)

    // 釋放設備內存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

總結

理解 CUDA 的 Grid 和 Block 概念對於有效利用 GPU 的並行計算能力至關重要。通過合理地組織線程，你可以將複雜的計算任務分解為數百萬個小任務，並在 GPU 上高效地執行它們。