84、C++项目实战：GPU编程（CUDA入门）

说实话，很多C++开发者听到“GPU编程”就觉得头大。我当年也是这样——总觉得那是搞图形学的人才会碰的东西。直到有一次，我接手了一个图像处理项目，CPU跑一张4K图片要好几秒，用户直接骂娘。嗯，从那以后，我就老老实实学起了CUDA。

今天这一章，咱们不扯虚的。直接上手，带你写一个能跑的CUDA程序。你不需要有显卡，但最好有个NVIDIA的GPU。如果没有，用Google Colab也行，我试过，免费的T4够用了。

CUDA到底是什么？

说白了，CUDA就是NVIDIA搞的一套并行计算平台。它允许你用C++的语法，直接调用GPU里的几千个小核心来干活。

CPU像是一个全能型选手，什么都能干，但一次只能干几件事。GPU呢？它像是一群只会做简单算术的工人，但数量多啊——几千个一起上，处理海量数据时，CPU根本追不上。

核心思想：把数据拆成小块，让每个GPU线程处理一小块。这就是所谓的“单指令多线程”（SIMT）模型。

先搭个环境

我个人习惯用Visual Studio + CUDA Toolkit。安装步骤其实很简单：

  1. 去NVIDIA官网下载CUDA Toolkit（我建议用11.x版本，稳定）
  2. 安装时选默认路径就行
  3. VS里新建项目，选“CUDA Runtime”模板

如果你用的是Linux，那就更简单了：

sudo apt install nvidia-cuda-toolkit
nvcc --version  # 检查是否装好

装好后，写个最简单的程序验证一下：

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

int main() {
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    std::cout << "检测到 " << deviceCount << " 块GPU" << std::endl;
    return 0;
}

如果输出不是0，恭喜你，环境搞定了。

第一个CUDA程序：向量加法

咱们从最经典的例子开始——两个数组相加。CPU版本大家都会写：

