第1章 CUDA基础回顾

第一章：CUDA基础回顾

各位同学，咱们开始上课。今天聊的是CUDA基础，这部分内容我估计不少人都接触过，但说实话，很多人在实际项目中栽跟头，往往就是基础没打牢。我自己带团队时，发现性能瓶颈排查到最后，十有八九都能追溯到对GPU架构或者内存模型的理解偏差上。

所以这一章，咱们不赶进度，把地基夯实了再说。

1.1 GPU架构概览

先聊聊GPU长什么样。你想想看，CPU和GPU最大的区别是什么？

CPU是「精兵强将」——核心少，但每个核心都极其复杂，有庞大的缓存、分支预测、乱序执行单元。它擅长处理串行任务，逻辑复杂多变的那种。

GPU呢？它是「人海战术」——核心多到吓人，但每个核心都很简单。它擅长做大量相同或相似的计算，也就是数据并行。

我打个比方。CPU就像几个顶尖大厨，能做满汉全席。GPU就像几百个帮厨，专门负责切土豆丝，一秒钟能切一吨。你说谁更适合做图像渲染、矩阵乘法？显然是后者。

具体到NVIDIA的GPU架构，这些年经历了Fermi、Kepler、Maxwell、Pascal、Volta、Turing、Ampere、Hopper……每一代都有改进。但核心思想没变：

SM（Streaming Multiprocessor）：这是GPU的基本计算单元。你可以把它理解成一个「小分队」，里面包含若干CUDA核心、共享内存、寄存器文件等。
Warp：32个线程组成一个warp，这是GPU调度和执行的基本单位。注意，不是单个线程，是32个一捆。
全局内存：显存，所有线程都能访问，但延迟很高。
共享内存：每个SM内部的快速存储，同一个block的线程可以共享。

重点记住：GPU的强项是「大量线程并行执行相同指令」。如果程序分支太多，导致warp内线程走不同路径，性能就会暴跌。这就是所谓的「warp divergence」。

我记得有一次帮客户优化一个流体力学模拟程序，他们用了大量if-else判断边界条件。结果一profiling发现，warp执行效率只有30%多。后来改成用位运算和查表法，效率直接拉到90%以上。嗯，这就是架构理解带来的收益。

1.2 CUDA编程模型

CUDA编程模型，说白了就是「怎么指挥GPU干活」。它遵循一个经典模式：

Host（CPU）负责控制流程，把数据准备好。
Device（GPU）负责大规模并行计算。
数据从Host拷贝到Device，计算完再拷回来。

代码结构大概是这样的：

// 1. 分配主机和设备内存
float *h_data, *d_data;
h_data = (float*)malloc(N * sizeof(float));
cudaMalloc(&d_data, N * sizeof(float));

// 2. 从主机拷贝到设备
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);

// 3. 调用核函数
myKernel<<<gridDim, blockDim>>>(d_data, N);

// 4. 从设备拷贝回主机
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);

// 5. 释放内存
cudaFree(d_data);
free(h_data);

这里有个关键点：<<<gridDim, blockDim>>> 这个语法，很多人一开始觉得怪。其实它就是在告诉GPU：「我要启动这么多线程块（gridDim），每个块里放这么多线程（blockDim）」。

个人习惯：我写CUDA代码时，喜欢把gridDim和blockDim定义成dim3类型，这样能清晰表达三维索引。虽然大部分场景只用一维或二维，但养成好习惯没坏处。

1.3 线程层次结构

CUDA的线程组织方式，是三层结构：

Thread（线程）：最小编程单元。
Block（线程块）：一组线程，可以协作，共享共享内存。
Grid（网格）：一组线程块，构成整个核函数。

每个线程都有唯一的ID，通过 threadIdx、blockIdx、blockDim、gridDim 这些内置变量来定位。

举个例子，假设你要处理一张1024x1024的图像，每个像素一个线程：

dim3 blockDim(32, 32);   // 每个块32x32=1024个线程
dim3 gridDim(32, 32);    // 32x32=1024个块
// 总线程数 = 1024 * 1024 = 1,048,576

__global__ void processImage(float *img, int width, int height) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < width && y < height) {
        // 处理像素 (x, y)
    }
}

为什么会这样设计？说白了，就是为了灵活性和可扩展性。同一个核函数，在GTX 1060上跑和在A100上跑，只要调整gridDim就行，代码不用改。

避坑指南：我曾经犯过一个低级错误——blockDim设得太大，超过了SM能同时容纳的线程数上限。结果核函数启动失败，报错说「invalid configuration argument」。后来查文档才知道，每个block的线程数不能超过1024（老架构是512）。所以，blockDim.x * blockDim.y * blockDim.z <= 1024，这是硬约束。

1.4 内存层次结构

内存这块，是CUDA优化的重中之重。我见过太多人，核函数写得飞快，结果数据搬运占了大头，白忙活。

GPU的内存层次，从快到慢、从近到远：

内存类型	位置	访问速度	作用域	生命周期
寄存器	SM内部	~1 cycle	单个线程	核函数执行期间
共享内存	SM内部	~5 cycles	同一个block	block生命周期
L1缓存	SM内部	~10 cycles	同一个SM	自动管理
L2缓存	芯片全局	~100 cycles	所有SM共享	自动管理
全局内存	显存	~400-800 cycles	所有线程	程序运行期间
本地内存	显存	~400-800 cycles	单个线程	核函数执行期间

看到没？全局内存比寄存器慢了三个数量级。所以优化的核心思路就一条：尽量把数据往快了放。

具体怎么做？

用共享内存做数据复用：比如矩阵乘法，把子矩阵加载到共享内存里，让block内所有线程反复使用，避免反复访问全局内存。
合并访问：全局内存访问时，尽量让连续线程访问连续地址。这样硬件会把多次访问合并成一次大传输，效率翻倍。
减少寄存器溢出：每个线程用的寄存器太多，编译器会把多余的变量挤到本地内存（其实还是显存），速度骤降。我一般会控制每个线程的寄存器用量，不超过32个。

一个小技巧：用 __shared__ 声明共享内存变量时，可以动态指定大小。比如 extern __shared__ float sdata[];，然后在核函数启动时用第三个参数传大小：kernel<<<grid, block, sharedMemSize>>>()。这样同一个核函数可以适配不同大小的数据块，很灵活。

嗯，内存这块内容很多，咱们后面讲TensorRT调优时还会反复提到。现在先有个整体印象就行。

好，第一章就到这里。下一章咱们聊CUDA流和事件，以及怎么用它们来隐藏数据传输延迟。到时候我会分享一个我实际项目中用多流把推理吞吐量翻倍的案例，挺有意思的。

📚 课程目录

第一章：CUDA基础回顾

1.1 GPU架构概览

1.2 CUDA编程模型

1.3 线程层次结构

1.4 内存层次结构