GPU 架构与 CUDA 编程：线程层次与内存模型

GPU 架构与 CUDA 编程核心：线程层次与内存模型

CUDA 编程的本质是让程序员高效驾驭 GPU 成千上万个轻量级线程。理解线程的组织方式（线程层次）与数据放置位置（内存模型）是写出高性能 CUDA 程序的基石。本章将从硬件架构出发，系统讲解这两个核心概念，并辅以直观的核函数示例。

1. GPU 的硬件架构缩影

在深入概念前，先看一张简化的 GPU 架构图。一块 NVIDIA GPU 由多个流式多处理器（Streaming Multiprocessor，SM） 组成，每个 SM 内部拥有大量 CUDA 核心、寄存器、共享内存、L1 缓存等。SM 是调度和执行线程的基本硬件单元。

当我们启动一个 CUDA 核函数时，会创建数十万甚至数百万个线程。这些线程被组织成线程块，并以线程块为单位被调度到不同的 SM 上执行。同一线程块内的线程会始终运行在同一个 SM 上，这为它们提供了极低延迟的数据共享通道。

2. 线程层次：网格、块与线程

CUDA 的线程层次用三个关键词概括：网格（Grid）、线程块（Block）、线程（Thread）。它们的关系就像是一个多维的军队编制。

2.1 线程块与线程

最基本的结构是 线程块（Block），它是一组可以彼此进行同步和通信的线程集合。在核函数调用中，我们需要指定每个块的线程数量，通常写成 blockDim.x、blockDim.y、blockDim.z，从而支持一维、二维或三维的块形状。

每个线程在块内有唯一的三维索引 threadIdx，取值范围从 (0,0,0) 到 (blockDim.x-1, blockDim.y-1, blockDim.z-1)。通过 threadIdx，我们可以将线程映射到需要处理的数据上，例如图像像素的 (行, 列)。

2.2 网格与块

多个线程块又组成一个 网格（Grid），代表一次核函数调用的全部线程。网格维度由 gridDim 定义，同样支持一、二或三维。每个线程块在网格中都有一个唯一的三维索引 blockIdx。因此，一个线程的全局唯一标识可以写为：

全局ID = blockIdx * blockDim + threadIdx （对各个维度分别计算）。

这种层次设计使得同样的核函数代码可以轻松扩展到不同规模的问题：只需改变网格和块的维度，无需改动核函数内部逻辑。

典型的使用模式：

// 一维示例：处理 N 个元素的向量
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < N) {
    // 由第 tid 个线程处理第 tid 个数据元素
}

2.3 线程束：硬件调度的最小单元

从硬件角度看，线程块内的线程以 线程束（Warp） 为单位执行。一个线程束通常包含 32 个线程，它们在 SM 上同时执行同一条指令，但各自处理不同的数据 (SIMT 模型)。线程块内的线程数量最好设计为 32 的倍数，以避免线程束浪费，并且 blockDim.x 推荐至少为 128 或 256 以掩藏内存延迟。

3. CUDA 内存模型：分层存储，各司其职

CUDA 程序的内存空间是分层的，不同层级内存的特性差异巨大。选择合适的内存层次直接决定了程序的数据吞吐效率。

3.1 全局内存是必经之路，但要合并访问

任何并行计算的数据最初都加载在全局内存中。访问全局内存时，最关键的原则是合并访问（Coalesced Access）：同一个线程束内的线程应访问连续的、对齐的内存地址，这样硬件可以将多次内存请求合并为一次或少数几次宽事务处理，极大提升带宽利用率。

例如，在访问一个浮点数组时，让 threadIdx.x 对应的线程访问连续的 array[threadIdx.x]，就是典型的合并访问。如果乱序访问，带宽将大幅下降。

3.2 共享内存：程序员的掌中缓存

共享内存是片上高速暂存器，速度仅次于寄存器。它允许同一线程块的线程协作加载数据，然后以极高的带宽频繁读写。常见的使用模式有：

