第19章多线程与GPU编程

课程目录第19章多线程与GPU编程

第19章多线程与GPU编程：CUDA基础概念、CPU线程与GPU线程协作、异构计算模型

说实话，我第一次接触GPU编程时，心里是有点抵触的。

你想啊，CPU多线程刚玩明白，又来一个GPU。但后来我在一个图像处理项目里，CPU跑了半天算不完一张4K图，换成CUDA后，几秒钟搞定。嗯，从那天起，我就知道这玩意儿必须学。

19.1 为什么需要GPU？CPU不够用吗？

CPU擅长处理复杂逻辑，分支预测、乱序执行、大缓存，样样精通。但CPU的核心数有限——你见过64核的CPU，但见过几千核的吗？

GPU不一样。它把大量晶体管用来做计算单元，而不是缓存和控制逻辑。说白了，GPU就是一台「简单计算器阵列」。

核心区别一句话：

CPU：适合串行任务、复杂逻辑、少量线程
GPU：适合并行任务、简单运算、海量线程

我在项目中遇到过这样一个场景：需要对100万条传感器数据做FFT。用CPU多线程，开了16个线程，跑了3秒多。换成GPU后，开了1万个线程，跑了0.1秒。你想想看，这差距有多大。

19.2 CUDA基础概念

CUDA是NVIDIA推出的并行计算平台。它允许你用C语言扩展，直接写GPU代码。我个人习惯把CUDA编程模型理解成三层结构：

19.2.1 线程层次结构

CUDA把线程组织成三层：

Thread（线程）：最基础的计算单元
Block（线程块）：一组线程，可以共享内存、同步
Grid（网格）：一组Block，构成整个任务

为什么这么设计？因为GPU硬件有SM（流多处理器），每个SM负责执行一个或多个Block。Block内的线程可以高效协作，Block之间则相对独立。

我的经验：Block大小一般设为128或256比较合适。太小了SM利用率低，太大了寄存器不够用。我曾经设过1024，结果性能反而下降——因为每个线程分到的寄存器太少，数据被迫 spill 到全局内存。

19.2.2 内存层次结构

CUDA的内存模型，说白了就是「谁离计算单元近，谁就快」：

内存类型	位置	访问速度	作用域
寄存器	SM内部	最快（1周期）	单个线程
共享内存	SM内部	快（~5周期）	Block内线程
全局内存	显存	慢（~400周期）	所有线程
常量内存	显存（带缓存）	较快	所有线程（只读）

嗯，这里要注意：全局内存虽然慢，但容量大。共享内存快，但每个Block只有几十KB。所以写CUDA程序，本质上就是在「快但小」和「慢但大」之间做权衡。

19.3 CPU线程与GPU线程协作

异构计算，说白了就是CPU和GPU一起干活。CPU负责控制流和复杂逻辑，GPU负责大规模并行计算。

一个典型的协作流程是这样的：

CPU准备数据，拷贝到GPU显存
CPU启动GPU内核（Kernel）
GPU执行计算，CPU可以干别的事（异步）
GPU计算完成，CPU把结果拷回来

我刚开始写CUDA时犯过一个低级错误：每次启动Kernel后，都调用 cudaDeviceSynchronize() 等GPU完成。结果CPU一直在空转。后来改成异步启动，CPU在GPU算的时候同时做数据预处理，整体吞吐量提升了30%。

避坑指南：我曾经在项目里把大量小数据频繁地在CPU和GPU之间拷贝，结果数据传输时间比计算时间还长。记住：尽量减少CPU-GPU数据传输，能一次拷完就别分多次。

19.4 异构计算模型

异构计算模型，我理解就是「各司其职」：

CPU（Host）：负责任务调度、I/O、复杂逻辑判断
GPU（Device）：负责数据并行、数值计算

为什么叫「异构」？因为CPU和GPU的架构完全不同。CPU有强大的分支预测和乱序执行，GPU有大量的ALU（算术逻辑单元）。你让CPU去算矩阵乘法，它也能算，但效率低。你让GPU去处理if-else分支多的逻辑，它也会卡住——因为GPU的线程是SIMT（单指令多线程）执行，分支会导致部分线程闲置。

举个实际例子。我在做一个实时视频处理项目时，把任务拆成两部分：

CPU负责：读取视频帧、解析协议、控制流程
GPU负责：每帧的像素处理（滤波、边缘检测、色彩转换）

这样CPU和GPU同时工作，互不阻塞。整个系统的吞吐量，比纯CPU方案高了10倍以上。

19.5 一个完整的CUDA示例

来看一个最简单的向量加法。别小看它，这是CUDA的「Hello World」：

// 向量加法的CUDA内核
__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1 << 20;  // 100万个元素
    size_t size = N * sizeof(float);

    // 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

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

    // 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 拷贝数据到GPU
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 启动内核
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 拷贝结果回CPU
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 清理
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
    free(h_A); free(h_B); free(h_C);

    return 0;
}

这段代码里，有几点我想强调一下：

blockIdx.x * blockDim.x + threadIdx.x 是CUDA里最常用的线程索引计算方式，建议背下来
if (i < N) 这个判断不能少，因为线程总数可能超过N
cudaMemcpy 是同步操作，会阻塞CPU直到拷贝完成

我的习惯：写CUDA代码时，我会先用 cudaMallocManaged 做统一内存管理，调试方便。等性能调优时再换成手动管理。这样开发效率高不少。

19.6 知识体系总览

下面这张图，是我自己总结的CUDA异构计算知识体系。你看一眼，就能明白本章讲了什么：

19.7 本章小结

CUDA编程，说白了就是学会「把问题拆成海量小任务，扔给GPU去算」。你不需要成为硬件专家，但得理解线程层次、内存模型、数据传输这几个核心概念。

我个人觉得，学CUDA最好的方式就是动手写。从向量加法开始，然后试试矩阵乘法，再做个图像滤镜。踩几次坑，你就记住了。

嗯，记住一点：CPU和GPU是搭档，不是对手。让CPU做它擅长的控制，让GPU做它擅长的计算，这才是异构计算的精髓。

公众号:蓝海资料掘金营，微信deep3321

📚 课程目录