第27章 GPU与异构计算入门

课程目录第27章 GPU与异构计算入门

GPU与异构计算入门：CUDA基础，数据并行模型，CPU与GPU协同优化

说实话，很多做嵌入式或者服务端C语言开发的朋友，一听到“GPU编程”就觉得那是做图形学的人搞的。我以前也这么想。直到有一次，我在一个实时视频处理项目里，CPU把四核跑满了，帧率还是上不去。老板催得紧，我硬着头皮上了CUDA。结果你猜怎么着？同样的算法，丢到GPU上，延迟从80毫秒降到了3毫秒。

从那天起我就明白了：异构计算不是锦上添花，而是性能瓶颈的破局点。今天我们就来聊聊，怎么用C语言配合CUDA，把GPU这个“大杀器”用起来。

一、数据并行模型：GPU为什么快？

CPU是“少而精”的架构——几个大核，每个核有复杂的控制单元和缓存。GPU呢？它是“多而糙”的——成百上千个小核，每个核简单得要命，但堆在一起，算力惊人。

说白了，CPU适合处理复杂的控制流（比如if-else嵌套、分支预测），GPU适合处理大规模的数据并行（比如对一万个像素做同样的滤波操作）。

1.1 数据并行的本质

数据并行，就是把一个大数据集拆成很多小块，每个小块交给一个线程去处理。所有线程执行相同的指令，只是操作的数据不同。

举个例子：

// CPU版本：循环处理
for (int i = 0; i < N; i++) {
    c[i] = a[i] + b[i];
}

// GPU版本：每个线程处理一个元素
__global__ void vec_add(float *a, float *b, float *c, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        c[i] = a[i] + b[i];
    }
}

你看，CPU版本是一个线程串行跑N次循环。GPU版本是N个线程同时跑，每个线程只算一次。这就是数据并行——用线程数量换时间。

核心思想： 数据并行 = 把问题拆成独立的小块 + 每个小块一个线程 + 所有线程同时执行。

1.2 线程层次结构

CUDA的线程组织方式，我刚开始学的时候也觉得绕。其实就三层：

线程（Thread）：最基础的执行单元，每个线程执行同一个kernel函数。
线程块（Block）：一组线程，可以共享内存，可以同步。一个block里的线程跑在同一个SM（流多处理器）上。
网格（Grid）：一组block，构成整个任务。

为什么要分block？因为硬件资源有限。一个SM能同时跑的线程数是有限的（比如1024个），但一个kernel可能有上百万个线程。所以CUDA把线程分组，每组一个block，硬件调度器把block分配到不同的SM上执行。

我的经验： block大小一般选128或256比较稳妥。太小了硬件利用率低，太大了可能超出SM的资源限制。我习惯先设256，跑一下profiler，再根据occupancy调整。

二、CUDA基础：从Hello World开始

先别急着写复杂算法。我们从一个最简单的kernel开始，感受一下CUDA的编程模型。

2.1 一个完整的CUDA程序

#include <stdio.h>

// 在GPU上运行的函数，用__global__修饰
__global__ void hello_from_gpu() {
    printf("Hello from thread %d in block %d\n",
           threadIdx.x, blockIdx.x);
}

int main() {
    // 启动kernel：1个grid，3个block，每个block 5个线程
    hello_from_gpu<<<3, 5>>>();

    // 等待GPU执行完毕
    cudaDeviceSynchronize();
    return 0;
}

编译运行：

nvcc hello.cu -o hello
./hello

你会看到15行输出，每个线程打印自己的blockId和threadId。注意，这些线程的执行顺序是不确定的——GPU调度器会乱序执行。

注意： printf在kernel里可以用，但仅限于调试。生产代码里别这么干，会严重拖慢性能。我曾经在线上代码里留了一个printf，结果帧率直接掉了一半……排查了半天才发现。

2.2 内存模型：CPU和GPU怎么传数据？

CPU有主存（DDR），GPU有显存（VRAM）。两者是物理隔离的。你要在GPU上算数据，得先把数据从主存拷贝到显存，算完再拷回来。

基本流程：

cudaMalloc：在显存上分配空间。
cudaMemcpy：把数据从主存拷到显存（方向：cudaMemcpyHostToDevice）。
启动kernel，在GPU上计算。
cudaMemcpy：把结果从显存拷回主存（方向：cudaMemcpyDeviceToHost）。
cudaFree：释放显存。

代码示例：

float *h_a, *h_b, *h_c;  // 主机端指针
float *d_a, *d_b, *d_c;  // 设备端指针
int N = 1024;
size_t size = N * sizeof(float);

// 1. 分配主机内存
h_a = (float*)malloc(size);
h_b = (float*)malloc(size);
h_c = (float*)malloc(size);

// 2. 分配设备内存
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);

// 3. 拷贝数据到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

// 4. 启动kernel
vec_add<<<(N+255)/256, 256>>>(d_a, d_b, d_c, N);

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

// 6. 清理
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
free(h_a); free(h_b); free(h_c);

