GPU 计算基础：GPU 架构简介、CUDA 编程模型、GPU 与 CPU 比较
GPU 架构到底长什么样？
刚开始接触 GPU 时，我总觉得它就是个跑得快的显卡，打游戏用的。直到我第一次用 CUDA 做基因组 k-mer 计数，才发现这东西的设计思路和 CPU 完全两个路子。
GPU 的核心思想是「用大量的小核心干简单活」。一个典型的 NVIDIA GPU 里有几十个 SM（流式多处理器），每个 SM 里面又有几百个 CUDA 核心。这些核心不能单独跑复杂的逻辑控制，但它们可以同时执行上万条线程。
我这里画了个简化图，带你快速过一遍概念：
我在实际优化宏基因组比对程序时，就踩过「线程发散」的坑。说白了，一个 Warp 里的 32 个线程如果走了不同的分支（比如 if-else），GPU 会串行执行，性能直接腰斩。嗯，这里要注意，写 CUDA 代码时要尽量保证同一 Warp 内的线程走相同路径。
CUDA 编程模型：核函数与线程层次
CUDA 把 GPU 当做一个协处理器。你写一个 核函数（用 __global__ 修饰），然后在主机端调用它，指定网格（grid）和块（block）维度。举个例子，给两个数组做逐元素加法：
// 核函数：每个线程处理一个元素
__global__ void vecAdd(float *a, float *b, float *c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// 主机端调用
int n = 1024;
dim3 blockSize(256); // 每个块 256 个线程
dim3 gridSize((n + blockSize.x - 1) / blockSize.x); // 算出需要多少个块
vecAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);调用时，<<<gridSize, blockSize>>> 指定了网格和块的维度。每个线程靠 blockIdx.x 和 threadIdx.x 得到自己的唯一 ID。比如线程块 0 里的线程 0～255，块 1 里的线程 256～511，以此类推。
小经验：我一般把 blockDim.x 设成 128 或 256 的倍数。为什么？因为 Warp 大小是 32，64、128、256 都能被 32 整除，避免资源浪费。另外，块内线程数不要超过 1024——这是硬件上限。
内存层次：全局、共享、寄存器
GPU 的内存模型是分层的。我刚开始做生物信息加速时，不知道有共享内存这回事，结果全局内存访问带宽成了瓶颈。后来我把频繁访问的数据放到共享内存里，速度提升了将近 3 倍。下面这个表帮你快速理清：
| 内存类型 | 作用域 | 延迟 | 容量 | 典型用途 |
|---|---|---|---|---|
| 全局内存 | 所有线程 | ~400-800 cycle | GB 级 | 主数据存储 |
| 共享内存 | 单个块 | ~5-30 cycle | 几十 KB | 块内数据重用、归约 |
| 寄存器 | 单个线程 | 1 cycle | 有限（通常 255 个/线程） | 局部变量 |
| 常量内存 | 所有线程 | ~400 cycle（缓存命中快） | 64 KB | 只读常量数据 |
我曾经犯过的错：把一个大数组全放在共享内存里，结果每个块只能跑很少线程，占用率极低。共享内存是稀缺资源，用之前先算算每个块要多少字节，别超过硬件限制（一般 SM 中 48 KB～96 KB）。
GPU vs CPU：两个世界的比较
CPU 擅长处理复杂逻辑和串行任务，延迟低；GPU 擅长处理大量并行简单任务，吞吐量高。很多初学者问：「GPU 频率比 CPU 低不少，为什么跑某些程序反而快？」
关键就在「并行度」这三个字。你的宏基因组序列比对，每条 reads 的处理相互独立——这就是 GPU 的天然粮仓。
| 对比维度 | CPU | GPU |
|---|---|---|
| 核心数 | 4～64 个大核 | 几百～几千个小核 |
| 控制单元 | 复杂（分支预测、乱序执行） | 简单（SIMT 调度） |
| 缓存 | 大（MB 级多级缓存） | 小（L1/共享内存为主） |
| 内存带宽 | 几十 GB/s | 几百 GB/s～TB/s |
| 单线程性能 | 极高 | 极低 |
| 并行吞吐量 | 限（受限于线程数） | 极高（只要数据够多） |
举个例子：比对 1000 万条 reads，CPU 上开 16 线程可能要半小时，GPU 上开几万个线程也许只要 2 分钟。但如果在 CPU 上执行一个递归二叉树遍历，GPU 就傻了——线程发散、内存访问不连续，效率还不如单核 CPU。
初始化 GPU 环境：一条命令搞定
课程后面我们会实际搭建环境。标准的 CUDA 工具包安装完后，检查一下驱动和版本：
nvidia-smi
nvcc --version如果能看到 GPU 型号、驱动版本、CUDA 版本，就说明环境没问题。我个人习惯把这两行命令加到 shell 启动脚本里，每次登录自动打印，确保别用错 CUDA 版本。
小提示：宏基因组数据量大，建议选显存 8 GB 以上的 GPU。用 nvidia-smi 看显存占用情况，别让程序因为显存不足而崩溃。
好，这一章的核心概念就到这儿。记住三个关键词：SM、Warp、SIMT。理解了它们，你就能明白为什么 GPU 适合宏基因组这种「数据多、逻辑简单」的任务。下一节咱们开始上手写第一个 CUDA 程序——你想想看，直接把序列比对塞上去，是不是很爽？
（别急，先把手里的架构图再看一眼，调好代码里的 blockSize 和 shared memory 用量，后面才不会翻车。）