跳到主要内容

线程层次结构

CUDA 采用层次化的线程组织结构，从最小单位线程到线程块再到网格，形成三级抽象。理解这个层次结构是编写高效 CUDA 程序的关键。

线程层次概述​

CUDA 的线程组织结构分为三个层次：

Grid（网格）
  └── Block（线程块）
        └── Thread（线程）
层次说明数量限制
Thread最小执行单元，执行核函数代码-
Block多个线程的集合，可同步和共享内存最多 1024 个线程
Grid多个线程块的集合，执行一个核函数最多 231−12^{31}-1 个 Block（一维）

线程（Thread）​

线程是 CUDA 程序的最小执行单元。当启动一个核函数时，GPU 会创建成千上万个线程，每个线程独立执行相同的代码。

线程索引​

每个线程都有一个唯一的标识符，通过内置变量 threadIdx 访问：

__global__ void kernel() {
    int tid = threadIdx.x;
    printf("Thread %d\n", tid);
}

threadIdx 是一个三维向量，包含三个分量：

  • threadIdx.x：线程在 X 方向的索引
  • threadIdx.y：线程在 Y 方向的索引
  • threadIdx.z：线程在 Z 方向的索引

为什么是三维？​

三维索引的设计源于实际应用需求：

维度适用场景
一维数组、向量运算
二维图像处理、矩阵运算
三维体积数据、CT 扫描、流体模拟

这种设计让代码更直观，不需要手动计算坐标转换。

线程块（Block）​

线程块是多个线程的集合。同一个 Block 内的线程可以：

  1. 通过共享内存交换数据
  2. 通过同步屏障进行同步
  3. 在同一个 SM（流多处理器）上执行

Block 索引​

每个 Block 在 Grid 中有一个唯一标识符 blockIdx：

__global__ void kernel() {
    int bid = blockIdx.x;
    printf("Block %d, Thread %d\n", bid, threadIdx.x);
}

Block 维度​

blockDim 表示每个 Block 中线程的数量：

__global__ void kernel() {
    printf("Block dimension: %d\n", blockDim.x);
}

Block 大小限制​

每个 Block 最多可以有 1024 个线程，但维度可以灵活配置：

dim3 block1D(1024);
dim3 block2D(32, 32);
dim3 block3D(8, 8, 16);

注意：32 × 32 = 1024，8 × 8 × 16 = 1024，都不能超过 1024。

Block 大小选择​

Block 大小的选择需要考虑以下因素：

  1. Warp 大小：GPU 以 Warp（32 个线程）为单位调度线程，Block 大小应该是 32 的倍数
  2. 资源限制：每个 SM 的寄存器和共享内存有限，Block 越大，能同时运行的 Block 越少
  3. 占用率：需要平衡 Block 大小和 SM 占用率

常用的 Block 大小：64、128、256、512、1024

int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
kernel<<<gridSize, blockSize>>>(args);

网格（Grid）​

网格是多个线程块的集合，是核函数启动时创建的最高层次结构。

Grid 维度​

gridDim 表示 Grid 中 Block 的数量：

__global__ void kernel() {
    printf("Grid dimension: %d\n", gridDim.x);
}

Grid 大小限制​

Grid 大小受硬件限制：

维度x 方向y 方向z 方向
最大 Block 数231−12^{31}-16553565535

计算全局线程索引​

一维 Grid 和 Block：

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

二维 Grid 和 Block：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int globalIdx = y * gridDim.x * blockDim.x + x;

三维 Grid 和 Block：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
int globalIdx = z * gridDim.x * blockDim.x * gridDim.y * blockDim.y 
              + y * gridDim.x * blockDim.x + x;

线程组织示例​

一维示例：向量运算​

__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1000000;
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
}

二维示例：矩阵运算​

__global__ void matrixAdd(float *A, float *B, float *C, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (x < width && y < height) {
        int idx = y * width + x;
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int width = 1024, height = 1024;
    
    dim3 blockSize(16, 16);
    dim3 gridSize((width + 15) / 16, (height + 15) / 16);
    
    matrixAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, width, height);
}

三维示例：体积数据处理​

