3、CUDA编程模型入门：主机与设备、核函数概念、线程层次结构、第一个CUDA程序

好，咱们今天聊聊CUDA编程模型。说实话，很多初学者一上来就被各种概念砸晕了——主机、设备、核函数、线程块……我当年刚接触时也懵过一阵。但别急，这些东西说白了就是一套规则，搞懂了就很简单。

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

先搞清楚两个角色：主机（Host） 和 设备（Device）。

主机：就是CPU及其内存。它负责控制流程、逻辑判断、I/O操作。
设备：就是GPU及其显存。它负责大规模并行计算。

我在项目中遇到过不少新手，把CPU该干的活也扔给GPU，结果性能反而更差。记住：CPU擅长复杂控制，GPU擅长简单重复。你想想看，让一个厨师去切一万根葱，他肯定不如一台切菜机快——GPU就是那台切菜机。

核心原则：数据在主机和设备之间传输是有代价的。尽量减少传输次数，能一次传完就别分十次。

3.2 核函数（Kernel）：GPU上执行的函数

核函数，就是跑在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];
    }
}

注意几个要点：

核函数必须返回 void
用 <<<grid, block>>> 语法调用
核函数内部不能调用主机函数（比如 printf 在早期版本不行，现在可以了但别依赖）

个人习惯：我写核函数时，总会在开头加一个边界检查 if (i < N)。为什么？因为线程总数可能超过数据量，不检查就会访问越界。我曾经因为这个bug排查了整整一个下午……

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

这是CUDA最核心的概念。我画个图帮你理解：

这张图展示了CUDA的线程层次：

Thread（线程）：最小的执行单元。每个线程执行相同的核函数，但处理不同的数据。
Block（线程块）：一组线程的集合。同一个block内的线程可以共享内存、同步执行。
Grid（网格）：一组block的集合。一个核函数对应一个grid。

你可能会问：为什么要搞这么复杂？直接开一万个线程不行吗？

嗯，这里要注意。GPU的硬件结构决定了：同一个block内的线程跑在同一个SM（流多处理器）上。如果block太大，资源不够用；如果block太小，硬件利用率低。我一般建议block大小取128或256，这是经验值。

避坑指南：我曾经把一个block设了1024个线程，结果程序直接崩了。后来查文档才发现，每个block的线程数上限是1024（不同架构略有差异）。而且不是设满就好，还要考虑寄存器、共享内存的占用。

3.4 第一个CUDA程序：向量加法

光说不练假把式。咱们写个完整的向量加法程序，把上面讲的概念串起来。

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

// 核函数：向量加法
__global__ void vecAddKernel(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 = 1000000;
    size_t size = N * sizeof(float);

    // 1. 分配主机内存
    float *h_A = (float*)malloc(size);
    float *h_B = (float*)malloc(size);
    float *h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 3. 从主机拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 配置核函数执行参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    // 5. 调用核函数
    vecAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 6. 将结果拷贝回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 7. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 8. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

这个程序虽然简单，但包含了CUDA编程的完整流程：

分配主机内存：用 malloc 或 cudaMallocHost（后者更快）
分配设备内存：用 cudaMalloc
数据传输：cudaMemcpy 是必经之路
配置执行参数：<<<blocksPerGrid, threadsPerBlock>>>
调用核函数：异步执行，CPU不会等GPU
结果回传：别忘了把结果拿回来
释放内存：不释放会内存泄漏，我见过有人跑一晚上把显存占满的

小技巧：我习惯在每次 cudaMemcpy 和 cudaMalloc 后面加错误检查。用 cudaGetLastError() 或 cudaPeekAtLastError()。别嫌麻烦，出bug时能省你半天时间。

3.5 编译与运行

编译CUDA程序用 nvcc 编译器：

nvcc -o vecAdd vecAdd.cu
./vecAdd

如果看到输出 C[0] = 3.0, C[1] = 3.0, ...，恭喜你，第一个CUDA程序跑通了！

说实话，我第一次跑通时还挺激动的。虽然只是简单的1+2=3，但背后是成千上万个线程在GPU上同时工作。这种感觉，嗯，挺奇妙的。

概念	说明	类比
Host	CPU及其内存	项目经理，负责调度
Device	GPU及其显存	工人，负责干活
Kernel	GPU上执行的函数	施工图纸
Thread	单个执行线程	一个工人
Block	线程的集合	一个班组
Grid	Block的集合	整个工地

好了，这一章的内容就这些。记住：主机负责控制，设备负责计算，核函数是桥梁，线程层次是骨架。把这几个概念刻在脑子里，后面的路就好走了。

📚 课程目录

3、CUDA编程模型入门：主机与设备、核函数概念、线程层次结构、第一个CUDA程序

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

3.2 核函数（Kernel）：GPU上执行的函数

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

3.4 第一个CUDA程序：向量加法

3.5 编译与运行