📚 课程目录第4章线程层次结构

线程层次结构：线程、块与网格

CUDA 的并行执行模型建立在一种三层线程组织之上：线程（Thread）是最小的执行单元；若干线程组成一个线程块（Block）；若干线程块再组成一个网格（Grid）。这种层次结构并非人为设计的繁琐概念，而是直接映射到 GPU 物理硬件架构的关键抽象——它决定了线程如何在 SM（流式多处理器）上调度、如何访问不同级别的存储器，以及如何通过索引定位自己负责的数据。

核心原理：从硬件到软件的映射

理解线程层次，首先要建立硬件与软件的对应关系。GPU 内部包含多个 SM，每个 SM 拥有若干 CUDA 核心。当内核启动时，网格被分配到 GPU 设备，而线程块则被调度到具体的 SM 上。一个 SM 可以同时驻留多个线程块，但一个线程块只能在一个 SM 上执行，不能跨 SM 拆分。线程块内部的线程通过共享内存（Shared Memory）协作，而不同线程块之间无法直接通信，只能通过全局内存间接交互。

关键映射关系：

Grid → GPU 设备：整个网格对应一次内核启动，由所有 SM 共同完成。
Block → SM：线程块是调度到 SM 的基本单位，SM 以块为单位接收任务。
Thread → CUDA Core：线程块内的线程被进一步分配到 SM 的各个 CUDA 核心上执行。

这种层次结构带来的直接好处是可扩展性：同一份内核代码，在只有 2 个 SM 的入门级 GPU 和拥有 132 个 SM 的旗舰级 GPU 上都能正确运行。硬件自动将线程块分配到可用的 SM 上，无需程序员干预。同时，块内线程的同步（__syncthreads()）和共享内存访问，为细粒度协作提供了高效手段。

线程索引与数据映射方法

每个线程在运行时都可以通过内置变量获取自己的位置：threadIdx（线程在块内的索引）、blockIdx（块在网格中的索引）、blockDim（块的维度大小）、gridDim（网格的维度大小）。这些变量可以是 dim3 类型，支持一维、二维或三维的索引空间，极大方便了不同维度数据的处理。

对于一维数组的处理，最常用的映射公式为：

// 内核定义
__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];
    }
}

// 内核启动配置
int N = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(A, B, C, N);

对于二维矩阵，可以使用二维的线程块和网格，使索引更直观：

// 二维矩阵加法
__global__ void matrixAdd(float *A, float *B, float *C, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < width && col < width) {
        int idx = row * width + col;
        C[idx] = A[idx] + B[idx];
    }
}

dim3 blockDim(16, 16);           // 每个块 16×16 = 256 线程
dim3 gridDim((width + 15) / 16, (width + 15) / 16);
matrixAdd<<<gridDim, blockDim>>>(A, B, C, width);

索引设计要点：

块大小通常取 128、256 或 512，这些值能较好地平衡 SM 的资源利用率和调度开销。
块数应远大于 SM 数量（通常为 SM 数的 10 倍以上），以保证负载均衡和延迟隐藏。
始终进行边界检查，因为数据长度不一定能被块大小整除。
优先使用一维索引，除非二维/三维索引能显著提升代码可读性。

案例：图像灰度化中的二维映射

假设有一张 1920×1080 的 RGB 图像需要转换为灰度图。每个像素由 R、G、B 三个分量组成，输出为单个灰度值。这是一个典型的二维数据并行问题。

我们采用 16×16 的线程块，每个线程处理一个像素。网格尺寸计算如下：

dim3 blockDim(16, 16);
dim3 gridDim((1920 + 15) / 16, (1080 + 15) / 16);
// gridDim = (120, 68)，共 8160 个线程块
// 总线程数 = 8160 × 256 = 2,088,960 > 2,073,600（像素数）

__global__ void rgbToGray(const unsigned char *rgb,
                          unsigned char *gray,
                          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) * 3;  // RGB 连续存储
        unsigned char r = rgb[idx];
        unsigned char g = rgb[idx + 1];
        unsigned char b = rgb[idx + 2];
        gray[y * width + x] = (unsigned char)(0.299f * r + 0.587f * g + 0.114f * b);
    }
}

在这个案例中，二维索引直接对应像素的 (x, y) 坐标，无需额外的坐标换算。每个线程独立处理一个像素，线程之间无通信，属于embarrassingly parallel（易并行）问题。通过合理设置块大小，可以确保所有 SM 都满载工作。

常见误区与调试复盘

误区	后果	正确做法
块大小设置过大（如 1024 以上）	SM 资源不足，内核启动失败或占用率低	控制在 128~512 之间，并考虑寄存器/共享内存限制
忘记边界检查	数组越界访问，产生未定义行为或崩溃	始终用 `if (idx < N)` 保护全局内存访问
混淆 `blockDim` 与 `gridDim`	索引计算错误，数据错位	牢记：`blockDim` 是块内线程数，`gridDim` 是网格内块数
网格维度不足	部分数据未被任何线程处理	使用 `(N + blockDim.x - 1) / blockDim.x` 向上取整
假设块内线程数等于 SM 线程数	误解调度机制，性能调优方向错误	理解块是调度单位，SM 可同时驻留多个块

典型排障案例：某同学编写向量加法时，将 blockDim.x 误写为 gridDim.x，导致所有线程计算出的索引都落在前几个元素上，后段数据未被处理。调试时使用 cuda-memcheck 未发现越界，但输出结果中后半部分全为 0。最终通过打印 blockIdx.x、blockDim.x、threadIdx.x 的值对比，才定位到索引公式错误。建议在开发初期，用 printf 输出少量线程的索引值进行验证。

线程层次结构是 CUDA 编程的基石。掌握 threadIdx、blockIdx、blockDim 与 gridDim 的语义，并能熟练推导全局索引与数据的映射关系，是编写正确、高效内核程序的前提。建议初学者从一维向量加法开始，逐步过渡到二维矩阵运算，在反复练习中建立对线程组织方式的直觉。