GPU 计算基础

说到 GPU 加速多组学数据预处理，咱们得先把地基打牢。GPU 到底怎么工作的？凭什么比 CPU 快那么多？我个人觉得，理解了底层架构，写出来的代码才算心里有底。这一章咱们就聊聊 GPU 的核心概念：SM、Warp、内存层次，还有 CUDA 编程模型。

1.1 GPU 架构：SM 与 Warp

GPU 不是一个大核，它是一堆小核的集合。每个小核叫 CUDA Core，而组织这些小核的基本单位叫 SM（Streaming Multiprocessor）。我刚开始接触时把 SM 想成了 CPU 的核，其实差别大了去了。

  • SM：一个 SM 包含几十个 CUDA Core、共享内存、寄存器、调度器。GPU 可以有很多个 SM，比如 A100 上有 108 个 SM。
  • Warp：SM 执行的最小单位是 32 个线程组成的 Warp。所有线程执行同一条指令，但处理不同的数据——这就是 SIMT（单指令多线程）。
  • 为什么是 32？ 硬件设计决定的。一条指令发射出去，32 个线程同时跑。如果其中一个线程分支了，其他线程就被迫等待。这就引出了分支发散的问题。
避坑指南：我曾经写过一段代码，if (threadIdx.x % 2 == 0) 做一件事，else 做另一件事。结果同一个 Warp 里既有偶数又有奇数线程，指令被串行执行了，性能直接腰斩。所以，尽量让同一个 Warp 内的线程走同一条路径。

SM 里的 Warp 调度器负责快速切换 Warp。当一个 Warp 等待数据时，调度器立刻换另一个 Warp 上去执行。你不必像 CPU 那样手动做上下文切换，硬件替你干了。说白了，GPU 靠大量 Warp 隐藏延迟。

1.2 内存层次：从全局到共享

GPU 的内存比 CPU 复杂得多。你得搞清楚哪些快哪些慢，不然代码会跑得很慢。我在项目中优化过一个基因组比对算法，就是因为没用好 shared memory，带宽利用率不到 10%。

内存类型位置延迟（周期）容量作用域
Global Memory显存（DRAM）~400-800GB 级所有线程
Shared MemorySM 内~30-50KB 级（通常 48KB）线程块内
RegisterSM 内～0每个线程 255 个单个线程
L1 CacheSM 内~30KB 级SM 内所有线程
Constant Memory显存~10064KB所有线程（只读）

你想想看，Global Memory 访问一次要几百个周期，而 Shared Memory 只要几十个周期。如果一个数据被同一个线程块内的多个线程反复读，你就应该先把它加载到共享内存里。这也是 CUDA 优化的核心手段。

// 一个简单的 shared memory 使用例子
__global__ void shared_mem_sum(const float *in, float *out) {
    __shared__ float sdata[256];
    int tid = threadIdx.x;
    sdata[tid] = in[blockIdx.x * blockDim.x + tid];
    __syncthreads();  // 等所有线程写完后才能读
    for (int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }
    if (tid == 0) out[blockIdx.x] = sdata[0];
}
注意： 共享内存大小是有限的。如果线程块太大或每个线程用太多 shared memory，显卡会启不了内核。我曾经踩过这个坑，一个 1024 线程的块需求 shared memory 超出了 48KB，启动时直接报错。用 cudaOccupancyMaxPotentialBlockSize 可以算合适的尺寸。

1.3 CUDA 编程模型概览

CUDA 把 GPU 看作一个协处理器。你写一个 kernel 函数，然后告诉它启动多少个线程。线程按 Grid 和 Block 组织：

  • Grid 包含多个 Block。
  • Block 包含多个 Thread。
  • Block 内的线程可以通过 Shared Memory 通信，也可以通过 __syncthreads() 同步。
  • 不同 Block 之间没有直接通信手段。

启动一个 kernel 的语法：

// 定义 kernel
__global__ void vec_add(float *a, float *b, float *c) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) c[i] = a[i] + b[i];
}

// 调用 kernel
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c);
我经常被问到：blockDim 和 gridDim 是怎么定的？没有万能公式，但有一条经验：让 block 大小是 32 的倍数（比如 128 或 256），这样 Warp 不会浪费。同时保证每个 SM 有足够多的 block 来隐藏延迟，通常 4-8 个 block 起步。

1.4 GPU vs CPU：为什么 GPU 更适合并行处理

CPU 是低延迟的设计，GPU 是高吞吐的设计。来，看个对比：

特性CPUGPU
核心数量4-64数千（CUDA Core）
控制单元复杂（分支预测、乱序执行）简单，交给 Warp 调度
Cache大（MB 级）小（KB 级），靠显存带宽
内存带宽~50 GB/s~900 GB/s（A100）
适合场景串行、复杂逻辑数据并行、简单计算

你想想看，多组学数据预处理里，很多操作是“对每个元素做同样的事”——比如读取 FASTQ 文件，对每条序列做质量过滤，对齐到参考基因组。这些任务高度并行，GPU 天然适合。但转移数据到显存需要时间，所以 数据搬移的开销必须小于并行加速带来的收益。我在一个 RNA-seq 定量项目中，一开始直接逐条传输，结果 CPU 都跑完了 GPU 还在等数据。后来改成批传输，速度才提上去。

一个小技巧：用 cudaMemcpyAsync 结合流水线，让传输和计算重叠。后面章节会细讲。
GPU 计算基础知识框架 GPU 计算基础 GPU 架构 SM: 流式多处理器 Warp: 32线程/指令 内存层次: Global/Shared/Register CUDA 编程模型 Grid / Block / Thread Kernel 启动 <<<grid, block>>> 内存管理: cudaMalloc / cudaMemcpy GPU vs CPU 差异 并行度: 千核 vs 十核 内存带宽: 高 vs 低 延迟容忍: 隐藏 vs 显式 核心要点 利用 SM 和 Warp 隐藏延迟，善用 Shared Memory 减少 Global Memory 访问 CUDA 编程模型围绕 Grid/Block/Thread 组织，启动内核时注意 Block 尺寸和 Occupancy

这张图把本章的三个核心模块串起来了。后面我们会在实际的多组学数据预处理中，反复用到这些概念：比如用共享内存加速碱基质量分数的统计，用 Warp 规约做并行 count。先把基础扎稳，后面的优化才能有的放矢。

← 上一章 📖 返回目录 下一章 →