GPU计算基础：CUDA架构、线程层级与内存模型
这一章咱们聊点实的。你学BLAST加速，如果不先把CUDA的线程模型和内存体系搞明白，后面调优就像蒙着眼开车。我当年第一次写内核，线程块大小随便填了个128，结果跑出来比CPU还慢…后来才意识到线程组织方式和内存访问模式才是性能的命门。
从一张图开始：CUDA的层次化执行模型
先把整体框架装进脑子里。下面这张图展示了从Grid→Block→Thread的层级，以及它们怎么映射到硬件上的SM（流多处理器）。你看一眼就能记住。
线程层级：从Grid到Thread
CUDA把线程组织成三层结构。你写内核的时候，通过<<指定每一层的规模。
| 层级 | 逻辑含义 | 硬件映射 | 最大规模（典型） |
|---|---|---|---|
| Grid | 所有线程的集合 | 对应整个GPU设备 | 最多 2³¹-1 个Block |
| Block | 一组协同线程 | 分配到同一个SM | 最多 1024 个线程 |
| Thread | 最小执行单元 | 一个CUDA Core执行 | 任意（受硬件限制） |
| Warp | 32个线程的硬分组 | SM以warp为单位调度 | 固定32 |
我的经验：块大小最好选32的倍数。比如128、256。为什么？因为warp是32线程一组，块如果不是32的倍数，最后一个warp就填不满，浪费计算单元。我早期犯过这个错，设了个100线程的块，结果最后一个warp只有4个活跃线程，利用率掉了30%+。
一个典型的内核启动长这样：
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
myKernel<<>>(d_input, d_output, n); 这里面threadsPerBlock是块内线程数，blocksPerGrid是块的总数。运行时，硬件会把block分配到有空闲SM的单位上。如果一个block需要的资源（寄存器、共享内存）太多，SM能同时容纳的block数就会减少，这直接影响占用率。
内存模型：速度与容量的博弈
GPU的内存不是一锅粥。从快到慢、从小到大有五层。调优的核心思路：把频繁访问的数据尽可能往快存里放。
| 内存类型 | 可见范围 | 容量（典型） | 延迟（周期） |
|---|---|---|---|
| 寄存器 | 单个线程 | 每个SM 65536个（32位） | 1 |
| 共享内存 | 单个Block | 每SM 48KB ~ 164KB | ~30 |
| L1 缓存 | 单个SM | 与共享内存共享容量 | ~30 |
| L2 缓存 | 整个芯片 | 几MB | ~200 |
| 全局内存 | 所有线程 | 几GB到几十GB | 400~800 |
| 本地内存 | 单个线程（寄存器溢出） | 自动，放在全局内存中 | 同全局内存 |
记住一个原则：全局内存是硬盘，共享内存是内存条，寄存器是CPU的L1缓存——数据越靠近计算单元，越快。代价是空间小、生命周期短。
拿BLAST的序列比对来说，一条查询序列通常只有几百个碱基，而参考数据库可能有几百GB。你不能把整个库搬进共享内存，但可以把当前要比较的那一小段先塞进去，让一个block里的线程共同使用。这就是空间局部性的经典应用。
我曾经踩过的坑：共享内存虽然快，但有bank conflict。共享内存内部被分成32个bank（每个bank 4字节宽）。如果同一个warp里多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。解决方法：加padding或者调整数据布局。比如二维数组每行多加一个元素，让相邻线程错开bank。
// 错误：二维数组每行宽为N，相邻线程访问同一bank
__shared__ float s_data[32][32];
s_data[threadIdx.x][threadIdx.y] = ...; // 冲突！
// 正确：每行加1 padding
__shared__ float s_data[32][33];
s_data[threadIdx.x][threadIdx.y] = ...; // 无冲突怎么选？经验法则
这里没有银弹，但有几个常用套路可以抄作业：
- 块大小：选 128、256 或 512（保持32倍数）。太小（<64）会让SM里block太少，隐藏延迟能力不足；太大（>512）容易超寄存器限制，导致kernel启动失败。
- 共享内存分配：先算一算你的内核用了多少寄存器，再定共享内存。你可以用
--ptxas-options=-v编译，看寄存器用量和共享内存用量。 - 全局内存访问：尽量让warp内的32个线程访问连续的地址（合并访问）。比如一个线程读一个float，32个线程正好读128字节，一次内存事务搞定。要是跳跃式访问，就得发好几次事务，带宽直接腰斩。
我个人习惯：写新kernel之前，先拿cudaOccupancyMaxPotentialBlockSize算一下最佳block size。这个函数会考虑寄存器和共享内存的约束，给出理论最高占用率。很多时候，直觉选的256并不是最优的。
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);
printf("推荐 block size: %d\n", blockSize);你想想看，如果你能把BLAST里最热的那个循环——得分矩阵的计算——完全放到共享内存里，再配上合适的block大小，那加速比轻松翻倍不是梦。后面的章节我们会用实际数据来验证这一点。