__global__ void volumeProcess(float *volume, int nx, int ny, int nz) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int z = blockIdx.z * blockDim.z + threadIdx.z;
    
    if (x < nx && y < ny && z < nz) {
        int idx = z * nx * ny + y * nx + x;
        volume[idx] *= 2.0f;
    }
}

int main() {
    int nx = 128, ny = 128, nz = 128;
    
    dim3 blockSize(8, 8, 8);
    dim3 gridSize((nx + 7) / 8, (ny + 7) / 8, (nz + 7) / 8);
    
    volumeProcess<<<gridSize, blockSize>>>(d_volume, nx, ny, nz);
}

Warp 与 SIMT 执行模型​

Warp 概念​

Warp 是 GPU 执行线程的基本单位，包含 32 个线程。同一个 Warp 中的线程：

  1. 以 SIMD（单指令多数据）方式执行
  2. 同一时刻执行相同的指令
  3. 处理不同的数据

Warp 分组​

Block 中的线程按照 threadIdx 顺序分组为 Warp：

int warpId = threadIdx.x / 32;
int laneId = threadIdx.x % 32;

Warp 分化​

当同一个 Warp 中的线程执行不同的代码路径时，会发生 Warp 分化：

__global__ void kernel(int *data) {
    int idx = threadIdx.x;
    
    if (idx % 2 == 0) {
        data[idx] = data[idx] * 2;
    } else {
        data[idx] = data[idx] + 1;
    }
}

这种情况下，Warp 会串行执行两个分支，降低性能。优化方法：

__global__ void kernel(int *data) {
    int idx = threadIdx.x;
    
    if (idx < 16) {
        data[idx] = data[idx] * 2;
    } else {
        data[idx] = data[idx] + 1;
    }
}

硬件映射关系​

线程层次结构与 GPU 硬件的对应关系：

软件概念硬件概念说明
ThreadCUDA Core执行单元
WarpWarp Scheduler调度单位（32 线程）
BlockSM同一 Block 在同一 SM 上执行
GridGPU Device整个核函数

SM 资源分配​

一个 SM 可以同时执行多个 Block，受以下资源限制：

  1. 最大线程数：每个 SM 最多 2048 个线程
  2. 最大 Block 数：每个 SM 最多 32 个 Block
  3. 寄存器数量：每个 SM 64K 个 32 位寄存器
  4. 共享内存：每个 SM 48KB-164KB

占用率计算​

占用率 = 活跃 Warp 数 / 最大 Warp 数

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);

int maxThreadsPerSM = prop.maxThreadsPerMultiProcessor;
int activeThreads = blocksPerSM * threadsPerBlock;
float occupancy = (float)activeThreads / maxThreadsPerSM;

实用工具函数​

计算网格大小​

int ceilDiv(int a, int b) {
    return (a + b - 1) / b;
}

int gridSize = ceilDiv(n, blockSize);

获取线程全局索引​

__device__ int getGlobalIdx() {
    return blockIdx.x * blockDim.x + threadIdx.x;
}

__device__ int2 getGlobalIdx2D() {
    return make_int2(
        blockIdx.x * blockDim.x + threadIdx.x,
        blockIdx.y * blockDim.y + threadIdx.y
    );
}

边界检查宏​

#define IN_BOUNDS_1D(n) (blockIdx.x * blockDim.x + threadIdx.x < n)
#define IN_BOUNDS_2D(w, h) \
    (blockIdx.x * blockDim.x + threadIdx.x < w && \
     blockIdx.y * blockDim.y + threadIdx.y < h)

__global__ void kernel(float *data, int n) {
    if (IN_BOUNDS_1D(n)) {
        int idx = blockIdx.x * blockDim.x + threadIdx.x;
        data[idx] *= 2.0f;
    }
}

小结​

本章详细介绍了 CUDA 的线程层次结构：

  1. Thread：最小执行单元，通过 threadIdx 标识
  2. Block：线程集合，可同步和共享内存，最多 1024 线程
  3. Grid：Block 集合，执行一个核函数
  4. Warp：32 个线程，GPU 的调度单位
  5. 硬件映射：Thread → Core，Block → SM，Grid → GPU

理解线程层次结构是编写高效 CUDA 程序的基础。下一章将介绍 内存模型，学习如何高效地管理 GPU 内存。