GPU并行计算基础：CUDA与C++的集成、GPU线程模型、CPU与GPU的协同工作

说到并行计算，CPU多线程我们聊了很多。但说实话，当你面对海量数据——比如处理一张4K图像、训练神经网络、或者做实时物理模拟——CPU那几十个核心根本不够看。这时候，GPU就该登场了。

我第一次接触CUDA是在做视频实时滤镜的项目。当时用CPU做高斯模糊，一帧要算200多毫秒，根本跑不动。换成CUDA之后，直接降到5毫秒。嗯，那种感觉，就像骑自行车突然换成了火箭。

CUDA是什么？为什么和C++集成这么自然？

CUDA是NVIDIA推出的并行计算平台。说白了，它让你能用C++写代码，然后在GPU上跑。不是学一套新语言，而是C++的扩展。

你想想看，CUDA C++就是在标准C++基础上加了几个关键字和API。比如：

• __global__ — 声明一个函数在GPU上执行，由CPU调用
• __device__ — 声明一个函数在GPU上执行，由GPU调用
• __host__ — 声明一个函数在CPU上执行（默认就是）
• cudaMalloc() / cudaMemcpy() — 管理GPU内存

我在项目中见过不少新手，以为CUDA是另一门语言，结果被吓退了。其实不是。你只要会C++，CUDA就是多学几个API的事。

GPU线程模型：从线程到网格

CPU的线程模型，每个线程是独立的、重量级的。GPU完全不同。GPU的线程是轻量级的，成千上万个线程同时跑。

CUDA把线程组织成三层结构：

层级	名称	说明
最外层	Grid（网格）	一个核函数启动，对应一个Grid
中间层	Block（线程块）	Grid由多个Block组成，Block内线程可共享内存
最内层	Thread（线程）	实际执行计算的最小单位

举个例子。你要处理一张1024x1024的图像，每个像素一个线程。你可以这样组织：

dim3 grid(32, 32);      // 32x32 = 1024个Block
dim3 block(32, 32);     // 每个Block有32x32 = 1024个线程
// 总线程数 = 1024 * 1024 = 1,048,576

kernel<<<grid, block>>>(d_data, width, height);

这里有个坑。我曾经以为Block越大越好，结果发现每个Block的线程数上限是1024（取决于GPU架构）。而且共享内存也有限。你想想看，一个Block里1024个线程抢那几十KB的共享内存，很容易就爆了。

CPU与GPU的协同工作

CPU和GPU不是竞争关系，而是协作关系。CPU擅长控制逻辑、分支预测、串行任务。GPU擅长数据并行、简单计算、吞吐量优先。

典型的协同流程是这样的：

1. CPU准备数据 — 从文件、网络或传感器读取数据
2. CPU拷贝数据到GPU — 通过PCIe总线传输
3. GPU执行核函数 — 成千上万个线程并行计算
4. CPU拷贝结果回内存 — 或者直接在GPU上做后续处理
5. CPU处理最终结果 — 比如写入文件、显示到屏幕

这里有个关键点：数据传输是瓶颈。PCIe的带宽虽然高（比如PCIe 4.0 x16是32GB/s），但跟GPU显存带宽（HBM2e可达2TB/s）比，差了两个数量级。

一个完整的CUDA C++示例

来看一个向量加法的例子。这是CUDA的"Hello World"。

#include <iostream>
#include <cuda_runtime.h>

// GPU核函数
__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;  // 1,048,576个元素
    size_t bytes = n * sizeof(float);

    // 1. CPU分配内存
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);

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

    // 2. GPU分配内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. CPU -> GPU 拷贝
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    // 5. GPU -> CPU 拷贝结果
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        std::cout << "c[" << i << "] = " << h_c[i] << std::endl;
    }

    // 7. 释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

这段代码，你仔细看。核心就三步：分配内存、拷贝数据、启动核函数。跟CPU编程的思维模式很像，只是多了"设备"这个概念。

GPU线程模型的核心概念图

下面这张图，我画了CUDA的线程层次结构。你一看就明白。

这张图里，Grid是最大的单位，里面有很多Block。每个Block里有很多Thread。你写代码的时候，就是通过<<<grid, block>>>来指定这个结构。

CPU与GPU协同工作的流程图

下面这张图，展示了CPU和GPU如何配合完成一次完整的计算任务。

你看，CPU和GPU各司其职。CPU做准备工作，GPU做重活。数据传输是中间的红线，也是性能瓶颈所在。

几个实战中容易踩的坑

我做了几年CUDA开发，踩过的坑不少。挑几个典型的说说：

• 忘记同步 — 核函数是异步的。你启动完核函数，CPU不会等它执行完。如果你紧接着就拷贝结果，拿到的可能是垃圾数据。记得加cudaDeviceSynchronize()。
• 内存泄漏 — GPU内存不会自动释放。你cudaMalloc了，必须cudaFree。我见过一个项目，跑了三天后OOM，就是因为没释放。
• 线程索引越界 — 总线程数可能超过数据量。比如你数据只有1000个，但启动了1024个线程。多出来的24个线程会访问非法内存。所以核函数里一定要加if (i < n)判断。
• 过度使用共享内存 — 共享内存很快，但容量很小。每个Block只有几十KB。你如果每个线程用太多，Block就调度不了那么多线程，并行度下降。

总结一下

GPU并行计算，说白了就是用数量换质量。CPU是几个大力士，GPU是一群小蚂蚁。单个线程很弱，但成千上万个一起上，力量就大了。

CUDA和C++的集成非常自然。你只要记住几个关键字和API，就能写出能跑的代码。但要想写得好，还得理解线程模型、内存层次、数据传输这些底层机制。

我个人建议，初学者先从简单的向量加法开始，跑通了再去做图像处理、矩阵运算。别一上来就搞复杂的，容易劝退。

---

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