4. 标量指令与向量指令
好，咱们今天聊聊GPU里最核心的一对概念——标量指令和向量指令。很多人刚接触GPU编程时，总觉得向量指令就是“快”，标量指令就是“慢”。其实没那么简单。我做了这么多年GPU架构，见过太多因为没搞懂这两者关系，导致性能翻车的案例。
说白了，标量指令是给单个数据干活，向量指令是一次性给一堆数据干活。但GPU的厉害之处，就在于怎么把这两者巧妙地结合起来。
4.1 标量指令：地址计算与控制流
标量指令处理的是单个数据元素。你想想看，GPU里有成千上万个线程，每个线程都需要知道自己该访问哪个地址、该往哪跳转。这些活儿，就是标量指令干的。
4.1.1 地址计算指令
地址计算是标量指令的重头戏。我遇到过不少开发者，觉得地址计算不就是加加减减吗？有什么好优化的？嗯，这里要注意——地址计算的效率，直接决定了访存带宽能不能跑满。
常见的地址计算指令包括：
- 基址+偏移：最基础的模式，比如
ld.global.f32 %r0, [%base + %offset] - 索引计算：处理多维数组时，需要把 (i, j, k) 转成线性地址
- 对齐检查：确保地址满足对齐要求，否则会触发慢速路径
我曾经在一个图像处理项目里，发现性能始终上不去。查了半天，原来是地址计算指令生成了太多冗余的移位和掩码操作。后来我把地址计算逻辑手动优化了一下，性能直接提升了30%。
关键点：标量地址计算指令虽然简单，但它们是访存流水线的“咽喉”。每多一条地址计算指令，就可能多一个时钟周期的延迟。
4.1.2 控制流指令
控制流指令是标量指令的另一个重要角色。分支、循环、函数调用，这些都靠它。
GPU处理控制流的方式很有意思。它不像CPU那样有复杂的分支预测器，而是采用一种叫“预测执行”的机制。说白了，就是两条路都走，最后选一条。
我建议你记住这个原则：在GPU里，分支越少越好，分支越统一越好。为什么？因为同一个warp里的32个线程，如果走不同的分支，那它们就得串行执行。你想想看，32个线程本来可以并行跑，现在变成串行，性能直接掉到1/32。
避坑指南：我曾经在一个深度学习训练脚本里，看到有人用 if (threadIdx.x % 2 == 0) 这样的分支。这会导致warp里的线程一半走true，一半走false，性能损失巨大。正确的做法是用向量化的方式处理，或者重新组织数据布局。
4.2 向量指令：算术、逻辑与类型转换
向量指令才是GPU真正发力的地方。一条向量指令，可以同时对多个数据元素执行相同的操作。
4.2.1 算术指令
向量算术指令包括加法、乘法、乘加（FMA）、倒数、平方根等。其中FMA是最常用的，因为它把乘法和加法合并成一条指令，既减少了指令数，又提高了精度。
举个例子，在NVIDIA的GPU上，一条 HFMA2 指令可以同时做两个半精度浮点的乘加运算：
// 伪代码：同时计算 a*b+c 和 d*e+f
HFMA2 R0, R1, R2, R3 // R0.x = R1.x*R2.x + R3.x
// R0.y = R1.y*R2.y + R3.y我个人习惯在写shader时，尽量把计算组织成向量形式。比如处理颜色时，把RGBA四个通道一起算，而不是一个一个来。这样编译器更容易生成向量指令。
4.2.2 逻辑指令
向量逻辑指令处理的是位运算和比较操作。比如AND、OR、XOR、比较大小等。这些指令在图形渲染和通用计算中都很常见。
我记得有一次做物理模拟，需要判断大量粒子的位置是否在某个区域内。如果用标量指令一个个判断，那效率太低了。后来我用向量比较指令，一次比较8个粒子，速度提升了将近8倍。
4.2.3 类型转换指令
类型转换是容易被忽视但非常重要的向量指令。GPU里常见的数据类型有：
| 源类型 | 目标类型 | 典型指令 |
|---|---|---|
| float32 | float16 | CVT.F32.F16 |
| float16 | float32 | CVT.F16.F32 |
| int32 | float32 | CVT.F32.S32 |
| float32 | int32 | CVT.S32.F32 |
这里有个坑：类型转换指令的延迟通常比普通算术指令高。我曾经在优化一个神经网络推理引擎时，发现大量时间花在了float32到float16的转换上。后来我调整了数据流水线，让转换操作和计算操作重叠执行，才把性能提上来。
小技巧：如果你的数据精度要求不高，尽量用float16或int8。这样一条向量指令可以处理更多数据元素。比如float16的向量宽度通常是float32的两倍。
4.3 指令发射与调度
好了，现在我们知道标量指令和向量指令各自干什么了。但问题是，GPU怎么决定什么时候发射哪条指令？这就是指令发射与调度要解决的问题。
4.3.1 指令发射的基本流程
GPU的指令发射单元，说白了就是一个“指挥官”。它从指令缓存里取出指令，检查指令的依赖关系，然后决定什么时候把指令发射到执行单元。
这个过程大致是这样的：
- 取指：从指令缓存中取出下一条指令
- 译码：解析指令的操作码和操作数
- 依赖检查：检查指令需要的寄存器是否已经被前面的指令占用
- 发射：如果依赖满足，就把指令发送到对应的执行单元
我建议你记住一个关键点：指令发射的瓶颈往往不在发射单元本身，而在数据依赖。如果一条指令需要等待前一条指令的结果，那它就只能干等着。
4.3.2 标量与向量的调度差异
标量指令和向量指令的调度策略是不同的。标量指令通常由标量处理器（SP）执行，而向量指令由向量处理器（VP）执行。这两者可以并行工作。
举个例子，一个典型的调度场景是这样的：
// 标量指令：计算地址
IADD R_addr, R_base, R_offset
// 向量指令：加载数据（依赖R_addr）
LD.GLOBAL.F32 V_data, [R_addr]
// 向量指令：计算（依赖V_data）
VADD V_result, V_data, V_const在这个例子里，标量指令 IADD 先发射，计算地址。然后向量加载指令 LD.GLOBAL.F32 发射，但需要等待地址计算完成。最后向量加法指令 VADD 发射，需要等待加载完成。
我曾经在调试一个性能问题时，发现指令发射单元经常处于“空转”状态。原因是标量指令和向量指令之间的依赖太紧密，导致流水线经常停顿。后来我通过重新排列指令顺序，把不依赖的指令提前发射，才把吞吐量提上去。
核心思想：指令调度的本质，就是让执行单元尽量“不闲着”。标量指令和向量指令如果能并行执行，那是最好的。如果不行，那就尽量让它们错开依赖关系。
4.3.3 常见的调度策略
不同的GPU厂商有不同的调度策略。我简单说几种常见的：
- 静态调度：编译器在编译时就决定指令的发射顺序。优点是简单，缺点是灵活性差。
- 动态调度：硬件在运行时根据依赖关系动态调整指令发射顺序。优点是灵活，缺点是硬件复杂度高。
- 混合调度：结合静态和动态调度的优点。编译器做初步优化，硬件做最终调整。
我个人比较喜欢混合调度。因为编译器可以做一些全局优化，而硬件可以处理一些运行时的不确定性。两者配合，效果最好。
注意：不要过度依赖硬件动态调度。我曾经见过一个案例，开发者写了大量依赖链很长的代码，指望硬件自动优化。结果硬件调度器被复杂的依赖关系搞懵了，性能反而更差。记住，编译器优化和手写优化同样重要。
4.4 小结
好了，这一章的内容就到这里。我们讲了标量指令（地址计算、控制流）、向量指令（算术、逻辑、类型转换），以及指令发射与调度的基本概念。
最后送你一句话：标量指令是骨架，向量指令是肌肉，指令调度是大脑。三者缺一不可。下次写GPU程序时，多想想你的代码里标量和向量指令的比例是否合理，指令依赖是否太紧密。相信我，这些小细节会带来大不同。