📚 课程目录第3章 GPU计算基础

3、GPU计算基础：CUDA架构与编程

好，咱们正式开始碰GPU计算的核心——CUDA。很多做生信的朋友，第一次接触GPU时都有个错觉：把这玩意儿当CPU用，塞进去一堆数据等着跑快就完事了。结果呢？要么跑不起来，要么比CPU还慢。我当年也踩过这个坑，折腾了整整一周才发现问题出在哪儿。

CUDA，说白了就是NVIDIA给自家GPU开的“后门”。它允许你用C/C++直接操作GPU的硬件资源，而不是只把GPU当成一个黑盒显示器。对于基因组注释这种计算密集型任务，搞懂CUDA，就等于拿到了性能调优的钥匙。

3.1 一张图看懂CUDA全景

在讲具体细节之前，我习惯先给全局。下面的图展示了CUDA编程的核心脉络，我建议你盯着它看30秒，再往下读。

这张图分了三个层次：硬件架构、线程模型、内存层次。你看，从GPU设备一路往下，到SM（流处理器簇），再到CUDA核心，硬件是物理基础。中间那层是咱们编程最常打交道的部分——网格、块、线程束、线程。最底下是内存，不同的存储区域，速度和容量天差地别。我每次调优新模型，都会对着这个图来回确认问题出在哪一层。

3.2 CUDA线程模型：从Grid到Thread

聊CUDA编程，首先得搞明白它的线程组织方式。CPU上你开几个线程，系统就给你分几个核。GPU不一样，它会生成成千上万个轻量级线程，然后靠硬件调度器批量喂给计算单元。

为什么能这么做？因为GPU的线程切换代价极低，本质就是寄存器级别的上下文切换。CPU切一次线程可能要几微秒，GPU呢？你想想看，单周期就能切走。

层级结构是这样的：

Grid：整个任务，对应一个kernel函数的一次调用。
Block：Grid里的一个工作组，会被调度到某一个SM上执行。
Warp：32个线程绑在一起，是硬件调度和指令执行的最小单位。
Thread：单个执行流，拥有自己的寄存器和PC指针。

关键思路：编程时你只管声明“我要启动多少个线程”，硬件调度器会自动把它们分组喂给SM。你不需要关心哪个物理核执行哪个线程，这层抽象正是CUDA的核心理念——透明的大规模并行。

3.3 内存层次：你花了很多时间等数据

我在做基因组比对加速时遇到过一个大坑：算法在理论上是完美的O(n)复杂度，但跑在GPU上就是快不起来。后来我查了性能分析报告，发现有一半的周期花在了Global Memory访问上。说白了，计算不是瓶颈，搬数据才是。

CUDA内存的类型和差异

内存类型	位置	访问速度	容量	作用域
Global Memory	显存 (VRAM)	~200-800 GB/s (高带宽)	大 (4~80 GB)	所有线程 + 主机可访问
Shared Memory	SM内 (片上)	~TB/s级别 (极快)	极小 (48~164 KB/Block)	同Block内所有线程共享
Registers	SM内 (片上)	最快 (单周期)	有限 (一般几千/线程)	单个线程私有
Constant Memory	显存 (有L1缓存)	快 (只读缓存命中时)	64 KB	所有线程只读
Texture Memory	显存 (有专用缓存)	快 (对空间局部性好)	有限	所有线程只读

我的个人经验：做序列比对或神经网络推理时，尽量把频繁访问的参考数据塞进Constant Memory或Texture Memory。如果你发现Global Memory读占用了80%以上的时间，那就说明共享内存或常量缓存这块你还没用好。别急着优化计算逻辑，先优化数据搬运的路径。

3.4 第一行CUDA代码：向量加法

咱们不用那些花里胡哨的例子，就从最朴素的向量加法开始。这玩意儿在基因组注释里很常见——比如你要对成千上万个碱基的得分做批量加法。

先看内核函数（kernel）：

// 向量加法 kernel
__global__ void vec_add_kernel(float *d_a, float *d_b, float *d_c, int n)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        d_c[idx] = d_a[idx] + d_b[idx];
    }
}

