GPU架构与CUDA编程模型：从硬件到编程的完整链路
做量化交易这些年，我越来越觉得GPU是个好东西。尤其是做高频因子挖掘、蒙特卡洛模拟这些计算密集型任务时，CPU那点并行能力根本不够看。今天我们就来聊聊GPU的硬件架构和CUDA编程模型——说白了，就是搞清楚GPU到底是怎么工作的，以及我们怎么指挥它干活。
一、GPU硬件架构：三个核心概念
先看硬件。GPU不是一个大核，而是由很多小核组成的。这些小核的组织方式，决定了我们写代码的方式。
1. SM（Streaming Multiprocessor，流式多处理器）
SM是GPU的核心计算单元。你可以把它想象成一个“小工厂”，里面有很多计算核心、共享内存、寄存器等资源。一个GPU通常有几十到上百个SM。
我在项目中遇到过一个问题：一开始以为GPU核数越多越好，结果发现如果SM的资源（比如寄存器、共享内存）被某个线程块占满了，其他线程块就得排队等着。嗯，这里要注意——SM的资源是有限的，不是线程块越多越好。
关键点：SM是硬件调度的基本单位。一个SM可以同时运行多个线程块，但具体能运行多少个，取决于每个线程块占用的资源量。
2. Warp（线程束）
Warp是GPU执行指令的最小单位。一个Warp包含32个线程，它们在同一时刻执行同一条指令（SIMT模式）。
你想想看，如果这32个线程走的分支不一样（比如if-else），那就会发生“分支发散”——部分线程执行，部分线程等待，性能直接打折扣。我刚开始写CUDA时没注意这个，结果一个简单的条件判断让性能掉了30%。
避坑指南：我曾经因为没处理好Warp内的分支发散，导致回测速度比预期慢了一倍。后来把条件判断改成位运算，性能才恢复正常。记住：尽量让Warp内的32个线程走同一个分支。
3. Memory Hierarchy（存储层次）
GPU的存储系统分好几层，速度差异巨大。我画了个图，你看一眼就明白了。
这张图从上到下，速度越来越慢，容量越来越大。我个人习惯把最常用的数据放在寄存器或共享内存里，全局内存能少用就少用。量化交易中，如果你要处理海量的tick数据，一定要考虑数据在存储层次中的摆放——否则GPU再快，也被内存带宽卡死了。
二、CUDA编程模型：Grid、Block、Thread
硬件讲完了，我们看看怎么编程。CUDA把计算任务分成三层：Grid（网格）、Block（线程块）、Thread（线程）。
1. Thread（线程）
线程是最小的执行单元。每个线程执行一个kernel函数实例。在量化交易中，我经常让每个线程处理一个股票或一个时间窗口的数据。
2. Block（线程块）
Block是一组线程的集合。同一个Block内的线程可以共享数据（通过共享内存），也可以同步（通过__syncthreads()）。
我记得有一次做因子计算，需要同一个Block内的线程协作计算均值。如果不用共享内存，每个线程都得去全局内存读数据，慢得离谱。用了共享内存后，速度提升了5倍。
3. Grid（网格）
Grid是所有Block的集合。一个Grid对应一个kernel函数的执行。你可以把Grid理解成“整个任务”，Block是“子任务”，Thread是“具体干活的人”。
核心关系：Grid → Block → Thread，层层包含。每个Thread有唯一的ID（threadIdx），每个Block有唯一的ID（blockIdx），通过这两个ID可以定位到具体的数据。
三、Host与Device：CPU和GPU的分工
Host是CPU及其内存，Device是GPU及其显存。它们之间通过PCIe总线通信。
我刚开始做CUDA时，犯过一个低级错误：在kernel里直接访问Host内存。结果程序直接崩溃。后来才明白，Host和Device的内存是物理隔离的，必须显式地拷贝数据。
典型的数据流
- 分配内存：在Host和Device上分别分配内存
- 拷贝数据：从Host拷贝到Device（cudaMemcpyHostToDevice）
- 执行kernel：在Device上运行计算
- 拷贝结果：从Device拷贝回Host（cudaMemcpyDeviceToHost）
- 释放内存：清理资源
这个流程看起来简单，但实际项目中坑很多。比如，如果你频繁地在Host和Device之间拷贝数据，PCIe带宽会成为瓶颈。我建议尽量把计算留在Device上，减少来回拷贝。
个人经验：在量化回测中，我会把一整天的tick数据一次性拷贝到Device，然后在GPU上完成所有因子计算和信号生成，最后只把结果拷贝回来。这样PCIe的利用率最高。
四、一个简单的CUDA代码示例
光说不练假把式。我们写一个向量加法的例子，感受一下CUDA的编程风格。
// kernel函数：在Device上执行
__global__ void vecAdd(float *A, float *B, float *C, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}
int main() {
int N = 1024;
size_t size = N * sizeof(float);
// 1. 分配Host内存
float *h_A = (float*)malloc(size);
float *h_B = (float*)malloc(size);
float *h_C = (float*)malloc(size);
// 初始化数据...
// 2. 分配Device内存
float *d_A, *d_B, *d_C;
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);
// 3. 拷贝数据到Device
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// 4. 配置Grid和Block
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 5. 执行kernel
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// 6. 拷贝结果回Host
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// 7. 释放内存
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);
return 0;
}这段代码里，blockIdx.x * blockDim.x + threadIdx.x 是CUDA里最经典的线程索引计算方式。我个人习惯把blockDim设为256或512，这个值通常能充分利用SM的资源。
五、量化交易中的实际应用
说了这么多，到底怎么用？我举两个例子。
1. 并行因子计算
假设你有1000只股票，每只股票有5000根K线。你想计算每只股票的均线、波动率等因子。用CPU的话，你得循环1000次；用GPU的话，你可以让每个Block处理一只股票，每个Thread处理一根K线。这样一次kernel调用就搞定了。
2. 蒙特卡洛模拟
做期权定价时，经常需要做上万次路径模拟。每个路径之间是独立的，天然适合GPU并行。我做过一个项目，用GPU把100万次模拟从10分钟压缩到了10秒——嗯，这就是GPU的魅力。
注意：不是所有计算都适合GPU。如果计算量小、数据依赖强、或者频繁需要Host-Device通信，CPU反而更快。我建议先做profiling，再决定是否上GPU。
六、总结
这一章我们聊了GPU的硬件架构（SM、Warp、存储层次）和CUDA编程模型（Grid、Block、Thread），以及Host和Device的分工。说白了，就是搞清楚GPU怎么组织计算、怎么管理数据。这些是后续所有CUDA优化的基础。
我个人觉得，理解Warp和存储层次是最关键的——很多性能问题都出在这两个地方。下一章我们会深入讨论内存优化，到时候你会更深刻地体会到这一点。
一句话记住：GPU编程 = 理解硬件 + 合理组织线程 + 高效管理数据。三者缺一不可。