2. CUDA编程模型：主机与设备、核函数概念、线程层次结构、<<<>>>语法

好，咱们正式开始聊CUDA编程模型。说实话，我刚接触CUDA那会儿，最困惑的就是「这玩意儿到底怎么跟CPU配合工作？」。你想想看，CPU和GPU是两个完全不同的处理器，一个擅长串行逻辑，一个擅长并行计算。怎么让它们协同干活？这就是CUDA编程模型要解决的核心问题。

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

在CUDA的世界里，我们把CPU叫做主机（Host），把GPU叫做设备（Device）。这两个角色分工非常明确：

• 主机（CPU）：负责控制逻辑、串行任务、数据准备、启动核函数。说白了，它是「老板」，发号施令的。
• 设备（GPU）：负责大规模并行计算。它是「打工人」，专门干那些重复性高、数据量大的脏活累活。

我刚开始做项目时犯过一个低级错误：直接在GPU上分配内存，然后试图用CPU去访问它。结果程序直接崩了。后来才明白——主机和设备有各自独立的内存空间。CPU不能直接访问GPU的显存，反之亦然。数据必须通过PCIe总线显式地拷贝过去。

2.2 核函数概念：在GPU上运行的函数

核函数（Kernel Function）是什么？说白了，就是在GPU上执行的函数。它由CPU调用，但在GPU上运行。你写一个普通的C函数，前面加个__global__修饰符，它就变成了核函数。

举个例子：

// 这是一个核函数
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x;  // 获取当前线程的索引
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

注意几个关键点：

• __global__ 告诉编译器：这个函数要在GPU上跑，但由CPU来调用。
• 核函数返回类型必须是 void。你不能让它返回一个值，因为GPU上有成千上万个线程同时执行，谁返回？
• 核函数里不能使用C++标准库（比如 printf 在早期版本也不行，现在可以了，但我不建议你在核函数里频繁打印）。

2.3 线程层次结构：网格、线程块、线程

这是CUDA最核心的概念之一。我当年学的时候，觉得这东西有点绕。但一旦理解了，你就掌握了GPU并行计算的精髓。

CUDA把线程组织成三层结构：

1. 线程（Thread）：最基础的执行单元。每个线程执行核函数中的同一段代码，但处理不同的数据。
2. 线程块（Block）：一组线程的集合。同一个块内的线程可以共享数据（通过共享内存），也可以同步（通过 __syncthreads()）。
3. 网格（Grid）：一组线程块的集合。一个核函数对应一个网格。

你可以把网格想象成一个工厂，线程块是车间，线程是工人。每个工人干同样的活，但处理不同的零件。

为什么要搞这么复杂？因为GPU的硬件架构决定了这种层次结构。一个线程块内的线程会被调度到同一个流多处理器（SM）上执行，这样它们才能高效地共享数据。而不同线程块之间是独立的，可以运行在不同的SM上，也可以运行在同一个SM上——这由硬件调度器决定，你不需要操心。

2.4 <<<>>>语法：启动核函数的「开关」

好了，现在你有了核函数，也理解了线程层次结构。怎么把两者结合起来？答案就是 <<<...>>> 语法。

这个语法看起来有点奇怪，像个三明治。它的作用就是告诉CUDA运行时：「我要启动这个核函数，用这么多线程块，每个线程块里放这么多线程。」

基本格式：

kernel_name<<<grid_dim, block_dim>>>(arguments);

其中：

• grid_dim：网格的维度，即线程块的数量。可以是整数（一维），也可以是 dim3 类型（二维或三维）。
• block_dim：每个线程块的维度，即每个块内的线程数量。同样可以是整数或 dim3。

举个例子：

// 启动核函数：使用 256 个线程块，每个块 512 个线程
vecAdd<<<256, 512>>>(d_A, d_B, d_C, N);

// 也可以使用 dim3 定义二维网格和二维线程块
dim3 gridDim(16, 16);   // 16x16 = 256 个线程块
dim3 blockDim(32, 32);  // 32x32 = 1024 个线程/块
kernel2D<<<gridDim, blockDim>>>(d_data, width, height);

我个人习惯用一维的网格和线程块来处理数组，用二维来处理图像或矩阵。三维用得比较少，除非是体数据或物理模拟。

2.5 完整的CUDA程序示例

把上面所有知识点串起来，写一个完整的向量加法程序。你感受一下：

#include <stdio.h>

// 核函数：向量加法
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 全局线程索引
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

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] = i * 1.0f;
        h_B[i] = 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;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // 5. 把结果从设备拷贝回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

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

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

    return 0;
}

这个流程你记一下：分配内存 → 拷贝数据到设备 → 启动核函数 → 拷贝结果回主机 → 释放内存。几乎所有CUDA程序都遵循这个模式。

2.6 总结

嗯，这一章的内容就这些。总结一下：

• 主机与设备：CPU是主机，GPU是设备，各自有独立内存。
• 核函数：用 __global__ 修饰，在GPU上执行，由CPU调用。
• 线程层次结构：网格 → 线程块 → 线程。块内线程可共享数据，块间独立。
• <<<>>>：启动核函数的语法，指定网格和线程块的维度。

这些概念是CUDA编程的基石。你如果能把它们理解透，后面的内容学起来会轻松很多。下一章咱们聊聊内存管理——这可是性能优化的重头戏。