3、CUDA编程基础回顾：CUDA编程模型，Kernel Launch，Grid-Block-Thread层次，内存管理API

好，咱们正式开始聊CUDA。说实话，很多做AI推理优化的朋友，一开始都觉得CUDA是搞图形学的人学的。其实不然。你想想看，现在大模型这么火，哪个推理框架底层离得开CUDA？我当年刚接触GPU编程时，也踩过不少坑。今天这一节，咱们就把最核心的几个概念捋清楚。

3.1 CUDA编程模型：说白了就是“主机+设备”

CUDA的编程模型，本质上是一个异构计算模型。什么叫异构？就是CPU和GPU各干各的，协同工作。

Host（主机）：指的是CPU及其内存（系统内存）。负责控制逻辑、串行任务、数据准备。
Device（设备）：指的是GPU及其显存（全局内存）。负责大规模并行计算。

我个人习惯把Host比作“项目经理”，Device比作“施工队”。项目经理把图纸（数据）准备好，告诉施工队怎么干，然后施工队一拥而上，快速干完。项目经理再回来收尾。

核心要点：CUDA程序从CPU启动，在GPU上执行并行计算，结果再传回CPU。这个流程，你写任何CUDA程序都绕不开。

这里我画了一张图，帮你直观理解这个模型：

3.2 Kernel Launch：启动GPU干活

在CUDA里，GPU上执行的函数叫Kernel。你写一个函数，前面加上 __global__ 修饰符，它就变成了一个Kernel。调用的时候，用三括号语法 <<<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];
    }
}

// 在主机端调用
int main() {
    // ... 分配内存、拷贝数据 ...
    int N = 1000000;
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // ... 拷贝结果回主机 ...
}

这里要注意，<<<grid, block>>> 里的两个参数，分别指定了网格和线程块的维度。我刚开始写的时候，经常搞混这两个参数。其实你记住：grid 是线程块的集合，block 是线程的集合。

避坑指南：我曾经在项目里，因为忘记检查 blocksPerGrid 的计算是否正确，导致数组尾部元素没有被处理。那个bug查了我整整一个下午。所以，一定要确保总线程数 >= 数据量，并且在Kernel里做边界检查。

3.3 Grid-Block-Thread层次：GPU的“组织架构”

CUDA的线程层次结构，是理解GPU并行的关键。说白了，就是GPU怎么组织成千上万个线程来干活。

Thread（线程）：最小的执行单元。每个线程执行Kernel函数中的同一段代码，但处理不同的数据。
Block（线程块）：一组线程的集合。同一个Block内的线程可以共享显存（Shared Memory），并且可以同步。
Grid（网格）：一组Block的集合。一个Kernel对应一个Grid。

你想想看，为什么要有这个层次？因为GPU的硬件架构就是这样的。一个Block内的线程，会被调度到同一个Streaming Multiprocessor（SM）上执行。这样它们才能高效地共享数据。

我画个图帮你理解：

在Kernel内部，我们可以通过内置变量来获取当前线程的位置：

内置变量	含义	示例值
`threadIdx.x`	当前线程在Block内的索引	0, 1, 2, ..., blockDim.x-1
`blockIdx.x`	当前Block在Grid内的索引	0, 1, 2, ..., gridDim.x-1
`blockDim.x`	每个Block包含的线程数	256, 512, 1024
`gridDim.x`	Grid包含的Block数	由用户指定

计算全局唯一线程ID的公式：int tid = blockIdx.x * blockDim.x + threadIdx.x;。这个公式你最好背下来，写Kernel时几乎每次都要用。

3.4 内存管理API：cudaMalloc 和 cudaFree

GPU不能直接访问CPU的内存。所以，我们需要在显存上分配空间，把数据拷过去，算完再拷回来。这就用到了CUDA的内存管理API。

最核心的两个函数：

cudaMalloc(void **devPtr, size_t size)：在显存上分配 size 字节的空间，返回指针。
cudaFree(void *devPtr)：释放显存空间。

配合 cudaMemcpy 使用，完成数据搬运：

float *h_A = (float*)malloc(N * sizeof(float)); // 主机内存
float *d_A = nullptr;

cudaMalloc(&d_A, N * sizeof(float));            // 显存分配
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice); // 拷入

// ... 执行Kernel ...

cudaMemcpy(h_A, d_A, N * sizeof(float), cudaMemcpyDeviceToHost); // 拷出
cudaFree(d_A);                                     // 释放显存
free(h_A);                                         // 释放主机内存

重要提醒：cudaMalloc 的第一个参数是 void**，也就是二级指针。我见过很多新手写成 cudaMalloc(d_A, ...)，结果程序直接崩溃。记住，要传指针的地址。

另外，每次调用 cudaMalloc 和 cudaMemcpy 后，最好检查一下返回值：

cudaError_t err = cudaMalloc(&d_A, N * sizeof(float));
if (err != cudaSuccess) {
    printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
    // 处理错误
}

嗯，这里要注意。在压测场景下，频繁地 cudaMalloc 和 cudaFree 是非常耗时的。这就是为什么后面我们要讲显存池——提前分配好一大块显存，重复利用，避免反复调用这些API。

个人经验：我在做某个推理服务优化时，发现每次请求都调用 cudaMalloc 分配临时缓冲区，导致延迟抖动很大。后来改用显存池预分配，P99 延迟直接降了 40%。所以，能复用就复用，别偷懒。

好了，CUDA编程基础就回顾到这里。这些内容虽然基础，但它们是后面所有优化技巧的基石。你把这些概念吃透了，后面讲显存池、流水线、异步操作时，才能跟得上节奏。

📚 课程目录

3、CUDA编程基础回顾：CUDA编程模型，Kernel Launch，Grid-Block-Thread层次，内存管理API

3.1 CUDA编程模型：说白了就是“主机+设备”

3.2 Kernel Launch：启动GPU干活

3.3 Grid-Block-Thread层次：GPU的“组织架构”

3.4 内存管理API：cudaMalloc 和 cudaFree