第1章：CUDA编程模型入门

大家好，欢迎来到《实时系统GPU加速性能调优实战》的第一章。我是你们的老朋友，一个在GPU计算领域摸爬滚打了十几年的工程师。今天咱们聊聊CUDA编程模型的基础，这部分内容看似简单，但却是后续所有性能调优的根基。我个人习惯把这一章叫做「地基」，地基打不牢，后面盖再高的楼也容易塌。

1.1 CUDA平台概述：GPU计算的「操作系统」

CUDA是什么？说白了，它就是NVIDIA给GPU装上的一个「操作系统」。没有CUDA之前，你想用GPU做通用计算？那得用OpenGL或者DirectX这种图形API，把数据伪装成纹理，把计算伪装成渲染——想想就头大。

CUDA的出现，让GPU从「图形加速卡」变成了「通用并行计算设备」。你不需要懂图形学，只需要写C/C++风格的代码，就能让成千上万个线程同时跑起来。我在项目中遇到过不少新手，上来就问：「CUDA是不是只能用在NVIDIA的卡上？」嗯，这个问题问得好。答案是：是的，CUDA是NVIDIA的专有技术。但别急着失望，目前市面上绝大多数深度学习、科学计算场景，用的都是NVIDIA的GPU。

核心要点：CUDA = 并行计算架构 + 编程模型 + 指令集。它让你能用C/C++编写在GPU上运行的代码。

1.2 主机与设备：CPU和GPU的分工协作

在CUDA的世界里，有两个角色：主机（Host）和设备（Device）。主机就是CPU及其内存，设备就是GPU及其显存。它们各司其职，配合默契。

  • 主机（Host）：负责控制逻辑、串行任务、I/O操作。说白了，就是「大脑」，做决策的。
  • 设备（Device）：负责大规模并行计算。就是「肌肉」，干苦力的。

你想想看，一个程序里，CPU负责把数据准备好，然后「一声令下」让GPU去算，GPU算完再把结果传回来。这个过程中，数据需要在主机内存和设备显存之间来回搬运。我曾经在项目中犯过一个低级错误：每次迭代都把数据从主机拷贝到设备，算完再拷回来，结果90%的时间都花在数据传输上了。后来改成一次拷贝、多次计算，性能直接翻了10倍。

避坑指南：数据传输是GPU加速的「隐形杀手」。尽量减少主机与设备之间的数据交换，能一次传完就别分多次。

1.3 核函数（Kernel）：GPU上执行的函数

核函数，就是你在GPU上运行的函数。它的定义方式很简单：在函数前面加上 __global__ 修饰符。调用的时候，用特殊的语法 <<<grid, block>>> 来指定线程的组织方式。

来看一个最简单的例子：

// 定义一个核函数
__global__ void hello_from_gpu() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    // 调用核函数：1个grid，1个block，block里有10个线程
    hello_from_gpu<<<1, 10>>>();
    
    // 等待GPU执行完毕
    cudaDeviceSynchronize();
    return 0;
}

这段代码会输出10行「Hello from GPU thread 0」到「Hello from GPU thread 9」。注意，核函数是异步的——你调用它之后，CPU不会等它执行完就继续往下跑了。所以我们需要 cudaDeviceSynchronize() 来同步。

我记得第一次写核函数时，忘了加同步，结果CPU跑完了，GPU还在那慢慢算，程序就退出了。嗯，这个坑我踩过，你们就别再踩了。

小技巧：核函数里不能调用C++标准库的大部分函数（比如malloc、printf在特定架构下有限支持）。尽量只做纯计算，别搞花里胡哨的。

1.4 线程层次结构（Thread Hierarchy）：组织你的「千军万马」

GPU能同时跑成千上万个线程，但怎么组织它们？CUDA给出了一个清晰的层次结构：线程（Thread）→ 线程块（Block）→ 线程网格（Grid）。

层级 说明 最大数量（典型）
Thread 最小的执行单元，每个线程有自己的寄存器和局部内存 取决于硬件
Block 一组线程，可以共享内存，可以同步 每个Block最多1024个线程（现代GPU）
Grid 一组Block，构成整个计算任务 每个维度最多2^31-1

你想想看，一个Grid就像一支军队，Block就是各个连队，Thread就是每个士兵。连队内部可以互相通信（共享内存），但连队之间是独立的。这种设计有什么好处？

  • 可扩展性：Block可以独立调度到不同的SM（流多处理器）上。你用一个GPU跑，或者用一百个GPU跑，代码不用改。
  • 局部性：同一个Block里的线程可以高效共享数据，减少全局内存访问。

来看一个二维线程索引的例子：

__global__ void matrix_add(float *A, float *B, float *C, int N) {
    // 计算线程在全局的索引
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < N && col < N) {
        C[row * N + col] = A[row * N + col] + B[row * N + col];
    }
}

// 调用方式
dim3 blockDim(16, 16);  // 每个Block有16x16=256个线程
dim3 gridDim((N + 15) / 16, (N + 15) / 16);  // 计算需要的Block数量
matrix_add<<<gridDim, blockDim>>>(A, B, C, N);

这里用到了几个内置变量：

  • threadIdx：线程在Block内的索引（可以是1D、2D、3D）
  • blockIdx：Block在Grid内的索引
  • blockDim：Block的维度大小
  • gridDim：Grid的维度大小

我个人习惯用2D的Block来处理矩阵，因为这样索引更直观。但要注意，Block的维度不是越大越好。我曾经在项目中把Block设成1024个线程，结果因为寄存器不够用，性能反而下降了。嗯，这里要记住：线程数不是越多越好，够用就行。

黄金法则：Block大小通常是32的倍数（因为warp大小是32），推荐128或256。太小了浪费SM资源，太大了可能寄存器溢出。

本章小结

这一章我们聊了CUDA编程模型的核心概念：

  1. CUDA是NVIDIA的并行计算平台，让GPU能做通用计算。
  2. 主机（CPU）和设备（GPU）分工协作，数据传输是性能瓶颈。
  3. 核函数用 __global__ 定义，用 <<<grid, block>>> 调用。
  4. 线程层次结构：Thread → Block → Grid，决定了你的并行度。

下一章，我们会深入内存层次结构，聊聊全局内存、共享内存、寄存器这些「兵家必争之地」。到时候你会发现，很多时候性能瓶颈不在计算，而在内存访问。咱们下章见！

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