• 数据块划入：将全局内存中的一小块数据协同加载到共享内存，让线程块内的线程反复使用，避免重复访问全局内存。
• 块内归约：在共享内存上进行求和、最大值等归约操作，通过 __syncthreads() 保证所有线程写入完成后再读取。

声明共享内存使用 __shared__ 关键字：

__global__ void kernel() {
    __shared__ float tile[BLOCK_SIZE][BLOCK_SIZE];
    // 协同加载数据到 tile，然后同步
}

3.3 寄存器和本地内存

寄存器是每个线程私有的最快存储，核函数中的标量变量通常都会放入寄存器中。但注意，编译器在寄存器压力过大会自动将变量“溢出”到本地内存，而本地内存实际上是位于显存中，访问代价高昂。因此，应避免在核函数内使用太多大数组或深层嵌套调用，以控制寄存器用量。

4. 案例：矩阵乘法中的线程层次与内存使用

把线程层次和内存模型融合到一个经典的矩阵乘法核函数中，会非常清晰。

假设我们要计算 C[M][N] = A[M][K] * B[K][N]，我们采用二维线程块，每个块处理一个 TILE_SIZE x TILE_SIZE 的小分块。

并行策略：

• 线程网格二维，x 维度对应 N / TILE_SIZE，y 维度对应 M / TILE_SIZE。
• 每个二维线程块包含 TILE_SIZE * TILE_SIZE 个线程。
• 每个线程计算 C 矩阵的一个元素。

内存使用：

• 每个线程块在共享内存中声明两个分块：__shared__ float As[TILE_SIZE][TILE_SIZE] 和 Bs[TILE_SIZE][TILE_SIZE]。
• 所有线程协作将 A、B 的对应分块从全局内存中合并载入到共享内存。
• 然后循环移位，加载下一个分块，并在每个分块累加部分乘积累积到寄存器变量 cValue。

简化核函数骨架：

__global__ void matMul(float* A, float* B, float* C, int M, int N, int K) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    float cValue = 0.0f;
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    int numTiles = (K + BLOCK_SIZE - 1) / BLOCK_SIZE;
    for (int t = 0; t < numTiles; ++t) {
        // 协作加载当前分块
        int aRow = row;
        int aCol = t * BLOCK_SIZE + threadIdx.x;
        As[threadIdx.y][threadIdx.x] = (aRow < M && aCol < K) ? A[aRow * K + aCol] : 0.0f;

        int bRow = t * BLOCK_SIZE + threadIdx.y;
        int bCol = col;
        Bs[threadIdx.y][threadIdx.x] = (bRow < K && bCol < N) ? B[bRow * N + bCol] : 0.0f;

        __syncthreads();

        // 分块内乘积累加
        for (int k = 0; k < BLOCK_SIZE; ++k) {
            cValue += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();
    }

    if (row < M && col < N)
        C[row * N + col] = cValue;
}

在这个例子中，线程的二维索引自然地映射到结果矩阵的行与列；共享内存分块大大减少了全局内存的读取次数；每个线程累加值 cValue 存放在寄存器中，循环变量等也优先使用寄存器。最终代码性能远超朴素版本的矩阵乘法。

5. 小结与最佳实践

• 线程层次：网格 → 线程块 → 线程。合理选择 blockDim 和 gridDim 以覆盖全部数据，同时确保线程块大小为 32 的倍数，推荐 128~512。
• 内存模型：从最快的寄存器开始，善用共享内存作为用户可控的缓存；全局内存务必保证合并访问；常量内存用于广播只读数据。
• 协同与同步：同一线程块内通过 __syncthreads() 保证数据可见性；线程块之间独立无同步。
• 思考维度：始终用线程的全局索引去思考问题的划分，用共享内存去减少全局通信。

掌握线程层次与内存模型，你就拥有了打开 CUDA 性能之门的第一把钥匙。接下来的实践，请不断尝试用不同的块维度和内存策略，测量程序的带宽和计算吞吐，在实验中深化理解。