第二章：CUDA编程模型——主机与设备、核函数与线程层次

各位同学，欢迎来到第二章。上一章我们聊了GPU为什么快，这一章咱们直接上手写代码。说实话，CUDA编程模型是整个课程的地基，你把它搞明白了，后面所有优化技巧都是锦上添花。我见过太多人一上来就调优，结果连核函数怎么启动都没搞对——嗯，那基本就是白忙活。

2.1 主机与设备：CPU和GPU的分工

先问个问题：你的程序里，CPU和GPU到底谁说了算？

答案是：CPU是老板，GPU是打工的。CPU（我们叫它主机，Host）负责逻辑控制、I/O、串行任务。GPU（我们叫它设备，Device）只干一件事——拼命算。

我刚开始学CUDA时犯过一个低级错误：试图在GPU上直接调用printf打印调试信息。结果呢？编译通过了，运行啥也没打印。后来才明白，GPU没有独立的标准输出，你得先把数据拷回主机再打印。这个坑，我替你们踩过了。

主机和设备之间的数据传递，靠的是PCIe总线。说白了就是一条高速公路，但带宽再高也有延迟。所以我的个人习惯是：尽量减少主机和设备之间的数据拷贝，能一次拷完就别分十次。

核心原则：

  • CPU负责控制流和少量串行计算
  • GPU负责大规模并行计算
  • 数据传递通过PCIe，开销不小，能省则省

2.2 核函数定义与调用：GPU上跑的“函数”

核函数（Kernel）就是跑在GPU上的函数。定义它很简单，加个__global__修饰符就行。但有几个硬性规定，你记牢了：

  • 返回类型必须是void
  • 不能使用可变参数（比如...）
  • 不能是类的成员函数（静态成员函数可以）
  • 调用的其他函数也必须是__device__或__global__

来看个最简单的例子：

// 核函数定义：给数组每个元素加1
__global__ void add_one(float *d_data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        d_data[idx] += 1.0f;
    }
}

// 主机端调用
int main() {
    float *d_data;
    int n = 1024;
    // ... 分配显存、拷贝数据 ...
    
    // 启动核函数：1个block，每个block里n个线程
    add_one<<<1, n>>>(d_data, n);
    
    // ... 拷贝结果回主机 ...
    return 0;
}

注意看那个<<<1, n>>>，这就是CUDA里最标志性的语法。我当年第一次看到这个三括号时，心想：这什么鬼？后来用习惯了，反而觉得它挺直观——左边是grid大小，右边是block大小。

个人经验：核函数里一定要做边界检查。就是那个if (idx < n)。为什么？因为线程总数可能比数据量多，不做检查就会访问越界。我曾经在一个项目里忘了加，结果数据量不是线程数的整数倍，跑出来的结果时对时错，排查了整整两天。

2.3 <<<>>>语法详解：启动配置的奥秘

三括号语法是CUDA的启动配置，格式是：

<<<grid_dim, block_dim, shared_mem_bytes, stream>>>

四个参数，但后两个可以省略。我来拆解一下：

参数 含义 类型 默认值
grid_dim grid的维度，即有多少个block dim3 或 int 无（必填）
block_dim block的维度，即每个block有多少线程 dim3 或 int 无（必填）
shared_mem_bytes 每个block动态分配的共享内存大小（字节） size_t 0
stream CUDA流，用于异步执行 cudaStream_t 0（默认流）

举个例子：

// 一维grid，一维block
kernel<<<32, 256>>>(...);

// 二维grid，二维block
dim3 grid_dim(4, 4);
dim3 block_dim(16, 16);
kernel<<<grid_dim, block_dim>>>(...);

// 带共享内存和流
kernel<<<grid_dim, block_dim, 1024, stream1>>>(...);

这里有个关键点：block_dim的三个维度乘积不能超过1024（对于现代GPU，老架构是512或768）。也就是说，一个block里最多1024个线程。你想想看，如果每个block开1024个线程，那grid_dim就决定了总线程数。比如<<<100, 1024>>>就是10万个线程。

避坑指南：我曾经在项目里把grid_dim和block_dim写反了，写成<<<256, 32>>>，结果每个block只有32个线程，浪费了大量SM的计算资源。性能直接掉了一半。所以写的时候一定要想清楚：左边是block数量，右边是每个block的线程数。

2.4 线程层次结构：Grid、Block、Thread

CUDA的线程组织方式，说白了就是三级结构：

  • Thread（线程）：最小的执行单元，每个线程执行核函数的一份拷贝
  • Block（线程块）：一组线程，可以协作（通过共享内存和同步）
  • Grid（线程网格）：一组block，构成整个核函数

我画了一张图，帮你理解这个层次关系：

Grid（线程网格） Block (0,0) Block (1,0) Block (2,0) T(0,0) T(1,0) T(2,0) T(3,0) T(0,1) T(1,1) T(2,1) T(3,1) T(0,2) T(1,2) T(2,2) T(3,2) ... ... ... = Block = Thread 内置变量：gridDim, blockIdx, blockDim, threadIdx

每个线程在核函数里可以通过内置变量知道自己是谁：

  • threadIdx：线程在block内的索引（可以是三维）
  • blockIdx：block在grid内的索引（可以是三维）
  • blockDim：block的维度，即每个block有多少线程
  • gridDim：grid的维度，即有多少个block

计算全局唯一ID的公式，我建议你背下来：

// 一维情况
int tid = threadIdx.x + blockIdx.x * blockDim.x;

// 二维情况
int tid_x = threadIdx.x + blockIdx.x * blockDim.x;
int tid_y = threadIdx.y + blockIdx.y * blockDim.y;
int global_idx = tid_y * gridDim.x * blockDim.x + tid_x;

为什么要有三维？说白了就是为了方便处理图像、体素这类天然带空间结构的数据。你处理一张2D图片，用二维grid和二维block，索引直接对应像素坐标，多直观。

我的建议：刚开始学，先用一维。把一维玩熟了，再上二维、三维。我见过有人一上来就搞三维grid，结果索引算错了，数据全乱套。记住：简单就是美。

2.5 一个完整的例子：向量加法

理论说再多，不如跑一个完整的例子。咱们写个向量加法，把两个数组逐元素相加：

#include <stdio.h>
#include <cuda_runtime.h>

// 核函数：向量加法
__global__ void vec_add(float *a, float *b, float *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;
    size_t bytes = n * sizeof(float);
    
    // 主机端分配
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);
    
    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = (n - i) * 1.0f;
    }
    
    // 设备端分配
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    // 拷贝数据到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    
    // 启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 等待GPU完成
    cudaDeviceSynchronize();
    
    // 拷贝结果回主机
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }
    
    // 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

注意那个blocks_per_grid的计算方式：(n + threads_per_block - 1) / threads_per_block。这是CUDA里最经典的向上取整写法。为什么？因为数据量不一定能被线程数整除，多出来的数据需要额外的block来处理。核函数里的if (idx < n)就是干这个的。

关键点总结：

  1. 主机负责控制，设备负责计算
  2. 核函数用__global__定义，返回void
  3. <<<grid, block>>>左边是block数，右边是每个block的线程数
  4. 线程层次：grid → block → thread，通过内置变量定位
  5. 边界检查是必须的，别偷懒

好了，这一章的内容就到这儿。记住：CUDA编程模型的核心就是理解主机和设备的分工，以及线程的组织方式。你把这个搞透了，后面学共享内存、流、原子操作什么的，都会轻松很多。

想要更多资料，联系可可

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