📚 课程目录第4章 CUDA编程基础

4. CUDA编程基础

说实话，我第一次接触CUDA的时候，心里是有点发怵的。那时候我刚从CPU编程转到GPU，满脑子都是"这玩意儿到底怎么并行？"后来踩了不少坑，才慢慢摸清门道。今天我就把这些年积累的经验，跟你好好聊聊。

4.1 CUDA编程模型：主机与设备

CUDA的编程模型，说白了就是"CPU+GPU"协同工作。CPU叫主机（Host），GPU叫设备（Device）。我习惯把CPU想象成项目经理，GPU就是一群埋头苦干的程序员。

典型的CUDA程序流程是这样的：

CPU把数据准备好，拷贝到GPU显存
CPU告诉GPU："来，帮我算这个"
GPU吭哧吭哧算完
CPU再把结果拷回来

嗯，这里要注意：数据拷贝是开销最大的操作。我在项目中遇到过，有人把每次迭代的数据都来回拷贝，结果性能还不如CPU。你想想看，这多冤啊。

  核心原则：尽量减少CPU和GPU之间的数据传输。能一次拷完就别分多次。

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

这是CUDA最核心的概念，也是新手最容易懵的地方。我刚开始学的时候，Grid和Block的关系绕了好久才搞明白。

简单来说：

Thread（线程）：最基础的计算单元，每个线程执行相同的核函数
Block（线程块）：一组线程的集合，可以互相协作
Grid（网格）：一组Block的集合，构成整个计算任务

你可以把GPU想象成一个工厂：

Grid = 整个工厂
Block = 车间
Thread = 工人

每个工人（Thread）干一样的活，但处理不同的数据。这就是所谓的SIMT（单指令多线程）模型。

我的经验：Block的大小一般设为32的倍数，因为GPU的warp（线程束）大小就是32。设成128或256通常效果不错。

来看个具体的例子：

// 定义Grid和Block的维度
dim3 gridDim(2, 2);      // 2x2的网格，共4个Block
dim3 blockDim(16, 16);   // 16x16的Block，共256个线程

// 启动核函数
myKernel<<<gridDim, blockDim>>>(d_data, N);

这段代码启动了4个Block，每个Block有256个线程，总共1024个线程并行执行。每个线程可以通过内置变量知道自己是谁：

threadIdx.x：线程在Block内的索引
blockIdx.x：Block在Grid内的索引
blockDim.x：Block的维度大小
gridDim.x：Grid的维度大小

我经常用这个公式计算全局线程ID：

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

这个公式你最好记牢了，几乎每个CUDA程序都会用到。

4.3 内存层次结构：Global、Shared、Local

GPU的内存体系比CPU复杂得多。我刚开始做项目时，就因为搞混了内存类型，性能差了10倍不止。

先看这张图，帮你理清思路：

三种内存各有特点，我总结一下：

内存类型	容量	速度	作用域	典型用途
Global Memory	GB级	慢（~400 cycles）	所有线程	存储数据、模型参数
Shared Memory	几十KB	快（~20 cycles）	同一Block	数据复用、线程协作
Local Memory	寄存器溢出时	同Global	单个线程	自动使用，尽量避免

避坑指南：我曾经在训练一个大型模型时，发现性能远低于预期。排查了半天，原来是某个中间变量太大，导致寄存器溢出，数据被放到了Local Memory里。Local Memory的访问速度和Global一样慢，但很多人意识不到这一点。解决办法是：要么减少每个线程的变量数量，要么手动使用Shared Memory。

4.4 CUDA核函数编写入门

核函数（Kernel Function）就是在GPU上执行的函数。写法上有个标志性的关键字——__global__。

来看一个最简单的例子：向量加法。

// 核函数定义
__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 = 1 << 20;  // 约100万个元素
    float *d_A, *d_B, *d_C;
    
    // 在GPU上分配内存
    cudaMalloc(&d_A, N * sizeof(float));
    cudaMalloc(&d_B, N * sizeof(float));
    cudaMalloc(&d_C, N * sizeof(float));
    
    // 拷贝数据到GPU
    cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);
    
    // 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 拷贝结果回CPU
    cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);
    
    // 释放GPU内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    
    return 0;
}

这段代码有几个关键点：

__global__：告诉编译器这是个核函数，在GPU上执行
<<<grid, block>>>：配置启动参数，指定线程组织方式
cudaMalloc/cudaMemcpy：显存管理和数据传输
if (i < N)：边界检查，防止越界

我的习惯：写核函数时，我总会加边界检查。虽然理论上线程数可以整除数据量，但实际中总会有边界情况。这个习惯帮我避免了好几次段错误。

核函数有几个限制要记住：

不能调用C++标准库函数（除了少数数学函数）
不能使用静态变量
不能使用变长参数
返回值必须是void

说白了，核函数就是"轻装上阵"。它只负责计算，不负责管理资源。

最后分享一个调试技巧：我经常用cudaGetLastError()来检查核函数是否启动成功。如果返回cudaSuccess，说明启动没问题。但要注意，这个函数只能检查启动错误，不能检查运行时的逻辑错误。

vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("核函数启动失败: %s\n", cudaGetErrorString(err));
}

嗯，CUDA编程的基础就这些。说白了就是三件事：理解线程怎么组织、搞清内存怎么用、写好核函数。把这些搞明白了，后面的优化技巧才能派上用场。