关键点： 数据传输是PCIe总线走的，带宽远低于显存带宽。所以尽量减少CPU和GPU之间的数据交换。能一次拷完的，别分十次拷。

三、CPU与GPU协同优化

很多人以为，把代码丢到GPU上就万事大吉了。其实不然。CPU和GPU是协同工作的，不是替代关系。

3.1 异步执行与流

CUDA支持流（Stream）机制。你可以把多个操作（比如数据拷贝、kernel执行）放到不同的流里，让它们并行执行。

举个例子：

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

// 流1：拷贝数据A，计算A
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream1);
kernel_A<<<grid, block, 0, stream1>>>(d_a);

// 流2：拷贝数据B，计算B
cudaMemcpyAsync(d_b, h_b, size, cudaMemcpyHostToDevice, stream2);
kernel_B<<<grid, block, 0, stream2>>>(d_b);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

这样，数据拷贝和kernel执行可以重叠，隐藏传输延迟。我在做视频编码的时候，用流把输入帧的拷贝和上一帧的计算重叠起来，整体吞吐量提升了30%。

3.2 任务划分：什么该给GPU？

不是所有任务都适合GPU。我总结了一个简单的判断标准：

适合GPU的任务	不适合GPU的任务
大规模数据并行（图像、矩阵）	逻辑复杂的分支（大量if-else）
计算密集型（浮点运算多）	I/O密集型（频繁读写文件）
数据局部性好（连续内存访问）	递归、链表遍历
可分解为独立子任务	强依赖的串行任务

举个例子：图像滤波、FFT、矩阵乘法，这些是GPU的强项。而解析JSON、字符串匹配、数据库查询，这些还是留给CPU吧。

我的建议： 先用CPU实现一个正确版本，然后用profiler分析热点。如果某个函数占了总运行时间的80%以上，而且数据是并行的，那就值得移植到GPU。别一上来就全盘GPU化，容易翻车。

3.3 协同工作模式

实际项目中，我常用的模式是：

CPU负责控制流和I/O：读取数据、解析参数、调度任务。
GPU负责计算密集型部分：把数据打包成连续块，一次性传给GPU。
CPU做后处理：GPU算完的结果，CPU做最后的整理、输出。

这种模式的好处是：CPU和GPU各司其职，互不阻塞。CPU在等GPU结果的时候，可以去做别的事情（比如准备下一批数据）。

四、知识体系总览

下面这张图，是我自己梳理的CUDA异构计算知识结构。你可以把它当作学习路线图：

五、避坑指南与实战建议

最后，分享几个我踩过的坑，希望能帮你少走弯路。

5.1 别忽略错误检查

CUDA API调用几乎都会返回错误码。很多人（包括以前的我）懒得检查，结果程序跑出来结果不对，排查半天才发现是cudaMalloc失败了。

我习惯写一个宏：

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", \
                    __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(EXIT_FAILURE); \
        } \
    } while(0)

// 使用
CUDA_CHECK(cudaMalloc(&d_a, size));

5.2 小心bank conflict

共享内存（shared memory）是GPU上的一块高速缓存，但访问它时如果多个线程访问同一个bank的不同地址，就会发生bank conflict，导致串行化。解决办法是：让相邻线程访问相邻地址，或者用padding技巧。

5.3 我曾经犯过的错

有一次，我在kernel里用了大量的if (threadIdx.x % 2 == 0)这样的分支。结果性能惨不忍睹。为什么？因为GPU是SIMT架构，一个warp（32个线程）里的线程必须执行相同的指令。如果分支分歧，一部分线程执行if，另一部分等待，相当于一半的算力浪费了。

后来我把数据重新排列，让同一个warp里的线程处理相同类型的数据，分支就消除了，性能直接翻倍。

记住： GPU讨厌分支分歧，喜欢连续内存访问。写kernel的时候，脑子里要时刻想着“我这个warp里的32个线程在做什么”。

六、小结

GPU异构计算，说白了就是把对的事交给对的人。CPU擅长复杂控制，GPU擅长简单重复。两者配合好了，性能可以上一个台阶。

我个人觉得，学习CUDA最好的方式就是动手写。从向量加法开始，然后试试矩阵乘法，再然后做图像滤波。每写一个kernel，都用nvprof或者Nsight看看occupancy、带宽利用率。慢慢你就会找到感觉。

嗯，今天就聊到这里。记住：数据并行是思想，CUDA是工具，协同优化是艺术。多练，多调，多思考。

📚 课程目录

GPU与异构计算入门：CUDA基础，数据并行模型，CPU与GPU协同优化

一、数据并行模型：GPU为什么快？

1.1 数据并行的本质

1.2 线程层次结构

二、CUDA基础：从Hello World开始

2.1 一个完整的CUDA程序

2.2 内存模型：CPU和GPU怎么传数据？

三、CPU与GPU协同优化

3.1 异步执行与流

3.2 任务划分：什么该给GPU？

3.3 协同工作模式

四、知识体系总览

五、避坑指南与实战建议

5.1 别忽略错误检查

5.2 小心bank conflict

5.3 我曾经犯过的错

六、小结