第3章 CUDA编程基础

3. CUDA编程基础：Host与Device代码编写、内存管理、Kernel函数定义与启动、向量加法实战

好，咱们今天聊聊CUDA编程最核心的那点事儿。说白了，就是搞清楚CPU和GPU怎么分工，数据怎么搬，以及怎么让GPU真正跑起来干活。

我记得刚接触CUDA那会儿，最头疼的就是“Host”和“Device”这两个概念。Host就是CPU及其内存，Device就是GPU及其显存。你想想看，它们俩物理上是分开的，不能直接访问对方的内存。所以，写CUDA程序，本质上就是在写两套代码：一套跑在CPU上，负责调度和数据搬运；另一套跑在GPU上，负责大规模并行计算。

3.1 Host与Device：CPU和GPU的分工

我个人习惯把Host代码看作是“项目经理”，Device代码看作是“一线工人”。项目经理不干具体的活，但他知道活怎么分、资源怎么调。一线工人只管闷头执行，效率极高。

Host代码：在CPU上执行，负责逻辑控制、串行任务、数据准备、调用Kernel。用标准的C/C++编译器（如nvcc）编译。
Device代码：在GPU上执行，负责高度并行的计算任务。用CUDA C/C++编写，通过nvcc编译，其中包含__global__、__device__等关键字。

核心要点：Host和Device的内存是隔离的。你不能在CPU代码里直接访问GPU的指针，反之亦然。所有数据交换都必须通过显式的API调用完成。

3.2 内存管理：cudaMalloc 与 cudaMemcpy

数据搬运是CUDA编程里最容易出错的环节。我在项目中遇到过好几次，程序跑起来没报错，但结果就是不对。查了半天，发现是cudaMemcpy的方向搞反了，或者大小算错了。

CUDA提供了几个关键API来管理设备内存：

API函数	功能	注意事项
`cudaMalloc`	在GPU显存上分配内存	返回的是设备指针，不能在Host端直接解引用
`cudaMemcpy`	在Host和Device之间拷贝数据	需要指定拷贝方向（HostToDevice, DeviceToHost, DeviceToDevice）
`cudaFree`	释放GPU显存	必须与cudaMalloc成对使用，否则内存泄漏

来看一个典型的内存管理流程：

// 1. 在Host端分配内存
float *h_a = (float*)malloc(N * sizeof(float));
float *h_b = (float*)malloc(N * sizeof(float));
float *h_c = (float*)malloc(N * sizeof(float));

// 2. 在Device端分配内存
float *d_a, *d_b, *d_c;
cudaMalloc((void**)&d_a, N * sizeof(float));
cudaMalloc((void**)&d_b, N * sizeof(float));
cudaMalloc((void**)&d_c, N * sizeof(float));

// 3. 将数据从Host拷贝到Device
cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);

// 4. 调用Kernel（稍后讲）

// 5. 将结果从Device拷贝回Host
cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);

// 6. 释放内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(h_a);
free(h_b);
free(h_c);

避坑指南：我曾经犯过一个低级错误——忘记检查cudaMalloc的返回值。当显存不足时，cudaMalloc会返回错误，但程序不会崩溃，后续的cudaMemcpy会访问非法地址，导致各种诡异问题。所以，每次调用CUDA API后，最好都检查一下返回值，或者用cudaGetLastError()捕获错误。

3.3 Kernel函数定义与启动

Kernel函数就是跑在GPU上的那个“一线工人”。它的定义和普通C函数很像，但有几个关键区别：

用__global__修饰符声明，表示这是一个在Device上执行、从Host端调用的函数。
返回类型必须是void。
不能使用可变参数（如...）。
可以调用__device__函数（在Device上执行、从Device端调用的函数）。

启动Kernel的语法也很特别：

kernel_name<<<gridDim, blockDim>>>(arguments);

这里的<<<gridDim, blockDim>>>就是传说中的“执行配置”。gridDim指定了有多少个线程块（Block），blockDim指定了每个线程块里有多少个线程（Thread）。

嗯，这里要注意：gridDim和blockDim可以是dim3类型，支持一维、二维、三维的线程组织方式。我个人习惯用一维来处理数组，用二维来处理图像或矩阵。

每个线程在Kernel内部可以通过内置变量知道自己的身份：

threadIdx.x：线程在块内的索引
blockIdx.x：块在网格内的索引
blockDim.x：每个块的线程数
gridDim.x：网格内的块数

所以，一个线程的全局唯一ID可以这样计算：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

3.4 向量加法实战

理论说再多，不如动手写一个。向量加法就是CUDA界的“Hello World”。咱们来一步步实现它。

先定义Kernel函数：

__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];
    }
}

注意那个if (i < n)的判断。为什么要有它？因为线程总数（gridDim.x * blockDim.x）可能大于数组长度n。如果不加这个判断，多余的线程会访问越界内存，导致段错误或者计算结果错误。

然后是在Host端调用：

int n = 1000000;
size_t size = n * sizeof(float);

// ... 分配内存、初始化数据、拷贝到Device ...

// 设置执行配置
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

// 启动Kernel
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

// 同步等待Kernel执行完成
cudaDeviceSynchronize();

// ... 将结果拷贝回Host、验证结果、释放内存 ...

这里blocksPerGrid的计算方式很常见：(n + threadsPerBlock - 1) / threadsPerBlock，确保有足够的线程覆盖所有数据元素。

个人经验：线程块大小（threadsPerBlock）的选择有讲究。我一般选128或256，这是大多数GPU的warp大小（32）的整数倍，能保证线程束的充分利用。太小了会导致硬件利用率低，太大了又可能受限于寄存器数量。你可以用cudaOccupancyMaxPotentialBlockSize这个API来帮你算最优值。

最后，别忘了同步。Kernel启动是异步的，Host代码会立即继续执行。如果你在Kernel还没完成时就读取结果，拿到的就是垃圾数据。cudaDeviceSynchronize()会阻塞Host，直到所有Device操作完成。

好了，向量加法就这么简单。但你别小看它，这里面包含了CUDA编程的所有核心要素：Host/Device分工、内存管理、Kernel定义与启动、线程索引计算。把这个吃透了，后面的复杂算法都是在这个基础上堆起来的。

这张图把整个流程串起来了。你写CUDA程序时，脑子里就要有这个图：数据从Host搬到Device，GPU并行处理，结果再搬回来。每一步都对应着具体的API调用，缺一不可。

📚 课程目录

3. CUDA编程基础：Host与Device代码编写、内存管理、Kernel函数定义与启动、向量加法实战

3.1 Host与Device：CPU和GPU的分工

3.2 内存管理：cudaMalloc 与 cudaMemcpy

3.3 Kernel函数定义与启动

3.4 向量加法实战