4、CUDA编程模型：Host与Device、Kernel函数、线程层次结构
好，咱们今天聊聊CUDA编程模型的核心骨架。说白了，就是CPU和GPU怎么配合干活，以及GPU内部那成千上万个线程到底是怎么组织起来的。这部分内容我建议你当成「肌肉记忆」来练，因为后面所有的高性能优化，都建立在这个模型之上。
4.1 Host与Device：CPU是老板，GPU是工人
在CUDA的世界里，CPU被称为Host，GPU被称为Device。Host负责逻辑控制、串行任务、I/O操作；Device负责大规模并行计算。我刚开始接触时犯过一个低级错误——试图在GPU里直接打印调试信息，结果啥也没输出。后来才明白，GPU的核函数里不能直接调用printf（除非用特殊的printf函数），因为Device没有标准输出流。
Host和Device之间通过PCIe总线通信。数据需要先从Host内存拷贝到Device显存，计算完再拷回来。这个拷贝过程是瓶颈，我曾在项目中遇到过：一个核函数只跑了0.5毫秒，但数据拷贝花了10毫秒。嗯，这就是典型的「计算没吃饱，传输先累倒」。
核心原则：尽量减少Host与Device之间的数据传输。能一次拷完就别分十次，能复用显存就别反复申请释放。
4.2 Kernel函数：在GPU上跑的函数长啥样
Kernel函数就是跑在GPU上的函数。定义时用__global__修饰符，调用时用尖括号语法<<<grid, block>>>指定线程组织方式。来看个最简单的例子：
// 定义一个Kernel函数
__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];
}
}
// 在Host端调用
int main() {
// ... 分配内存、拷贝数据 ...
vecAdd<<<256, 256>>>(d_A, d_B, d_C, N);
// ... 拷贝结果回Host ...
}注意几个关键点：
__global__函数必须返回void，不能有返回值- Kernel是异步执行的——调用后CPU会继续往下跑，不会等GPU算完
- 如果需要同步，得手动调用
cudaDeviceSynchronize()
个人习惯：我写Kernel时，总会在调用后加一句cudaGetLastError()检查错误。因为Kernel启动失败不会抛异常，只会默默返回错误码。我曾经因为这个习惯，在项目上线前抓到了一个因显存不足导致的Kernel启动失败。
4.3 线程层次结构：Grid、Block、Thread
这是CUDA最巧妙的设计之一。GPU的线程不是平铺的，而是分层的：
- Thread：最基础的执行单元，每个线程执行同一段Kernel代码
- Block：一组线程的集合，共享一块共享内存，可以互相同步
- Grid：一组Block的集合，构成整个Kernel启动的线程总量
你想想看，为什么要有这个层次？因为GPU的硬件架构就是这样的：一个SM（流多处理器）可以同时运行多个Block，但Block内的线程必须在一个SM内。Block内的线程可以通过共享内存通信，Block之间则不行。
| 层级 | 硬件对应 | 通信方式 | 同步方式 |
|---|---|---|---|
| Thread | CUDA Core | 寄存器、本地内存 | 无（线程独立） |
| Block | SM | 共享内存、__syncthreads() | 块内同步 |
| Grid | GPU设备 | 全局内存 | 全局同步（需Kernel结束） |
每个线程都有一个唯一的ID，通过内置变量获取：
threadIdx.x：线程在Block内的索引blockIdx.x：Block在Grid内的索引blockDim.x：Block内线程的数量gridDim.x：Grid内Block的数量
计算全局线程ID的公式：int tid = threadIdx.x + blockIdx.x * blockDim.x;。这个公式我建议你手写十遍，因为几乎所有Kernel里都会用到。
我曾经踩过的坑：有一次我写了一个处理图像的Kernel，图像宽度是1024，我设了Block大小为1024，心想一个Block搞定。结果运行时报错——因为每个Block的线程数上限是1024（对于某些老架构是512）。后来我改成Block大小256，Grid大小4，问题解决。所以，一定要查你目标GPU的规格，别想当然。
4.4 线程索引的维度扩展
上面只用了x维度，但CUDA支持最多三维的线程索引：threadIdx.x、threadIdx.y、threadIdx.z。这在处理二维图像或三维体数据时特别方便。
// 二维Block示例：处理一张图像
__global__ void processImage(unsigned char* img, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x; // 线性化索引
if (x < width && y < height) {
// 对像素进行操作
img[idx] = img[idx] * 2; // 亮度增强
}
}
// 调用时指定二维Grid和Block
dim3 blockSize(16, 16); // 每个Block 16x16=256个线程
dim3 gridSize((width + 15) / 16, (height + 15) / 16);
processImage<<<gridSize, blockSize>>>(d_img, width, height);这里有个小技巧：用dim3类型来定义多维尺寸。如果只传一个整数，默认其他维度为1。比如<<<256, 256>>>等价于<<<dim3(256,1,1), dim3(256,1,1)>>>。
4.5 实战建议：如何选择Block和Grid大小
这个问题没有标准答案，但我可以分享一些经验法则：
- Block大小最好是32的倍数——因为warp（线程束）是32个线程一组，32的倍数能充分利用硬件
- Block大小建议128~512——太小则SM利用率低，太大则可能受限于寄存器/共享内存
- Grid大小要足够大——至少是SM数量的几倍，这样才能隐藏内存延迟
- 总线程数要覆盖所有数据——记得加边界检查，防止越界
我的调试习惯：先用小数据量（比如N=1024）跑通逻辑，再用cuda-memcheck检查内存访问越界。确认无误后，再上大数据量调优。别一上来就搞几百万数据，调试起来你会崩溃的。
好了，这一章的核心内容就是这些。Host与Device的分工、Kernel函数的写法、线程的三层结构——这三个概念就像CPU的寄存器、内存、缓存一样基础。你把这些搞透了，后面学共享内存、原子操作、流并行时就会轻松很多。