GPU计算基础：CUDA架构与内存层次
各位同学好，我是老张。今天咱们聊聊GPU计算最核心的东西——CUDA架构和它的内存体系。说实话，我刚开始接触GPU编程时，也被这些概念搞得晕头转向。但后来我发现，只要理解了这几个关键点，后面写代码就会顺畅很多。
1. CUDA架构：GPU的大脑是怎么工作的？
CUDA是NVIDIA推出的并行计算平台。说白了，它让GPU不仅能做图形渲染，还能干通用计算的活。我最早用CUDA做量化回测时，第一反应是「这玩意儿能比CPU快多少？」结果一跑，快了将近50倍。嗯，当时我就知道，这条路走对了。
CUDA架构的核心是流多处理器（SM）。你可以把SM想象成一个微型计算中心。每个SM里有很多CUDA核心，它们能同时执行大量线程。我习惯把SM比作一个工厂车间，CUDA核心就是车间里的工人。工人越多，并行能力越强。
关键点：CUDA程序是运行在GPU上的，但控制逻辑还是在CPU上。CPU负责发号施令，GPU负责埋头苦干。这种模式叫「异构计算」。
我在项目中遇到过一个问题：一开始把大量数据搬运到GPU，结果发现搬运时间比计算时间还长。后来才明白，GPU擅长的是「计算密集型」任务，而不是「数据搬运密集型」任务。这个坑，我踩过。
2. GPU内存层次结构：数据住哪儿？
GPU的内存不是铁板一块。它分了好几个层次，每个层次的速度、容量、用途都不一样。我画了一张图，帮你理清这个结构。
这张图你看明白了吗？从上到下，速度越来越快，容量越来越小。我刚开始写CUDA代码时，总喜欢把所有数据都扔到全局内存里，结果性能惨不忍睹。后来才学会把热点数据搬到共享内存里，速度一下就上来了。
2.1 全局内存
全局内存是GPU里最大的存储空间，所有线程都能读写。但它的延迟很高，大概几百个时钟周期。我在做量化回测时，经常把历史行情数据放在全局内存里。但要注意，频繁访问全局内存会严重拖慢速度。
避坑指南：我曾经在全局内存上做过一个愚蠢的操作——在循环里反复读取同一个变量。后来发现，把这个变量缓存到寄存器里，速度提升了3倍。记住：全局内存不是用来频繁读写的。
2.2 共享内存
共享内存是每个线程块私有的存储空间。它的速度比全局内存快一个数量级，但容量有限，通常只有几十KB。我习惯把共享内存当作「线程块内的缓存」来用。
举个例子，在计算股票收益率时，如果多个线程需要访问同一段历史数据，我会先把数据加载到共享内存里，然后让线程们从共享内存里读。这样比每个线程都去全局内存里读快得多。
2.3 寄存器
寄存器是速度最快的存储，每个线程独享。但数量非常有限，每个线程通常只有几十个。我写代码时，会尽量把循环变量、临时计算结果放在寄存器里。
小技巧：如果你发现某个变量在循环里被反复使用，不妨用register关键字提示编译器把它放到寄存器里。虽然编译器不一定听你的，但至少是个好习惯。
3. 线程与块：并行计算的基本单位
CUDA的并行模型是「线程-块-网格」三层结构。我刚开始学的时候，觉得这玩意儿太抽象了。后来我把它想象成军队的编制：
- 线程：单个士兵，执行具体任务
- 线程块：一个班，共享资源（共享内存），协同作战
- 网格：整个连队，由多个班组成
你想想看，如果让一个将军直接指挥每个士兵，那肯定乱套。但通过班长（线程块）来管理，就清晰多了。CUDA也是这个道理。
3.1 线程的组织方式
在CUDA里，我们用blockIdx和threadIdx来定位每个线程。比如，一个线程块里有256个线程，那threadIdx就是从0到255。我写代码时，经常这样分配任务：
// 每个线程处理一个数据点
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
output[idx] = input[idx] * 2.0f;
}这段代码很经典。它把N个数据点均匀分配给所有线程。每个线程只处理自己负责的那一个点。这就是所谓的「数据并行」。
3.2 线程块的大小怎么选？
这个问题我经常被问到。说实话，没有标准答案。但有几个经验法则：
- 线程块大小最好是32的倍数（因为warp是32个线程一组）
- 常用值：128、256、512
- 不要太大，也不要太小。太大浪费资源，太小无法充分利用SM
我的经验：在做量化回测时，我通常用256个线程一个块。这个大小在大多数GPU上都能获得不错的性能。当然，具体还是要调优。
3.3 Warp：真正的执行单元
这里有个重要概念——warp。warp是32个线程组成的一组，它们是真正同时执行的。我习惯把warp比作「一列并行执行的队伍」。如果warp里的线程执行了不同的分支（比如if-else），那性能就会下降。这叫「warp divergence」。
我在项目中遇到过这个问题：一个简单的if判断，让warp里的线程走了不同的分支，结果性能直接腰斩。后来我把代码重构，让同一个warp里的线程走相同的分支，速度就回来了。
4. 总结一下
好了，今天的内容就这些。我帮你理一下重点：
| 概念 | 要点 | 我的建议 |
|---|---|---|
| CUDA架构 | SM是核心，CUDA核心是工人 | 理解SM的并行能力 |
| 全局内存 | 容量大，速度慢 | 少用，尽量缓存 |
| 共享内存 | 容量小，速度快 | 多用，线程块内共享 |
| 寄存器 | 最快，线程独享 | 优先使用 |
| 线程与块 | 线程是士兵，块是班 | 块大小选256左右 |
| Warp | 32线程一组，真正并行 | 避免分支发散 |
这些概念是GPU编程的基石。你如果能把它们吃透，后面写优化代码就会得心应手。我记得自己刚开始学的时候，光是理解内存层次就花了一周。但一旦搞懂了，后面的路就顺了。
嗯，今天就到这儿。下一章咱们会深入聊聊如何用CUDA写第一个量化回测程序。到时候我会手把手带你写代码，把今天学的知识用起来。