1. GPU计算基础：GPU架构简介（CUDA核心、显存层次）

各位同学好。开始讲单细胞多组学加速之前，我们得先摸清楚手里的“工具”——GPU到底是个什么东西。我当年刚开始接触GPU计算时，也以为它就是块“能跑图形的卡”，直到第一次在单细胞数据上跑了个聚类，CPU跑了三天三夜没出来，换上GPU一个多小时搞定……嗯，从那以后我彻底改变了看法。

说白了，GPU和CPU的设计哲学完全不同。CPU是“全能冠军”，擅长处理各种复杂逻辑；GPU是“人海战术”，靠成千上万个简单核心并行干活。你的单细胞数据里几万个细胞、几万个基因，正好适合GPU这种“一个核心盯一个细胞”的暴力模式。

1.1 GPU架构核心概念

先看一张总体图，帮你把本章知识点串起来：

GPU加速单细胞多组学分析——计算基础框架 GPU架构简介 CUDA核心 · SM单元 全局内存 · 共享内存 L1/L2缓存 · 寄存器 显存带宽是关键瓶颈 CUDA编程模型 Grid · Block · Thread Kernel函数 · 内存层次 host/device数据转移 同步与流管理 CuPy & Numba CuPy：GPU版NumPy Numba JIT编译加速 @cuda.jit装饰器 CPU/GPU无缝切换 单细胞多组学加速应用 降维 · 聚类 · 差异表达 批次校正 · 细胞类型注释 本章打好GPU基础，后续所有单细胞分析加速案例都建立在这个框架之上

这张图算是本章的“路线图”。我们从GPU硬件架构出发，理解它为什么适合并行计算；然后看CUDA怎么让软件“指挥”硬件；最后用CuPy和Numba两个Python工具，让你不用写太多C代码也能用上GPU。我个人习惯先看整体再抠细节，所以你先有个印象就行。

1.2 CUDA核心与显存层次

先聊硬件。NVIDIA的GPU里塞了几千个CUDA核心，这些核心被分成一组一组的，叫SM（Streaming Multiprocessor）。每个SM内部有自己的共享内存、寄存器、调度器。你想想看，一个GPU里几十个SM，每个SM管一百多个核心，这些核心同时干活——这就是并行能力的来源。

显存层次这块我得多说两句。我遇到过不少同学，写了个GPU程序发现反而比CPU慢，跑个profile一看——数据传输占了90%的时间。为什么会这样？说白了，GPU的显存层次比CPU复杂，你得搞清楚数据该放哪儿。

内存类型 位置 访问速度 容量 作用
全局内存 显存芯片 ~200 GB/s 8~80 GB 存原始数据、计算结果
共享内存 SM内部 ~1.5 TB/s 48~164 KB/SM 线程间数据共享、减少全局访问
寄存器 每个核心 最快 256个/线程 局部变量、临时结果
L1/L2缓存 芯片内部 ~1 TB/s 各几MB 自动缓存全局内存数据

核心要点：单细胞数据动辄几万个细胞×几万个基因，一个浮点矩阵就是几个GB。这些数据必须放在全局内存里。但如果你能在每个SM的共享内存里做局部聚合计算，速度会快一个数量级。

注意：我曾经有个项目，把整个表达矩阵一次性从CPU拷贝到GPU，10GB数据花了15秒。后来改成流式传输，边拷贝边计算，总时间降到3秒。数据传输是GPU计算的隐藏瓶颈，千万留意。

1.3 CUDA编程模型

CUDA的编程模型，其实就三句话：Grid一>Block一>Thread。你写一个函数（叫Kernel），然后告诉CUDA“开多少个线程来跑”。这些线程按Grid和Block组织，每个线程有唯一的ID。

我们看个最简单的例子——向量加法：

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 调用方式：
int threads_per_block = 256;
int blocks = (n + threads_per_block - 1) / threads_per_block;
vec_add<<<blocks, threads_per_block>>>(d_a, d_b, d_c, n);

这里每个线程只处理一个元素。单细胞分析里很多操作都能这样“暴力并行”——比如对每个细胞做标准化、对每个基因算方差。

