19、GPU并行编程入门：CUDA基础概念、GPU线程模型、CPU与GPU协作、cuBLAS简单示例

说到并行编程，很多人第一反应就是多线程、线程池、锁这些东西。但说实话，CPU的并行能力是有天花板的。你想想看，一颗顶级的服务器CPU也就几十个核心，撑死了上百个硬件线程。那如果我们要处理上百万个数据点呢？比如矩阵乘法、图像处理、深度学习训练——这时候，GPU就登场了。

我个人第一次接触CUDA是在做图像处理项目的时候。当时要在4K视频上做实时滤镜，CPU怎么优化都差一口气。后来把核心计算扔到GPU上，速度直接翻了20倍。嗯，从那以后我就再也没小看过GPU编程。

CUDA是什么？

CUDA，全称是Compute Unified Device Architecture，是NVIDIA推出的并行计算平台。说白了，就是让GPU不光能画图，还能干通用计算的活儿。

你可能会问：GPU和CPU到底有什么区别？我打个比方：CPU就像几个顶尖的数学家，能解非常复杂的方程；GPU就像几千个小学生，每个人只会做加减法，但人多力量大，几万个加减法同时算，速度碾压数学家。

CUDA的核心思想就是：把数据量大、计算逻辑简单的任务，扔给GPU去并行处理。

GPU线程模型：从Grid到Thread

GPU的线程模型，说白了就是三层结构：Grid → Block → Thread。我刚开始学的时候觉得这玩意儿绕，后来发现其实很好理解。

想象一下：你要处理一张1024×1024的图片，每个像素做一个简单的颜色变换。CPU的做法是循环遍历，一个像素一个像素算。GPU的做法呢？

• Thread（线程）：每个线程处理一个像素。这是最小的执行单元。
• Block（线程块）：一组线程打包在一起。同一个Block里的线程可以共享内存，也可以同步。我建议Block大小一般取32的倍数，比如128、256。为什么？因为GPU的warp（线程束）是32个线程一组，取32的倍数效率最高。
• Grid（网格）：所有Block的集合。Grid有多大，取决于你的数据量有多大。

举个例子，处理1024×1024的图片，你可以这样划分：

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

每个线程通过内置变量知道自己该处理哪个像素：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

这个公式我建议你背下来。我在项目中见过太多人写错坐标，结果图像处理出来全是乱的。

CPU与GPU协作：内存管理是关键

CPU和GPU是两套独立的硬件，各自有自己的内存。CPU的内存叫Host Memory，GPU的内存叫Device Memory。数据不能直接共享，必须显式地拷贝。

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

1. 在GPU上分配显存（cudaMalloc）
2. 把CPU上的数据拷贝到GPU（cudaMemcpy，方向：HostToDevice）
3. 在GPU上启动核函数（Kernel）进行计算
4. 把结果从GPU拷回CPU（cudaMemcpy，方向：DeviceToHost）
5. 释放GPU显存（cudaFree）

这里有个坑，我必须要说：cudaMemcpy是同步操作，它会阻塞CPU直到拷贝完成。而Kernel启动是异步的，CPU会继续往下执行。如果你在Kernel启动后立刻去读结果，读到的可能是旧数据。

一个完整的CUDA示例：向量加法

来看一个最简单的例子，两个数组相加。这个例子虽然简单，但包含了CUDA编程的所有核心步骤。

#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;  // 1048576个元素
    size_t bytes = n * sizeof(float);

    // 1. 分配CPU内存
    float *h_a = new float[n];
    float *h_b = new float[n];
    float *h_c = new float[n];

    // 初始化数据
    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. 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

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

    // 5. 同步并拷贝结果回CPU
    cudaDeviceSynchronize();
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

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

    // 6. 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    delete[] h_a;
    delete[] h_b;
    delete[] h_c;

    return 0;
}

这个代码里，我特别想强调一下if (i < n)这个判断。为什么要有它？因为GridSize是向上取整的，最后一个Block可能只有部分线程在处理有效数据。不加这个判断，就会访问越界。嗯，这个细节很容易被忽略。

cuBLAS：别自己造轮子

如果你要做矩阵运算，我强烈建议直接用cuBLAS。这是NVIDIA官方提供的线性代数库，针对GPU做了极致优化。你自己手写的矩阵乘法，大概率跑不过cuBLAS。

举个例子，用cuBLAS做矩阵乘法：

#include <cublas_v2.h>

// 假设 A(MxK), B(KxN), C(MxN)
// cuBLAS默认是列主序，这里用行主序需要转置
cublasHandle_t handle;
cublasCreate(&handle);

float alpha = 1.0f, beta = 0.0f;
cublasSgemm(handle,
    CUBLAS_OP_N, CUBLAS_OP_N,  // 不转置
    M, N, K,                    // 矩阵维度
    &alpha,                     // alpha
    d_A, M,                     // A矩阵及其leading dimension
    d_B, K,                     // B矩阵及其leading dimension
    &beta,                      // beta
    d_C, M);                    // C矩阵及其leading dimension

cublasDestroy(handle);

你看，几行代码就搞定了。我自己以前手写过矩阵乘法，各种优化技巧都用上了，性能还是只有cuBLAS的60%左右。所以我的建议是：能用库就别自己写，除非你是为了学习。

GPU编程的思维转变

从CPU编程转到GPU编程，最大的挑战不是语法，而是思维方式的转变。我总结了几点：

• 从串行到并行：不要想着"怎么循环"，要想着"怎么让每个线程处理一个数据"
• 关注内存带宽：GPU的计算能力很强，但数据搬运是瓶颈。尽量减少CPU和GPU之间的数据拷贝
• 避免分支发散：同一个warp里的线程如果走不同的分支，性能会下降。尽量让相邻线程执行相同的代码路径
• 善用工具：nvidia-smi看显存使用，nsys做性能分析，cuda-gdb调试。这些工具能帮你省下大量时间

这张图展示了CPU和GPU协作的完整流程。你看，数据从CPU内存搬到GPU显存，GPU算完再搬回来。中间的内存管理是重中之重。我个人习惯在项目里用cudaMallocManaged做统一内存管理，省去手动拷贝的麻烦，但性能上会有一点损失。具体用哪种，看你的场景。

---

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