CUDA 编程模型：Thread / Block / Grid

CUDA 的编程模型是其最核心的抽象，理解三级线程层次结构是写出正�?CUDA 程序的基础�?

三级线程层次

CUDA 将并行执行的线程组织为三个层次：

Grid（网格）
└── Block（线程块）�?N
    └── Thread（线程）× M

对应硬件�?
Grid     �?整个 GPU（所�?SM�?
Block    �?单个 SM（共享内存范围）
Thread   �?单个 CUDA Core

可视化结�?

Grid (gridDim.x=3, gridDim.y=2)
┌─────────────────────────────────────�?
�? Block(0,0)  Block(1,0)  Block(2,0) �?
�? Block(0,1)  Block(1,1)  Block(2,1) �?
└─────────────────────────────────────�?

每个 Block (blockDim.x=4, blockDim.y=4)
┌───────────────────�?
�?T(0,0) T(1,0) T(2,0) T(3,0) �?
�?T(0,1) T(1,1) T(2,1) T(3,1) �?
�?T(0,2) T(1,2) T(2,2) T(3,2) �?
�?T(0,3) T(1,3) T(2,3) T(3,3) �?
└───────────────────�?

Kernel 函数

Kernel 是在 GPU 上执行的函数，用 __global__ 修饰�?

// Kernel 定义
__global__ void vectorAdd(float* A, float* B, float* C, int N) {
    // 计算全局线程索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查（线程数可能多于数据量�?
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// Kernel 启动
int N = 1024 * 1024;
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;  // 向上取整

vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

函数修饰�?

修饰�?	执行位置	调用位置	说明
__global__	GPU	CPU �?GPU	Kernel 函数
__device__	GPU	GPU	设备函数（不能直接从 CPU 调用�?
__host__	CPU	CPU	普�?CPU 函数（默认）
__host__ __device__	CPU + GPU	CPU + GPU	两端都可调用

内置变量

每个线程可以访问以下内置变量�?

__global__ void kernel() {
    // 线程�?Block 内的索引�?D�?
    threadIdx.x, threadIdx.y, threadIdx.z
    
    // Block �?Grid 内的索引�?D�?
    blockIdx.x, blockIdx.y, blockIdx.z
    
    // Block 的尺�?
    blockDim.x, blockDim.y, blockDim.z
    
    // Grid 的尺寸（�?Block 为单位）
    gridDim.x, gridDim.y, gridDim.z
}

计算全局索引

// 1D Grid, 1D Block
int idx = blockIdx.x * blockDim.x + threadIdx.x;

// 2D Grid, 2D Block（处理矩阵）
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
int idx = row * width + col;

// 3D（处理体数据�?
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

启动配置 <<<>>>

kernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args...);

参数	类型	说明
gridDim	dim3	Grid 的尺寸（Block 数量�?
blockDim	dim3	Block 的尺寸（Thread 数量�?
sharedMemBytes	size_t	动态共享内存大小（可选，默认0�?
stream	cudaStream_t	执行流（可选，默认0�?

dim3 类型

// dim3 可以用整数或三元组初始化
dim3 blockDim(16, 16, 1);   // 16×16 �?2D Block
dim3 gridDim(64, 64, 1);    // 64×64 �?2D Grid

// 等价写法
kernel<<<dim3(64,64), dim3(16,16)>>>(args);
kernel<<<64*64, 16*16>>>(args);  // 1D 简写（总数相同但语义不同！�?

Block 设计原则

线程数必须是 32 的倍数

// 好：256 = 8 × 32，完整的 Warp
dim3 blockDim(256);

// 差：100 个线�?�?4 �?Warp，最后一�?Warp 只有 4 个活跃线�?
dim3 blockDim(100);

推荐 Block 大小

常用选择�?28, 256, 512
最优值取决于�?
  - Kernel 的寄存器使用�?
  - 共享内存使用�?
  - 访存模式
  
经验法则：从 256 开始，�?Nsight 测量 Occupancy

最大限�?

限制	�?
每个 Block 最大线程数	1024
Block 每维最大尺�?	x: 1024, y: 1024, z: 64
Grid 每维最大尺�?	x: 2³¹-1, y: 65535, z: 65535
每个 SM 最�?Block �?	32 (Ampere)
每个 SM 最�?Warp �?	64 (Ampere)

同步机制

Block 内同�?

__global__ void kernel(float* data) {
    __shared__ float smem[256];
    
    // 每个线程加载数据到共享内�?
    smem[threadIdx.x] = data[blockIdx.x * blockDim.x + threadIdx.x];
    
    // 等待 Block 内所有线程完成加�?
    __syncthreads();
    
    // 现在可以安全读取其他线程写入的数�?
    float val = smem[(threadIdx.x + 1) % blockDim.x];
}

注意

__syncthreads() 只能同步同一�?Block 内的线程�?不能�?Block 同步*。跨 Block 同步需要通过 Kernel 结束（隐式全局同步）或使用 Cooperative Groups�?

Cooperative Groups（现代同步方式）

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel() {
    // Block 级同步组
    cg::thread_block block = cg::this_thread_block();
    block.sync();  // 等价�?__syncthreads()
    
    // Warp 级同步组
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);
    warp.sync();
    
    // Warp �?Shuffle 操作
    float val = warp.shfl_down(myVal, 1);  // 从相邻线程获取�?
}

完整示例：矩阵乘�?

#define TILE_SIZE 16

__global__ void matMul(float* A, float* B, float* C, int N) {
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载 Tile 到共享内�?
        tileA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        __syncthreads();
        
        // 计算部分�?
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

// 启动
dim3 blockDim(TILE_SIZE, TILE_SIZE);
dim3 gridDim(N / TILE_SIZE, N / TILE_SIZE);
matMul<<<gridDim, blockDim>>>(d_A, d_B, d_C, N);

---

下一篇：CUDA C/C++ 编程基础 →