第一章：GPU计算基础回顾
各位同学好，我是老张。在量化圈摸爬滚打了十几年，从最早的CPU多线程优化，到后来全面转向GPU加速，这条路我走得不算快，但踩过的坑确实不少。今天咱们先不急着上高深技巧，把地基夯实了再说。
1.1 CUDA编程模型——说白了就是“怎么指挥GPU干活”
CUDA的编程模型，核心就三个词：Grid、Block、Thread。你想想看，GPU里有成千上万个核心，你得告诉它们每个人该干什么。我刚开始接触时，总觉得这玩意儿跟CPU多线程差不多，后来发现完全不是一回事。
关键概念：一个Kernel函数启动时，会生成一个Grid，Grid里包含多个Block，每个Block里又包含多个Thread。Thread是实际执行计算的最小单位。
// 典型的CUDA Kernel启动方式
dim3 gridDim(32, 32); // 32x32 = 1024个Block
dim3 blockDim(16, 16); // 每个Block有16x16 = 256个Thread
kernel_name<<<gridDim, blockDim>>>(args);这里有个容易犯迷糊的地方：Block和Thread的维度不一定是二维的，也可以用一维或三维。我个人习惯用一维来处理向量运算，二维处理矩阵，三维处理体数据。嗯，这个选择其实会影响性能，后面会细说。
小技巧：Block里的Thread数量最好是32的倍数。为什么？因为线程束（Warp）的大小就是32。我曾经因为设了个31，结果性能直接掉了30%，排查了半天才发现是这问题。
1.2 内存层次结构——GPU的“仓库管理系统”
GPU的内存体系，说白了就是个金字塔结构。从上到下：寄存器、共享内存、L1/L2缓存、全局内存、常量内存、纹理内存。速度从快到慢，容量从小到大。
这张图我画了好几次才满意。你看，越靠近顶层的存储，速度越快但容量越小。做量化计算时，最常用的策略就是把热点数据搬到共享内存里。我记得有一次做期权定价的Monte Carlo模拟，把全局内存的频繁访问改成共享内存后，速度提升了将近5倍。
避坑指南：我曾经犯过一个低级错误——在Kernel里大量使用全局内存的随机访问。结果呢？内存带宽被严重浪费，性能惨不忍睹。记住：合并访问（Coalesced Access）是全局内存的命根子。
1.3 线程束调度原理——GPU的“隐形指挥官”
线程束（Warp）是GPU调度和执行的基本单位，大小固定为32个线程。你写代码时感觉是每个Thread独立运行，但实际上GPU是以Warp为单位来发指令的。
为什么会这样？因为GPU的设计哲学是“大量线程 + 简单控制”。一个SM（Streaming Multiprocessor）里有多个Warp Scheduler，它们轮流发射指令。如果一个Warp因为访存被卡住了，Scheduler会立刻切换到另一个Warp。这就是所谓的零开销上下文切换。
核心要点：Warp内的32个线程执行的是SIMT（单指令多线程）模式。如果出现分支分歧（比如if-else语句），Warp会串行执行所有分支路径，性能直接打折。
// 分支分歧的例子——尽量避免
__global__ void divergent_kernel(float* data, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N / 2) {
data[idx] = data[idx] * 2.0f; // 路径A
} else {
data[idx] = data[idx] * 0.5f; // 路径B
}
// 同一个Warp里的线程可能走不同路径，导致串行执行
}嗯，这里要注意：并不是所有分支都会导致性能问题。如果分支条件是基于线程ID的连续范围，比如threadIdx.x < 16，那Warp内的线程可能都走同一条路。我一般建议：能用三元运算符就别用if-else，能提前计算就别在Kernel里做条件判断。
| 特性 | CPU线程 | GPU线程束 |
|---|---|---|
| 调度单位 | 单个线程 | 32个线程为一组 |
| 上下文切换 | 开销大（微秒级） | 零开销（硬件自动切换） |
| 分支处理 | 每个线程独立 | Warp内串行化分支 |
| 适用场景 | 复杂逻辑、小规模并行 | 大规模数据并行、计算密集型 |
你看这个对比表就明白了。CPU擅长处理复杂的分支逻辑，GPU擅长处理大规模的同质化计算。做量化系统时，我通常把风控逻辑、订单路由这些复杂决策放CPU，把矩阵运算、蒙特卡洛模拟这些计算密集型的扔给GPU。
个人经验：在写Kernel时，我习惯先估算一下每个Block的Thread数。比如一个Block有256个线程，那就是8个Warp。如果SM有64个Warp槽位，那至少需要8个Block才能占满。这个“占满”的概念很重要——Occupancy（占用率）直接影响隐藏访存延迟的能力。
最后说一句：GPU计算不是银弹。我见过不少团队，一上来就把所有逻辑往GPU上搬，结果性能反而更差。关键是要理解数据局部性和计算密度。如果数据传输时间比计算时间还长，那还不如用CPU。
好了，这一章的基础回顾就到这里。记住这些核心概念，后面咱们才能玩转更高级的技巧。