3、GPU编程模型入门：并行计算基本概念、SIMT模型、线程层次结构

好，咱们进入正题。这一章我打算聊聊GPU编程模型的核心思想。说实话，很多初学者一上来就盯着CUDA语法看，结果越看越懵。为什么？因为没搞懂GPU到底是怎么工作的。

我当年刚接触GPU编程时，也踩过不少坑。有一次写了个并行归约算法，跑出来的结果总是错的。查了两天，最后发现是没理解线程同步的粒度问题。嗯，从那以后我就明白了——搞懂编程模型，比会写几行CUDA代码重要得多。

3.1 并行计算的基本概念

先说说并行计算。说白了，就是同时干多件事。

CPU是「能干很多种活，但一次只能干几件」。GPU是「只能干一种活，但一次能干成千上万件」。你想想看，处理一张1080p的图片，每个像素都要做同样的运算——这不就是GPU的强项吗？

并行计算一般分两种：

  • 任务并行：不同的核心干不同的事。比如一个核做图像滤波，另一个核做特征提取。
  • 数据并行：所有核心干同样的事，但处理不同的数据。GPU最擅长的就是这个。

我在做嵌入式视觉项目时，90%以上的加速都来自数据并行。比如YOLO的卷积层，每个输出像素的计算逻辑完全一样，只是输入数据不同。这种场景，GPU简直是天生的匹配。

核心观点：GPU加速的本质，就是把「串行循环」变成「并行展开」。一个for循环跑1000次，在GPU上可能就是1000个线程同时跑一次。

3.2 SIMT模型：单指令多线程

接下来聊SIMT。这个名词听起来高大上，其实理解起来不难。

SIMT的全称是Single Instruction, Multiple Threads。翻译过来就是「一条指令，多个线程同时执行」。

打个比方：

  • 你是一个教官，喊了一声「向右转」。
  • 100个士兵同时执行这个指令。
  • 这就是SIMT——指令只有一个，执行者是多个。

GPU里也是这样。一个指令发射出去，32个线程（也就是一个Warp）同时执行。这32个线程干的是同一件事，只是处理的数据不同。

为什么会这样设计？因为GPU的硬件资源有限。如果每个线程都独立取指令、译码、执行，那控制单元的成本就太高了。共享指令单元，才能把更多的晶体管留给计算单元。

个人经验：我在优化一个3x3卷积核时，发现性能始终上不去。后来一查，是因为线程分支太多。32个线程里，有的走if分支，有的走else分支，导致指令流水线被「串行化」了。说白了，SIMT最怕的就是分支发散。

这里有个关键点要注意：

情况 性能影响 说明
所有线程走同一分支 无影响 SIMT正常执行，32个线程同步
部分线程走不同分支 性能下降 不同分支会被串行执行，先执行if，再执行else
每个线程都不同 性能极差 几乎退化为串行，32个线程依次执行

我曾经在一个项目中，把条件判断挪到了循环外面，性能直接提升了3倍。嗯，这就是理解SIMT模型带来的好处。

3.3 线程层次结构：Grid / Block / Warp

好，接下来是重头戏。GPU的线程不是乱糟糟的一团，而是有严格的层次结构。理解这个结构，你才能写出高效的并行程序。

3.3.1 三层结构概览

从大到小，依次是：

  • Grid（网格）：最顶层，包含多个Block。一个Kernel启动时，对应一个Grid。
  • Block（线程块）：中间层，包含多个线程。同一个Block内的线程可以共享数据。
  • Warp（线程束）：最底层，32个线程一组。这是硬件调度的基本单位。

我习惯这么理解：

  • Grid就像一栋楼
  • Block就像楼里的房间
  • Warp就像房间里的小组
  • 线程就是每个人

3.3.2 Grid和Block的配置

写CUDA代码时，第一件事就是配置Grid和Block的大小。比如：

// 处理一张1024x1024的图像
dim3 blockSize(16, 16);    // 每个Block有16x16=256个线程
dim3 gridSize(64, 64);     // Grid有64x64=4096个Block

kernel<<<gridSize, blockSize>>>(input, output);

这里要注意：

  • Block的总线程数不能超过1024（不同架构限制不同）
  • Grid的维度可以是一维、二维或三维
  • Block的维度最好选16、32这样的2的幂次

避坑指南：我曾经把Block大小设成了32x32=1024个线程，刚好卡在上限。结果在旧架构的Jetson Nano上跑，程序直接崩溃。后来查文档才知道，Maxwell架构的Block最大线程数是1024，但还要考虑寄存器占用。所以，别卡着上限设，留点余量。

3.3.3 Warp：硬件调度的最小单元

Warp是GPU硬件真正执行指令的单位。32个线程组成一个Warp，它们共享程序计数器，执行同一条指令。

举个例子：

  • 你启动了一个Block，里面有128个线程
  • 硬件会把128个线程分成4个Warp（128 / 32 = 4）
  • 每个Warp独立调度，互不影响

这里有个重要的性能要点：

Warp的占用率。说白了，就是SM（流多处理器）上同时活跃的Warp数量。占用率越高，越能隐藏内存访问延迟。

我一般这么估算：

SM上的Warp数量 = (Block数 × 每个Block的Warp数) / SM数

比如一个SM最多能放64个Warp，你只放了32个，那占用率就是50%。这时候内存延迟就藏不住了，性能会打折扣。

3.3.4 线程索引的计算

每个线程在Grid中的位置，由三个内置变量决定：

  • threadIdx：线程在Block内的索引
  • blockIdx：Block在Grid内的索引
  • blockDim：Block的维度大小

计算全局索引的公式：

// 一维情况
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维情况（处理图像时常用）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;

我建议你把这个公式记牢。每次写Kernel时，第一件事就是算好线程的全局索引。索引算错了，后面的计算全白搭。

小技巧：调试时，我经常先让每个线程打印自己的索引。虽然慢，但能快速确认索引映射是否正确。等逻辑对了，再把打印去掉。

3.4 总结与思考

这一章的内容，说白了就三件事：

  1. 并行计算：GPU擅长数据并行，不擅长任务并行
  2. SIMT模型：32个线程共享指令，最怕分支发散
  3. 线程层次：Grid → Block → Warp，理解这个才能写出高效代码

我个人觉得，学GPU编程最难的不是语法，而是思维方式的转变。从「串行思维」切换到「并行思维」，需要一点时间。别急，多写几个Kernel，慢慢就习惯了。

下一章，我会带你手写第一个CUDA程序。到时候咱们把今天学的这些概念，真正用起来。

← 上一章 📖 返回目录 下一章 →