第2章 GPU计算基础

GPU计算基础：CUDA编程模型入门、GPU内存层次结构、线程块与网格概念

各位同学好，我是你们的老朋友。今天咱们来聊聊GPU计算最核心的几个概念。说实话，我当年刚接触CUDA时，也被线程块、网格这些名词搞得晕头转向。但别担心，我会用最接地气的方式，把这些硬骨头啃下来。

1. CUDA编程模型：从CPU到GPU的思维转变

先问大家一个问题：你写C++代码时，是不是默认代码是顺序执行的？嗯，CPU就是这种思维——一个核心跑一个线程，串行执行。但GPU不一样，它天生就是为并行而生的。

CUDA的编程模型，说白了就是「CPU负责调度，GPU负责干活」。CPU我们叫它Host，GPU叫Device。我习惯把CPU比作项目经理，GPU就是一群搬砖工人。项目经理把任务拆分成小块，然后喊一声：「兄弟们，开干！」

来看个最简单的例子：

// CPU端代码（Host）
int main() {
    int N = 1024;
    float *a, *b, *c;
    // 分配GPU显存
    cudaMalloc(&a, N * sizeof(float));
    cudaMalloc(&b, N * sizeof(float));
    cudaMalloc(&c, N * sizeof(float));
    
    // 调用GPU核函数
    vector_add<<<1, 256>>>(a, b, c, N);
    
    // 同步等待GPU完成
    cudaDeviceSynchronize();
    return 0;
}

// GPU端代码（Device）
__global__ void vector_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

注意那个 <<<1, 256>>> 语法，这是CUDA特有的。第一个参数是网格大小，第二个是线程块大小。我刚开始写时老搞混，后来记住一句话：「网格管全局，线程块管局部」。

核心要点：CUDA程序 = 串行代码（CPU） + 并行核函数（GPU）。核函数用 __global__ 修饰，调用时用 <<<grid, block>>> 指定并行规模。

2. GPU内存层次结构：速度与容量的博弈

GPU的内存，就像我们公司的仓库系统。有离工人最近的工具箱（寄存器），有小组共用的工作台（共享内存），还有整个工厂的大仓库（全局内存）。

我画了张图，帮你理清这个层次：

我在做实时风控系统时，就吃过内存选择的亏。一开始把所有数据都放全局内存，结果核函数跑得比蜗牛还慢。后来把热点数据挪到共享内存，性能直接提升了5倍。你想想看，这就是「用空间换时间」的典型例子。

我的经验：写CUDA代码时，先问自己三个问题：

这个数据每个线程都需要吗？→ 放常量内存
这个数据同一线程块内共享吗？→ 放共享内存
这个数据只属于当前线程？→ 放寄存器

3. 线程块与网格：GPU的并行组织方式

好了，现在咱们聊聊最核心的概念——线程块和网格。我习惯把GPU想象成一个大型工厂：

网格（Grid）：整个工厂，包含所有工人
线程块（Block）：每个车间，工人之间可以互相交流
线程（Thread）：每个工人，独立完成自己的任务

为什么要有这种分层设计？说白了，是为了硬件效率。GPU的SM（流多处理器）是以线程块为单位调度的。一个线程块内的线程，可以共享内存、同步执行。但不同线程块之间，老死不相往来。

来看个实际例子，假设我们要处理一张1024x1024的图片：

// 定义网格和线程块大小
dim3 grid(32, 32);      // 32x32 = 1024个线程块
dim3 block(32, 32);     // 每个线程块32x32 = 1024个线程

// 总线程数 = 1024 * 1024 = 1,048,576
image_process<<<grid, block>>>(d_input, d_output);

__global__ void image_process(float *in, float *out) {
    // 计算全局线程ID
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    int idx = y * 1024 + x;
    // 处理像素...
    out[idx] = in[idx] * 2.0f;
}

注意那个 blockIdx 和 threadIdx，这是CUDA内置变量。blockIdx告诉你当前在哪个车间，threadIdx告诉你在这个车间里的工位号。我刚开始总搞混，后来记住一句话：「blockIdx是车间号，threadIdx是工位号」。

曾经踩过的坑：线程块大小不是越大越好。每个SM的资源是有限的，线程块太大，会导致每个线程分到的寄存器变少，甚至发生寄存器溢出（spill），性能反而下降。我一般建议线程块大小取128或256，这是经验值。

4. 实战：在风控系统中如何选择线程块大小

回到咱们的实时风控场景。假设我们要对100万笔交易做实时评分，每笔交易的计算量差不多。你会怎么设计？

我的做法是这样的：

确定总任务量：100万笔交易，每笔一个线程
选择线程块大小：256（经验值，兼顾资源利用和调度效率）
计算网格大小：1000000 / 256 ≈ 3907个线程块

int total_transactions = 1000000;
int block_size = 256;
int grid_size = (total_transactions + block_size - 1) / block_size;

risk_scoring<<<grid_size, block_size>>>(d_transactions, d_scores, total_transactions);

这里有个小技巧：用 (N + block_size - 1) / block_size 来向上取整，确保不会漏掉任何交易。嗯，这个写法我用了好多年，从来没出过问题。

性能调优口诀：

线程块大小：128~256，别太大也别太小
网格大小：让GPU吃饱，但别撑到
共享内存：能放就放，但别放太多
全局内存访问：尽量合并，别乱跳

好了，这一章的内容就到这。CUDA编程模型、内存层次、线程块网格，这三个概念是GPU计算的基石。你把这些搞懂了，后面学起来就会轻松很多。