void add_cpu(int* a, int* b, int* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

GPU版本怎么写呢？看好了：

// 这是GPU上运行的函数，叫kernel
__global__ void add_gpu(int* a, int* b, int* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    int size = n * sizeof(int);

    // 分配GPU内存
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 把数据从CPU拷贝到GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 启动kernel：1个block，1024个线程
    add_gpu<<<1, 1024>>>(d_a, d_b, d_c, n);

    // 把结果拷回来
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 释放GPU内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

你可能会问：<<<1, 1024>>> 这玩意儿是什么鬼？

这是CUDA的语法，用来指定启动多少个线程。第一个数字是block数量，第二个是每个block里的线程数。这里我用了1个block，1024个线程，刚好处理1024个元素。

我的经验：刚开始学CUDA时，我总搞混block和thread的关系。你只要记住：block是组织单位，thread是干活单位。一个block里的线程可以共享数据，不同block的线程不能直接通信。

理解线程层次结构

GPU的线程组织方式，有点像军队编制：

  • 线程（Thread）：最基础的执行单元，每个线程执行相同的代码
  • 线程块（Block）：一组线程，可以共享内存和同步
  • 网格（Grid）：所有block的集合，对应一次kernel启动

我画了张图，帮你理解这个结构：

Grid（网格） Block (0,0) Block (1,0) Block (2,0) Thread 0 Thread 1 ... Thread 32 Thread 33 ... 共享内存（Shared Memory） 线程间可通信 一个Grid包含多个Block，每个Block包含多个Thread

内存管理：别踩坑

GPU编程里，内存管理是最容易出问题的地方。我曾经在项目里因为忘记释放GPU内存，导致程序跑着跑着就崩了，查了半天才发现是内存泄漏。

记住这几个关键函数：

函数 作用 注意事项
cudaMalloc 在GPU上分配内存 返回的是设备指针，不能在CPU上直接访问
cudaMemcpy 在CPU和GPU之间拷贝数据 要指定方向：HostToDevice 或 DeviceToHost
cudaFree 释放GPU内存 不释放会导致显存泄漏
cudaMallocManaged 统一内存（新手友好） 性能稍差，但不用手动拷贝

避坑指南：我曾经在项目里用 cudaMallocManaged 偷懒，结果在大数据量时性能惨不忍睹。后来老老实实改回手动管理内存，速度提升了3倍。记住：统一内存方便是方便，但别在高性能场景下用。

性能优化：从入门到进阶

写一个能跑的CUDA程序不难，但要写得快，就得注意几个点：

1. 线程块大小怎么选？

我一般用256或512。太小了浪费GPU资源，太大了可能超过硬件限制。你可以试试看：

// 常见的配置
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
add_gpu<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

2. 内存访问要连续

GPU喜欢连续的内存访问模式。如果你跳着访问，性能会暴跌。举个例子：

// 好的方式：连续访问
int idx = threadIdx.x + blockIdx.x * blockDim.x;
c[idx] = a[idx] + b[idx];  // 每个线程访问相邻地址

// 坏的方式：跳跃访问
int idx = threadIdx.x * 1024 + blockIdx.x;
c[idx] = a[idx] + b[idx];  // 线程之间隔了1024个元素

3. 减少CPU和GPU之间的数据传输

PCIe带宽有限，频繁拷贝数据会拖慢整体速度。我的做法是：尽量在GPU上完成所有计算，只把最终结果拷回来。

一个小技巧：如果你有多个小数据要传输，可以合并成一次大的传输。比如把10次1MB的传输，合并成1次10MB的传输，速度能快不少。

实战：图像灰度化

咱们来写个有点实际意义的程序——把彩色图像转成灰度图。假设图像是RGB格式，每个像素3个字节。

__global__ void rgb_to_gray(unsigned char* rgb, unsigned char* gray, int width, int height) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    
    if (x < width && y < height) {
        int idx = y * width + x;
        int r = rgb[idx * 3];
        int g = rgb[idx * 3 + 1];
        int b = rgb[idx * 3 + 2];
        // 标准灰度转换公式
        gray[idx] = 0.299f * r + 0.587f * g + 0.114f * b;
    }
}

// 启动方式
dim3 blockSize(16, 16);
dim3 gridSize((width + 15) / 16, (height + 15) / 16);
rgb_to_gray<<<gridSize, blockSize>>>(d_rgb, d_gray, width, height);

这里我用了二维的block和grid，因为图像本身就是二维的。每个线程处理一个像素，逻辑上很清晰。

调试与性能分析

写CUDA程序最头疼的就是调试。CPU上的断点到了GPU上就不灵了。我的经验是：

  • 先用小数据量测试：比如只处理10个元素，确保逻辑正确
  • 检查返回值：每个CUDA API调用后，检查返回值是不是 cudaSuccess
  • 用 cuda-memcheck：这个工具能帮你找出越界访问等问题
// 检查错误的宏
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            std::cerr << "CUDA错误: " << cudaGetErrorString(err) << std::endl; \
            exit(1); \
        } \
    } while(0)

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

性能分析工具：NVIDIA的Nsight工具很好用。它能告诉你每个kernel花了多少时间，内存传输占了多少，哪里是瓶颈。我每次优化前都会跑一遍，看看问题出在哪。

总结一下

CUDA入门其实没那么可怕。你只要记住三件事：

  1. 写kernel：用 __global__ 修饰，在GPU上跑
  2. 管理内存：分配、拷贝、释放，一步都不能少
  3. 启动线程：用 <<<grid, block>>> 指定线程数量

剩下的就是多写多练。我刚开始学的时候，光是向量加法就写了十几遍，每次都能发现新的问题。嗯，慢慢来，别着急。

对了，如果你在Windows上遇到编译问题，记得检查项目属性里的“CUDA C/C++”配置。我之前被这个坑过，折腾了一下午才发现是路径没配好。

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

← 上一章 📖 返回目录 下一章 →