3、GPU编程模型：CUDA编程模型概述、主机与设备、内核函数（Kernel）、线程层次结构（Grid、Block、Thread）

好，我们进入第三章。这一章，咱们聊聊CUDA编程模型。

说实话，很多初学者一上来就被CUDA的线程层次结构搞晕了。Grid、Block、Thread，这三个东西到底什么关系？我当年刚接触时也懵过一阵子。别急，我们一步步拆开来看。

3.1 CUDA编程模型概述

CUDA，全称是Compute Unified Device Architecture。说白了，就是NVIDIA给GPU写的一套编程框架。你想想看，GPU本来是为图形渲染设计的，现在要拿来做通用计算，总得有个统一的接口吧？CUDA就是干这个的。

我个人习惯把CUDA编程模型理解成“主从协作模式”。CPU是主机（Host），GPU是设备（Device）。CPU负责逻辑控制、串行任务，GPU负责大规模并行计算。两者各司其职，配合得好，性能就上去了。

核心思想：CUDA将GPU视为一个高度并行的协处理器，通过CPU发起计算任务，GPU执行大规模并行计算。

我在项目中遇到过不少新手，一上来就想把整个程序都扔到GPU上跑。嗯，这其实是个误区。GPU擅长的是数据并行，不是逻辑控制。你想想看，if-else分支多了，GPU的SIMT（单指令多线程）架构反而会效率下降。

3.2 主机与设备（Host & Device）

主机（Host）就是CPU及其内存系统。设备（Device）就是GPU及其显存系统。两者有各自独立的内存空间。

这就引出了一个问题：数据怎么在主机和设备之间传递？

CUDA提供了几个关键API：

  • cudaMalloc()：在设备上分配显存
  • cudaMemcpy()：在主机和设备之间拷贝数据
  • cudaFree()：释放设备上的显存

我曾经踩过一个坑：频繁地在主机和设备之间拷贝小数据。每次拷贝都有延迟开销，小数据量根本划不来。后来我改成批量拷贝，性能直接翻倍。避坑指南：尽量减少主机与设备之间的数据传输次数，能一次传完就别分十次。

注意：主机内存和设备内存是物理隔离的。你不能直接在GPU上访问CPU的内存指针，反之亦然。必须通过cudaMemcpy显式拷贝。

3.3 内核函数（Kernel）

内核函数，就是跑在GPU上的函数。用__global__修饰符声明。调用时用尖括号语法指定线程配置。

看个最简单的例子：

// 内核函数定义
__global__ void vecAdd(float* A, float* B, float* C, int n) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

// 主机端调用
int main() {
    // ... 数据准备 ...
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);
    // ... 结果回收 ...
}

这里有个关键点：内核函数是异步执行的。你调用完vecAdd<<<...>>>()后，CPU不会等它执行完就继续往下跑了。如果需要同步，得调用cudaDeviceSynchronize()。

我记得有一次调试，发现结果总是不对。查了半天，原来是忘了同步，CPU读到的是还没算完的中间结果。嗯，这个教训挺深刻的。

小技巧：内核函数里不要用printf做调试输出，会严重影响性能。我一般用cuda-gdb或者NVIDIA Nsight来调试。

3.4 线程层次结构（Grid、Block、Thread）

这是CUDA编程模型的核心，也是很多人容易搞混的地方。我们把它拆开来看。

Thread（线程）：最小的执行单元。每个线程执行内核函数的一份拷贝。

Block（线程块）：一组线程的集合。同一个Block内的线程可以共享显存（Shared Memory），也可以通过__syncthreads()进行同步。

Grid（网格）：一组Block的集合。一个内核函数对应一个Grid。

你想想看，这就像是一个三层组织架构：

  • Grid = 公司
  • Block = 部门
  • Thread = 员工

每个员工（Thread）干自己的活，部门内（Block）可以共享资源，公司（Grid）统筹全局。

在代码里，我们通过内置变量来获取当前线程的位置：

变量 含义
threadIdx 当前线程在Block内的索引
blockIdx 当前Block在Grid内的索引
blockDim Block的维度（包含多少个线程）
gridDim Grid的维度（包含多少个Block）

计算全局线程ID的公式：

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

这个公式我建议你背下来。几乎每个内核函数里都会用到。

我个人的习惯是，先把Block大小定下来，再根据数据总量算Grid大小。Block大小一般取32的倍数（因为warp大小是32），比如128、256、512。太小了浪费硬件资源，太大了又可能超过每个Block的最大线程数限制。

重要：Block内的线程数不能超过1024（对于现代GPU）。Grid内的Block数理论上可以很大，但受限于GPU的硬件资源。

举个例子，假设我们要处理100万个元素的向量加法：

int n = 1000000;
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;  // 向上取整

vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);

这里gridSize算出来是3907。每个Block处理256个元素，最后一个Block可能处理不满，所以内核函数里要有边界检查：if (i < n)。

说到边界检查，我曾经犯过一个低级错误：忘了加边界检查，结果最后一个Block的线程访问了越界内存，导致程序崩溃。嗯，从那以后，我写内核函数的第一件事就是加边界检查。

线程层次结构还可以是一维、二维、三维的。比如处理图像时，用二维索引会更直观：

__global__ void imageProcess(float* img, 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;
        // 处理像素
    }
}

dim3 blockSize(16, 16);
dim3 gridSize((width + 15) / 16, (height + 15) / 16);
imageProcess<<<gridSize, blockSize>>>(d_img, width, height);

你看，用二维Block来处理二维图像，代码逻辑清晰多了。

最后，总结一下这一章的核心要点：

  1. CUDA是主从协作模型，CPU负责控制，GPU负责并行计算
  2. 主机和设备有独立的内存空间，数据传输要显式拷贝
  3. 内核函数用__global__声明，用尖括号语法配置线程
  4. 线程层次结构：Grid → Block → Thread，三层组织
  5. 计算全局线程ID的公式要牢记
  6. 别忘了边界检查和同步

下一章，我们会深入讨论内存层次结构，看看怎么用好全局内存、共享内存和寄存器。到时候再聊。

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