3、GPU编程模型入门：CUDA/OpenCL基本概念、线程层次结构（grid、block、thread）

GPU之所以能实现数千个核心并行计算，其根本原因在于它采用了与CPU截然不同的执行模型。CPU依赖大容量缓存和复杂控制逻辑来加速单线程性能，而GPU则将大部分晶体管用于算术逻辑单元（ALU）和线程调度器，通过海量线程的并行执行来隐藏访存延迟。理解CUDA（Compute Unified Device Architecture）和OpenCL（Open Computing Language）中的线程层次结构，是掌握GPU编程的第一道门槛。本章将从编程模型的基本概念出发，深入剖析grid、block、thread三级层次的设计逻辑，并通过实际案例帮助读者建立从“硬件架构”到“软件抽象”的映射思维。

上图展示了CUDA/OpenCL中核心的线程层次结构。在GPU编程模型中，Grid是内核启动时创建的最大线程集合，它由多个Block组成；每个Block内部又包含若干Thread。这种三级结构并非随意设计，而是与GPU物理硬件架构（SM、CUDA Core、内存层级）紧密对应。下面我们将从核心原理、操作流程、案例拆解和常见误区四个维度展开。

3.1 核心原理：从CPU多线程到GPU海量线程的范式转换

CPU编程模型通常采用“任务并行”（Task Parallelism），即每个线程执行不同的指令流，线程间通过锁或原子操作进行通信。而GPU编程模型采用“数据并行”（Data Parallelism），即所有线程执行相同的指令（SIMT，Single Instruction Multiple Threads），但处理不同的数据元素。这种设计使得GPU可以高效地调度数千个轻量级线程，而无需像CPU那样为每个线程维护独立的寄存器上下文和复杂的分支预测器。

在CUDA中，kernel是运行在GPU上的函数，通过kernel<<<gridDim, blockDim>>>(args)语法启动。其中gridDim和blockDim可以是dim3类型，支持一维、二维或三维的线程组织。OpenCL则使用NDRange（N-Dimensional Range）概念，其global_work_size对应CUDA的Grid大小，local_work_size对应Block大小。两种模型在概念上高度一致，只是术语不同。

线程层次结构的设计带来了三个关键优势：可扩展性（程序无需知道实际GPU核心数即可运行）、局部性（Block内的线程可以共享数据，减少全局内存访问）、调度效率（硬件以Block为单位调度到SM上，Block内线程以Warp（32线程）为单位执行）。

3.2 判断方法与操作流程：如何确定线程层次配置

在实际编程中，选择合适的Grid和Block维度是性能优化的关键步骤。以下是标准的设计流程：

1. 确定数据规模：首先明确需要处理的数据总量N（例如图像像素数、矩阵元素数）。
2. 选择Block大小：通常设置为128或256个线程，这是经验上能较好平衡占用率和调度开销的值。Block大小必须是32（Warp大小）的倍数，以避免线程束分化。
3. 计算Grid大小：gridSize = (N + blockSize - 1) / blockSize，即向上取整。
4. 处理边界条件：由于N不一定能被blockSize整除，内核中需要添加if (idx < N)判断。
5. 考虑硬件限制：每个Block最大线程数不能超过1024（现代GPU），Grid的x维度最大为2^31-1，y/z维度最大为65535。

对于二维或三维数据（如图像、体素），应使用dim3组织线程。例如处理1024×1024的图像，可以设置blockDim(16, 16)，则每个Block包含256个线程，gridDim(64, 64)，共4096个Block。这种组织方式使得线程索引threadIdx.x和threadIdx.y能自然映射到像素坐标。

// CUDA示例：二维线程索引计算
__global__ void kernel(float* input, float* output, 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;
        output[idx] = input[idx] * 2.0f;
    }
}
// 启动配置
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (height + 15) / 16);
kernel<<<grid, block>>>(d_input, d_output, width, height);

3.3 案例拆解：矩阵乘法的线程层次设计

矩阵乘法是GPU编程的“Hello World”，也是检验线程层次设计能力的经典案例。假设我们要计算C = A × B，其中A为M×K矩阵，B为K×N矩阵，C为M×N矩阵。

初始方案：最简单的做法是每个线程计算C矩阵的一个元素。此时需要M×N个线程。若M=N=1024，则总线程数为1,048,576。我们设置Block为16×16=256线程，Grid为64×64=4096个Block。

性能瓶颈：每个线程需要从全局内存读取A的一行（K个元素）和B的一列（K个元素），共2K次全局内存访问。当K=1024时，每个线程需要2048次全局内存读取，而全局内存带宽远低于计算速度，导致计算单元大量空闲。

优化方案：利用Block内的共享内存。将矩阵分块为16×16的Tile，每个Block协作加载A的一个Tile和B的一个Tile到共享内存，然后进行16×16的局部乘法累加。这样每个全局内存数据被重复使用了16次（每个线程读取一次共享内存，但共享内存带宽远高于全局内存）。

关键参数与结果：在NVIDIA A100 GPU上，未优化版本性能约为0.8 TFLOPS，而使用共享内存分块后可达15 TFLOPS，提升近19倍。代价是代码复杂度显著增加，需要处理Bank Conflict、边界填充等问题。

3.4 常见误区与总结复盘

初学者在理解线程层次时容易陷入以下误区：

• 误区一：认为Block之间可以同步。实际上，Block之间无法直接同步，只能通过内核结束和重新启动来实现全局同步。跨Block通信必须通过全局内存，并需要原子操作或内存栅栏。
• 误区二：忽略Warp分化。当Block内线程执行不同分支时，所有Warp中的线程都会串行执行每个分支路径，导致性能急剧下降。应尽量让同一Warp内的线程执行相同指令。
• 误区三：混淆线程索引和全局索引。在二维网格中，全局索引的计算公式为int idx = blockIdx.x * blockDim.x + threadIdx.x，但若数据是行主序存储，还需加上blockIdx.y * gridDim.x * blockDim.x * blockDim.y等偏移。
• 误区四：忽视共享内存容量限制。每个SM的共享内存通常只有48KB（可配置到96KB），若每个Block使用过多共享内存，会减少SM上可驻留的Block数量，降低占用率。

总结复盘：线程层次结构是GPU编程的基石。理解Grid→Block→Thread的映射关系，不仅有助于编写正确程序，更是性能调优的前提。在实际项目中，建议先使用cudaOccupancyMaxPotentialBlockSize（CUDA）或clGetKernelWorkGroupInfo（OpenCL）查询硬件建议的Block大小，再结合具体算法特性进行调整。同时，使用NVIDIA Nsight或AMD ROCm Profiler等工具分析占用率、内存吞吐量和Warp状态，可以更科学地验证线程层次配置的合理性。