GPU计算基础：CUDA架构概览
做量化因子挖掘，说白了就是跟海量数据死磕。我刚开始用CPU跑多因子回测时，一个中证500的全市场计算就要等半小时。后来接触了CUDA，才真正体会到什么叫「降维打击」。这一章，咱们就聊聊CUDA到底是个什么玩意儿。
CUDA架构：不只是显卡那么简单
很多人以为GPU就是用来打游戏的。其实不然。CUDA是NVIDIA推出的并行计算平台，它把GPU变成了一个「超级计算器」。我习惯把CUDA架构理解成三个层次：
- 流多处理器（SM）：GPU的核心计算单元，每个SM包含多个CUDA核心
- CUDA核心：最基础的执行单元，负责算术运算
- 线程束（Warp）：32个线程为一组，这是调度和执行的基本单位
你想想看，一个高端GPU可能有几十个SM，每个SM又有上百个CUDA核心。这意味着什么？意味着你可以同时跑几千甚至上万个线程。做因子挖掘时，我经常把每只股票的计算任务分配给一个线程，几秒钟就能算完整个A股市场。
核心要点：CUDA架构的精髓在于「大规模并行」。它不是让单个线程跑得更快，而是让成千上万个线程同时干活。
线程层次结构：Grid, Block, Thread
这是CUDA编程里最基础也最容易搞混的概念。我刚开始学的时候，也花了点时间才理清楚。说白了，就是三层嵌套：
- Thread（线程）：最基础的执行单元，每个线程执行相同的核函数
- Block（线程块）：一组线程的集合，共享Shared Memory，可以同步
- Grid（网格）：一组Block的集合，构成整个计算任务
举个例子你就明白了。假设你要计算10000只股票的日收益率：
// 定义核函数
__global__ void calcReturn(float* price, float* ret, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
ret[idx] = (price[idx] - price[idx-1]) / price[idx-1];
}
}
// 调用核函数
int n = 10000;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
calcReturn<<<blocksPerGrid, threadsPerBlock>>>(d_price, d_ret, n);这里每个线程处理一只股票的计算。256个线程组成一个Block，Block不够了就多开几个Grid。嗯，这里要注意：Block的数量不是随便定的，它受GPU硬件限制。
我的经验：Block大小选128或256通常性能最好。太小了SM利用率低，太大了寄存器压力大。我曾经试过512，结果性能反而下降了。
内存层次结构：Global, Shared, Local, Register
GPU的内存体系比CPU复杂得多。我画了一张图，帮你理清这四层的关系：
这四层内存各有各的脾气。我挑重点说：
寄存器（Register）
最快，但数量有限。每个线程最多能用255个寄存器。做因子计算时，我习惯把频繁使用的中间变量放在寄存器里。比如计算移动平均时，把累加器声明为局部变量，编译器会自动分配到寄存器。
避坑指南：我曾经写过一个核函数，用了太多局部变量，导致寄存器溢出。结果性能从200μs掉到了800μs。后来用--maxrregcount=32限制寄存器数量，反而性能提升了。因为编译器被迫使用Shared Memory，减少了寄存器压力。
共享内存（Shared Memory）
这是CUDA编程的「杀手锏」。同一Block内的线程可以共享数据，延迟只有全局内存的十分之一。做因子挖掘时，我经常用它来缓存股票的历史数据。
__global__ void calcMA(float* price, float* ma, int n, int window) {
__shared__ float s_price[256]; // 共享内存缓存
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
s_price[threadIdx.x] = price[idx];
__syncthreads(); // 同步，确保所有线程都加载完毕
// 计算移动平均
float sum = 0;
for (int i = 0; i < window; i++) {
sum += s_price[threadIdx.x + i];
}
ma[idx] = sum / window;
}
}这里用__syncthreads()同步很重要。不然后面线程读到的数据可能还没写进去。我刚开始做时吃过这个亏，算出来的移动平均全是错的。
全局内存（Global Memory）
容量最大，但速度最慢。所有线程都能访问。做因子挖掘时，原始数据、计算结果都放在这里。优化全局内存访问的关键是「合并访问」——让相邻线程访问相邻地址。
优化技巧：我习惯把数据按「列优先」方式存储。比如10000只股票×100天的价格数据，按天存储（每列是一只股票），这样相邻线程访问相邻地址，带宽利用率能到80%以上。
局部内存（Local Memory）
这个其实是个「伪概念」。它物理上在全局内存里，只是逻辑上每个线程私有。当寄存器不够用时，编译器会自动把变量溢出到局部内存。性能很差，尽量避免。
实战：因子计算中的内存选择
说了这么多理论，咱们来点实际的。假设你要计算一个简单的动量因子：过去20天收益率的标准差。你会怎么分配内存？
| 数据 | 内存类型 | 原因 |
|---|---|---|
| 原始价格数据 | Global Memory | 数据量大，所有线程都需要 |
| 每只股票的20天价格 | Shared Memory | 同一Block内线程共享，减少全局访问 |
| 累加器、计数器 | Register | 频繁使用，需要最快速度 |
| 临时中间结果 | Local Memory（尽量避免） | 寄存器不够时才会用 |
你看，选对内存类型，性能差距能到10倍以上。我刚开始做因子挖掘时，一股脑全用全局内存，结果一个简单的因子要算5分钟。后来优化了内存使用，同样的计算30秒搞定。
总结一下：CUDA编程的核心就是「数据在哪里，计算就在哪里」。尽量把数据往寄存器或共享内存里放，减少全局内存访问。这个原则，贯穿整个因子挖掘GPU并行化的始终。
好了，这一章的内容就到这里。记住我说的：寄存器最快但最少，共享内存是宝，全局内存要省着用。下一章咱们聊聊具体的并行化策略，到时候我会分享更多实战中的坑和技巧。
公众号:蓝海资料掘金营，微信deep3321