CUDA的内存层次你得记清楚：host是CPU端，device是GPU端。数据从host传到device用cudaMemcpy，算完再传回来。我习惯的做法是：在分析开始时把所有数据一次性传上去，中间所有步骤都在GPU上完成，最后只把结果传回来。

个人经验：写CUDA Kernel时，我总会在开头加一个边界检查 if (idx < n)。你想想看，如果我们开了1000个线程但数组只有500个元素，不检查的话就会访问非法内存，程序崩溃得莫名其妙。这个习惯帮我省了不少debug时间。

1.4 CuPy与Numba快速入门

好，现在说说怎么在Python里用GPU。我知道大部分做生信的同学C/C++用得少，让直接写CUDA C代码有点劝退。所以CuPy和Numba就是来救场的。

CuPy——说白了就是GPU版的NumPy。接口几乎一模一样：

import cupy as cp
import numpy as np

# 创建数组
a_cpu = np.random.randn(10000, 10000).astype(np.float32)
a_gpu = cp.asarray(a_cpu)   # 传到GPU

# 在GPU上做矩阵乘法（单细胞里经常算细胞-细胞相似度）
b_gpu = cp.random.randn(10000, 500, dtype=cp.float32)
c_gpu = cp.dot(a_gpu, b_gpu)

# 把结果传回CPU
c_cpu = cp.asnumpy(c_gpu)

print(f"GPU矩阵乘法耗时: {cp.cuda.Event.elapsed_time(...):.2f} ms")

代码量几乎不变，但速度提升几十倍。我记得第一次在单细胞数据上跑PCA时，用NumPy算协方差矩阵花了3分钟，换成CuPy只用4秒——那种感觉，啧，真爽。

Numba则更灵活。你写一个普通Python函数，加个 @cuda.jit 装饰器，Numba会把它编译成CUDA Kernel：

from numba import cuda
import numpy as np

@cuda.jit
def normalize_cells(expression, means, stds):
    """对每个细胞的基因表达做Z-score标准化"""
    idx = cuda.grid(1)
    if idx < expression.shape[0]:
        for j in range(expression.shape[1]):
            expression[idx, j] = (expression[idx, j] - means[idx]) / stds[idx]

# 准备数据
expr = np.random.randn(50000, 20000).astype(np.float32)
means = np.mean(expr, axis=1)
stds = np.std(expr, axis=1)

# 配置线程
threads_per_block = 256
blocks = (expr.shape[0] + threads_per_block - 1) // threads_per_block

# 移到GPU
d_expr = cuda.to_device(expr)
d_means = cuda.to_device(means)
d_stds = cuda.to_device(stds)

# 启动Kernel
normalize_cells[blocks, threads_per_block](d_expr, d_means, d_stds)

这个例子就是单细胞标准化里很常见的操作——每个细胞独立做Z-score。Numba让你用Python的语法写并行逻辑，底层自动处理好内存管理和线程映射。

CuPy和Numba怎么选？我个人的经验是：能用CuPy就用CuPy，因为它帮你把矩阵运算优化好了，代码也干净；碰到CuPy没有的、需要自定义并行逻辑的操作，就上Numba。还有，注意数据类型——GPU对float32支持最好，别用float64，速度会差很多。

避坑提醒：我曾经用Numba写了个Kernel，调试了整整一天才发现问题——在GPU函数里用了Python的list。Numba不支持Python原生list，你得用numpy数组或者Numba的cuda.local.array。嗯，这个坑下面几章还会遇到。

1.5 本章小结

这一章我们走马观花看了GPU计算的核心要素：CUDA架构的分层设计、显存访问的成本差异、CUDA编程的Grid-Block-Thread模型，以及CuPy和Numba两个趁手的Python工具。说白了，GPU加速的本质就三件事：把数据搬上去、并行算、把结果搬下来。后面29章，我们会在这个基础上，一步步把单细胞多组学的核心分析流程——从数据加载、归一化、降维、聚类到差异表达——全部用GPU加速一遍。

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