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

好，我们正式开始CUDA编程的核心内容。这一章，我会带你理解CUDA编程模型中最基础、也最重要的几个概念。说白了，就是搞明白CPU和GPU怎么配合干活，以及GPU内部那成千上万个线程是怎么组织起来的。

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

在CUDA的世界里，CPU被称为主机（Host），GPU被称为设备（Device）。它们各司其职，协同工作。

• 主机（CPU）：负责逻辑控制、串行计算、I/O操作。它不适合做大规模并行计算。
• 设备（GPU）：负责大规模并行计算。它拥有成百上千个核心，适合处理数据并行任务。

我刚开始接触CUDA时，犯过一个低级错误：试图在GPU上直接调用printf打印调试信息。嗯，后来才知道，GPU的核函数里确实可以printf，但输出是异步的，而且顺序完全不可控。调试起来那叫一个痛苦。

2.2 核函数（Kernel）：在GPU上运行的函数

核函数，就是你在GPU上执行的函数。它用__global__修饰符声明，并且必须返回void。

看一个最简单的例子：

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

int main() {
    // 启动核函数，使用1个block，每个block里10个线程
    hello_from_gpu<<<1, 10>>>();
    
    // 等待GPU执行完成
    cudaDeviceSynchronize();
    
    return 0;
}

这里有个细节，我建议你注意：核函数的调用是异步的。也就是说，CPU发出调用指令后，不会等待GPU执行完，而是立即继续执行后面的代码。所以，如果你需要等待GPU完成，必须显式调用cudaDeviceSynchronize()。

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

这是CUDA编程模型中最精妙的设计之一。GPU的线程不是乱糟糟的一团，而是有清晰的层次结构。

层次	说明	类比
Grid（网格）	一个核函数启动的所有线程	整个工厂
Block（线程块）	一组可以协作的线程	车间
Thread（线程）	最小的执行单元	工人

你想想看，为什么要有这种层次结构？

原因很简单：GPU的硬件资源是有限的。一个Block内的线程可以共享内存（Shared Memory），也可以同步（__syncthreads()）。但不同Block之间的线程，是完全独立的，无法直接通信。

我个人习惯把Block看作一个「工作组」。组内的线程可以互相帮忙，组与组之间各干各的。这种设计，说白了就是为了让GPU能高效地调度成千上万个线程。

2.4 <<< >>> 语法：启动核函数的钥匙

这个三括号语法，是CUDA特有的。它用来指定核函数启动时的配置参数。

kernel_name<<<grid_dim, block_dim>>>(args...);

• grid_dim：网格的维度，即有多少个Block。可以是整数或dim3类型。
• block_dim：每个Block的维度，即有多少个Thread。同样可以是整数或dim3类型。

举个例子：

// 一维网格，一维Block
my_kernel<<<32, 256>>>(data);

// 二维网格，二维Block
dim3 grid_dim(4, 4);
dim3 block_dim(16, 16);
my_kernel<<<grid_dim, block_dim>>>(data);

这里有个重要的限制：每个Block的线程数不能超过1024（对于现代GPU）。这是硬件决定的。我记得有一次，一个同事写了个<<<1, 2048>>>，结果程序直接崩溃。嗯，这个坑我帮你踩过了。

2.5 线程索引：如何定位每个线程

每个线程都有自己唯一的ID。CUDA提供了几个内置变量来获取这个ID：

• threadIdx.x：线程在Block内的索引
• blockIdx.x：Block在Grid内的索引
• blockDim.x：每个Block的线程数
• gridDim.x：每个Grid的Block数

要计算全局唯一的线程ID，公式是：

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

这个公式，说白了就是：先找到你所在的Block的起始位置，再加上你在Block内的偏移量。

看一个完整的例子：

__global__ void vector_add(float *a, float *b, float *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

int main() {
    int n = 1000000;
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    
    // ... 分配内存、拷贝数据 ...
    
    vector_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);
    
    // ... 拷贝结果、释放内存 ...
    
    return 0;
}

注意那个if (tid < n)的判断。为什么需要它？因为数据量不一定正好是Block大小的整数倍。多出来的线程，必须被「关掉」，否则会访问越界内存。

2.6 小结

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

• 主机与设备的分工
• 核函数的定义与调用
• Grid-Block-Thread的三层线程层次结构
• <<< >>>语法的使用
• 线程索引的计算方法

这些概念，是后续所有CUDA优化的基础。你想想看，不理解线程怎么组织，怎么谈优化？下一章，我们会深入内存层次结构，看看怎么用好全局内存、共享内存和寄存器。

嗯，先消化这些吧。有问题随时问我。