第4章 CUDA编程模型

4. CUDA编程模型：主机与设备、Kernel函数、线程层次结构、Hello World实战

好，咱们正式开始写CUDA代码了。

这一章是整个异构计算的基石。说白了，你得先搞清楚CPU和GPU到底怎么配合干活，Kernel函数是个什么玩意儿，以及那套让GPU疯狂并行的线程层次结构。我当年刚接触CUDA时，就是被这些概念绕晕过。今天咱们一次性把它理清楚。

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

在CUDA的世界里，CPU叫主机（Host），GPU叫设备（Device）。它们各司其职，分工明确。

主机（CPU）：负责逻辑控制、串行任务、数据准备。它脑子好使，但干活慢。
设备（GPU）：负责大规模并行计算。它脑子简单，但干活快，一呼百应。

我个人习惯把CPU想象成项目经理，GPU就是手底下那几千号搬砖工。项目经理负责分派任务、检查结果，搬砖工只管闷头干。

核心流程：

CPU在主机端分配内存，准备好数据。
CPU把数据从主机内存拷贝到设备显存。
CPU调用Kernel函数，让GPU开始干活。
GPU执行完毕后，CPU再把结果从显存拷回内存。

嗯，这里要注意：主机和设备的内存是物理隔离的。你不能在CPU代码里直接访问GPU的显存，反之亦然。必须通过cudaMemcpy来搬运数据。我曾经有个同事，刚上手时忘了拷贝这一步，结果GPU算了个寂寞，CPU还傻等着——这种坑，踩一次就记住了。

4.2 Kernel函数：GPU上执行的函数

Kernel函数，就是跑在GPU上的函数。它用__global__修饰符声明，调用时用尖括号语法<<<grid, block>>>指定执行配置。

说白了，Kernel函数就是告诉GPU：「兄弟们，按这个配方，每人干一份活。」

Kernel函数的特点：

必须返回void类型。
用__global__修饰。
由CPU调用，在GPU上执行。
所有线程执行同一份代码，但处理不同的数据（SIMT模型）。

举个例子，一个最简单的Kernel函数长这样：

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

调用它的时候，你得告诉GPU开多少个线程：

// 启动10个线程
hello_from_gpu<<<1, 10>>>();
cudaDeviceSynchronize();  // 等待GPU执行完毕

你想想看，这10个线程会同时执行hello_from_gpu里的代码，但每个线程的threadIdx.x不同，所以打印出来的线程ID是0到9。这就是并行——代码只有一份，但执行上下文有10份。

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

这是CUDA最核心的设计之一。GPU的线程不是一盘散沙，而是有组织、有层次的。

层级	说明	类比
Thread（线程）	最小的执行单元，每个线程执行Kernel函数的一个实例。	一个搬砖工
Block（线程块）	一组线程的集合，可以协作（共享内存、同步）。	一个施工小队
Grid（线程网格）	一组Block的集合，对应一次Kernel调用。	整个工地

为什么要有这种层次结构？

原因很简单：GPU的硬件资源是有限的。一个Block里的线程可以共享一块叫「共享内存」的高速缓存，还能通过__syncthreads()同步。但不同Block之间是独立的，不能直接通信。这种设计让GPU能灵活调度——哪个Block先算完，就腾出资源给下一个，互不干扰。

我记得刚开始写CUDA时，总想把所有线程都塞到一个Block里。结果发现Block里的线程数上限是1024（现代GPU），而且共享内存也有限。后来才明白，Grid负责「任务划分」，Block负责「局部协作」。这个思路一旦理清，写代码就顺了。

避坑指南：

我曾经在项目中把Block维度设得太大，导致每个Block的共享内存不够用，程序直接崩了。后来我养成了一个习惯：先查GPU的硬件参数（用cudaGetDeviceProperties），再根据任务量合理划分Grid和Block的尺寸。

4.4 Hello World实战：第一个CUDA程序

光说不练假把式。咱们来写一个完整的CUDA Hello World程序。这个程序会启动若干个GPU线程，每个线程打印自己的身份信息。

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

// Kernel函数：每个线程打印自己的坐标
__global__ void hello_kernel() {
    // 获取线程在Block内的索引
    int tid = threadIdx.x;
    // 获取Block在Grid内的索引
    int bid = blockIdx.x;
    // 获取Block的维度
    int bdim = blockDim.x;

    // 计算全局线程ID
    int global_tid = bid * bdim + tid;

    printf("Hello from Block %d, Thread %d (Global ID: %d)\n",
           bid, tid, global_tid);
}

int main() {
    // 设置Grid和Block的尺寸
    int grid_size = 2;    // 2个Block
    int block_size = 4;   // 每个Block有4个线程

    printf("Launching kernel with %d blocks, %d threads per block\n",
           grid_size, block_size);

    // 调用Kernel函数
    hello_kernel<<<grid_size, block_size>>>();

    // 等待GPU执行完毕
    cudaDeviceSynchronize();

    printf("Kernel execution completed.\n");
    return 0;
}

编译和运行：

nvcc hello.cu -o hello
./hello

输出结果大致如下：

Launching kernel with 2 blocks, 4 threads per block
Hello from Block 0, Thread 0 (Global ID: 0)
Hello from Block 0, Thread 1 (Global ID: 1)
Hello from Block 0, Thread 2 (Global ID: 2)
Hello from Block 0, Thread 3 (Global ID: 3)
Hello from Block 1, Thread 0 (Global ID: 4)
Hello from Block 1, Thread 1 (Global ID: 5)
Hello from Block 1, Thread 2 (Global ID: 6)
Hello from Block 1, Thread 3 (Global ID: 7)
Kernel execution completed.

你看到了什么？

8个线程，分属2个Block，每个Block里4个线程。它们同时执行，但打印的内容不同。这就是CUDA并行最基本的模样。

几个关键点：

threadIdx.x：线程在Block内的索引，从0开始。
blockIdx.x：Block在Grid内的索引，从0开始。
blockDim.x：每个Block包含的线程数。
cudaDeviceSynchronize()：阻塞CPU，直到GPU上的所有任务完成。没有这行，CPU可能先跑完，程序就退出了，GPU的打印还没输出。

我个人习惯在调试阶段，每个Kernel调用后都加一句cudaGetLastError()检查错误。因为Kernel调用是异步的，错误不会立即抛出。加上这句，能帮你快速定位问题：

hello_kernel<<<grid_size, block_size>>>();
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}

嗯，这个习惯帮我省了不少调试时间。你想想看，如果Kernel启动参数写错了（比如Block尺寸超过了硬件限制），没有错误检查，程序可能默默跑完但结果全错——这种bug最难找。

本章小结

这一章我们搞定了三件事：

主机与设备：CPU负责调度，GPU负责计算，数据通过cudaMemcpy搬运。
Kernel函数：用__global__定义，用<<<grid, block>>>启动，所有线程执行同一份代码。
线程层次结构：Grid包含Block，Block包含Thread。合理划分才能发挥GPU的威力。

下一章，我们会深入线程索引的计算，以及如何用多维Grid和Block处理更复杂的数据结构。到时候你会发现，这套层次结构的设计其实非常优雅。