几个关键点解释一下：

__global__：告诉编译器和运行时，这个函数跑在GPU上，从CPU端调用。
blockIdx.x：当前Block在Grid中的索引。
blockDim.x：每个Block有多少线程。
threadIdx.x：当前线程在Block内的索引。
那个if (idx < n)是边界检查，防止多余线程踩到非法内存。

然后是在CPU端如何调用和传参：

#include <iostream>
#include <cuda_runtime.h>

int main() {
    int n = 1024 * 1024; // 100万个元素
    size_t bytes = n * sizeof(float);

    // 分配主机内存
    float *h_a = new float[n];
    float *h_b = new float[n];
    float *h_c = new float[n];

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f * i;
        h_b[i] = 2.0f * i;
    }

    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 拷贝数据从主机到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 设置网格和块大小
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;

    // 启动 kernel
    vec_add_kernel<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);

    // 等待 GPU 完成
    cudaDeviceSynchronize();

    // 拷贝结果回主机
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 检查
    for (int i = 0; i < 10; i++) {
        std::cout << "c[" << i << "] = " << h_c[i] << std::endl;
    }

    // 清理
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    delete[] h_a; delete[] h_b; delete[] h_c;

    return 0;
}

这段代码你直接复制到支持CUDA的开发环境里就能跑。注意那个cudaDeviceSynchronize()——我曾经有一次忘了加，结果在A100上调试了三个小时才发现CPU端的cudaMemcpy还在排队，而GPU已经算完了。加上它，保证主机和设备是同步的。

3.5 避坑指南：我踩过的三个雷

讲到这儿，我想分享几个实战中的教训。你想想看，这些坑如果没人告诉你，自己踩一遍可能要浪费好几天。

雷区一：线程块尺寸选择不当

我曾经为了图省事，把Block大小设成1024，结果在老的Tesla卡上执行效率极低。原因是每个SM能同时容纳的Warp数量是有限的。Block太大，资源竞争反而厉害。后来我形成了习惯：通常128-256线程/Block比较稳妥，然后通过Occupancy API再细调。

雷区二：Shared Memory踩脏

做完共线性打分时，我让同一个Block的不同线程往Shared Memory里写数据，没加__syncthreads()。结果每次跑出来的结果都不一样。这问题特别隐蔽——你以为代码没问题，但数据老是漂。记住：同一个Block内如果有线程写Shared Memory后，其他线程要读，必须加__syncthreads()障眼。

雷区三：if分支过度造成Warp串行化

CUDA以32个线程为一股执行指令。如果这32个线程走上了不同的if分支，所有分支会串行执行，性能瞬间暴跌。最惨的一次，我写了个复杂的序列比对打分函数，分支多了之后，性能跟单线程差不离。解决方案：尽量让数据对齐到Warp边界，使用三元运算符替换复杂if-else。

3.6 编译环境搭建

最后提一下怎么把代码变成可执行程序。如果你是新手，我建议直接用CMake管理。

一个最简的CUDA工程CMakeLists.txt：

cmake_minimum_required(VERSION 3.18)
project(gpu_genome_annot LANGUAGES CXX CUDA)

set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_ARCHITECTURES "80;89") # A100 和 H100

add_executable(vec_add vec_add.cu)

编译命令很简单：

mkdir build && cd build
cmake ..
make -j$(nproc)

几点说明：

CMAKE_CUDA_ARCHITECTURES写的是80和89，对应安培架构和Hopper架构。如果你用RTX 3090，改成86；T4是75。搞错了，程序可能编译通过但运行时崩溃。
如果你用命令行直接编译，用nvcc vec_add.cu -o vec_add就行，但复杂项目还是交给CMake省心。

嗯，这一节的内容就先到这儿。CUDA的基础概念、线程模型、内存层次、你会的第一段代码、还有几个我踩过的坑……这些框架一旦搭好，后面我们讲如何用GPU加速基因组注释模型的训练时，就能直接往这套体系里填东西了。