4、GPU架构基础：CUDA核心、SM（流多处理器）、内存层级
搞GPU加速BLAST，绕不开CUDA架构。说白了，你得知道手里的显卡到底长什么样、怎么干活、数据怎么流动。我刚开始接触GPU时，总把它当成「更强一点的CPU」，结果性能一塌糊涂。后来才明白，你得按GPU的脾气来写程序。
这一节，我们就扎进NVIDIA GPU的肚子里，看看CUDA核心、SM（流多处理器）和内存层级到底是怎么回事。
4.1 从「核心」说起
CPU的核心，一个顶一个，频率高、缓存大、擅长复杂控制。GPU的核心呢？数量多、频率低、擅长简单运算。NVIDIA管它们叫CUDA核心。
一个典型的CUDA核心能做什么？整数运算、单精度浮点、双精度浮点（部分型号）。有些新一代核心还负责Tensor Core（矩阵乘）和RT Core（光线追踪），但对我们做BLAST加速来说，最常用的是单精度浮点和整数运算。
我习惯把CUDA核心比作「流水线上的工人」。工人多了，并行度就上去了。但工人不能独立干活——他们需要被组织起来。
| 概念 | CPU核心 | CUDA核心 |
|---|---|---|
| 数量 | 4~64 | 1024~16384 |
| 频率 | 3~5 GHz | 1~2 GHz |
| 擅长的活 | 分支预测、乱序执行 | 并行算术、吞吐优先 |
| 功耗 | 单核几十瓦 | 单核不到一瓦 |
你想想看，BLAST里的序列比对，本质上是大量独立的打分计算。每个CUDA核心算一小段，合起来就是整个数据库的搜索。这也正是GPU的优势所在。
4.2 SM（流多处理器）—— 真正的执行单元
CUDA核心不是直接暴露给我们的。实际的组织单位是流多处理器（SM）。每个SM里包含若干组CUDA核心、共享内存、寄存器文件、调度器、加载/存储单元等。
我记得有个项目，我写了一个简单的向量加法kernel，在计算能力7.5的GPU上跑，占用了32个线程束（warp）。但性能始终提不上去。后来看nvidia-smi的占用率，才发现SM的占用率只有30%。
为什么？因为我没考虑到SM的资源限制。每个SM能同时驻留的线程块（block）数量、线程总数、共享内存总量都是有限的。你写得再漂亮，资源不够，占不满SM，性能就上不去。
| SM部件 | 作用 | 对BLAST的影响 |
|---|---|---|
| Warp调度器 | 每时钟周期选择可执行的warp | 隐藏访存延迟 |
| CUDA核心 | 执行算术指令 | 打分函数计算 |
| Special Function Unit | sin/cos/rsqrt等 | 不太常用 |
| 共享内存 | 低延迟片上缓存（可编程） | 存查询序列、局部打分矩阵 |
| 寄存器文件 | 每个线程私有的最快存储 | 存放循环变量、局部分数 |
| 加载/存储单元 | 访存全局/共享/常量/纹理内存 | 数据搬运瓶颈 |
一个关键指标：每个SM能同时处理的warp数量决定了指令级并行度。我曾经踩过一个坑：盲目增大线程块大小，结果每个线程用的寄存器太多，SM能驻留的线程块数下降，整体性能反而不如小线程块。嗯，这里要注意，不是越大越好。
个人小窍门：调优时先用--ptxas-options=-v编译，看每个kernel用了多少寄存器和共享内存。然后根据SM的资源上限，反推最优的block/thread配置。
4.3 内存层级 —— 得知道数据放哪儿
GPU的内存层级，就像城市的物流体系。寄存器和共享内存是「门口小卖部」，拿东西快但是容量小。全局内存是「郊区大仓库」，装得多但来回一趟几十个时钟周期。图像、常量、纹理内存各有各的用途。
我习惯把BLAST的流程映射到内存层级上：
- 寄存器：每个线程临时保存当前比对的得分、指针等。
- 共享内存：存放查询序列（query）的碱基/氨基酸、局部打分矩阵（如BLOSUM62的某一片段）。
- 全局内存：存放整个数据库序列、完整打分矩阵、结果输出。
- 常量内存：存放gap penalty、score matrix（如果全局查不频繁的话）。
举个实际数字：在RTX 3080上，每个SM有128KB共享内存，64个warp能驻留的最大线程数是1536。如果你把查询序列放在共享内存里，每个比对线程块能共享这份数据，省掉反复从全局内存读的代价。我曾经这样优化过BLASTP的本地化版本，提升了40%的吞吐。
避坑指南：共享内存不是无限用的。我之前以为「放点数据总比放全局快」，结果每个线程块用了96KB共享内存，SM只能同时驻留一个线程块，并行度大跌。算下来，全局内存虽然慢，但并行度高了反而更快。记住，共享内存是给「频繁复用但小量数据」用的。
寄存器也得小心。每个线程默认最多255个32位寄存器。但如果你写了一个很复杂的打分循环，编译器可能会给你分配大量寄存器，导致每个SM能驻留的线程数减少。我曾在GTA优化的kernel里发现单个线程用到64个寄存器，只好拆分成多个小kernel来保存中间结果。
4.4 一个直观的框架图
下面这张图把上面说的东西串起来。你可以看到，从GPU芯片 → SM → CUDA核心 → 内存层级的层次关系，以及每个层次在BLAST加速中的对应角色。
核心要点：GPU加速BLAST的关键在于三条：
1. 把数据库切分成小任务分配给大量CUDA核心；
2. 通过共享内存减少全局内存访问次数；
3. 每个线程的寄存器要省着用，别让SM驻留的线程束变少。
4.5 归纳一下，我常用的几个检查清单
- 查看GPU计算能力：
cudaGetDeviceProperties，记下SM数量、每个SM的共享内存大小、寄存器数量。 - 估算BLAST任务的并行度：数据库序列数 × 每个序列的分段数。确保远大于SM数量 × 每个SM的warp上限。
- 共享内存分配：把查询序列（通常几百个残基）塞进共享内存，一个block共享一份。如果序列太长，分多次迭代。
- 寄存器限制：用
__launch_bounds__或者-maxrregcount控制每个线程的寄存器数目，避免资源过用导致occupancy下降。 - 访问模式：全局内存尽量连续对齐（coalesced access），否则带宽利用率暴跌。
我曾在一个医疗基因比对项目里，把这三个层级的资源分配算得很细。最终，一个BLASTN的GPU实现在双路RTX 4090上跑出了比CPU基准快78倍的成绩。那些日子的调优过程，说起来就是反复在SM资源和内存访问之间做取舍。
好了，这就是GPU架构的基础三件套。你掌握了CUDA核心、SM和内存层级，后面讲BLAST kernel的具体实现时，就不会一头雾水了。