第一章:CUDA入门——从GPU架构到第一个核函数
各位同学好,我是你们这门课的主讲。在AI系统优化这个领域摸爬滚打了十几年,我见过太多人一上来就写CUDA代码,结果性能惨不忍睹。说白了,不理解硬件就去写代码,就像闭着眼睛开车——能跑,但迟早要翻。
这一章,我们先打好地基。我会带你看看GPU架构是怎么一路进化过来的,再聊聊CUDA编程模型的核心思想,最后手把手教你搭好开发环境。嗯,这些内容看起来基础,但相信我,后面30章的所有优化技巧,都建立在这些理解之上。
本章核心目标:理解GPU为什么快、CUDA怎么组织计算、以及如何让你的电脑准备好跑CUDA程序。
1.1 GPU架构演进史——从图形卡到通用计算怪兽
先问个问题:为什么GPU能比CPU快那么多?
我刚开始做GPU计算时,也困惑过这个问题。后来我明白了——CPU是「全能选手」,什么活都能干,但一次只能干几件事。GPU是「专业打手」,只会干一件事,但一次能干几千件。
GPU架构的演进,说白了就是「从画图到算数」的转变。
| 年代 | 架构代号 | 核心变化 | 我的评价 |
|---|---|---|---|
| 2006 | G80(Tesla) | 首次支持CUDA,统一着色器 | 里程碑,但当时没人看好 |
| 2010 | Fermi | 引入L1/L2缓存,ECC内存 | 真正可用的通用计算架构 |
| 2012 | Kepler | 动态并行,Hyper-Q | 我项目里用得最多的架构 |
| 2014 | Maxwell | 改进共享内存,降低功耗 | 能效比大幅提升 |
| 2017 | Volta | Tensor Core,独立线程调度 | AI计算的转折点 |
| 2020 | Ampere | 第三代Tensor Core,MIG | 现在的主流,性能翻倍 |
| 2022 | Hopper | Transformer Engine,DPX指令 | 专为大模型设计 |
这里有个关键概念——SM(Streaming Multiprocessor)。你可以把SM想象成一个「小工厂」,里面有很多计算单元(CUDA Core)。GPU就是由几十个这样的「小工厂」组成的「超级工厂」。
我个人习惯把SM比作「车间主任」——它负责调度线程、分配资源、管理内存。你写的每个CUDA线程,最终都会跑到某个SM上执行。
避坑指南:我曾经以为GPU架构越新越好,结果在Kepler上写的代码,放到Fermi上跑直接报错。不同架构的寄存器数量、共享内存大小都不一样,写代码时一定要查目标架构的规格表。
1.2 CUDA编程模型概述——你只需要理解三个层次
CUDA编程模型,说白了就是「怎么组织你的计算任务」。我见过很多新手一上来就写__global__ void kernel(),但连grid和block是什么都没搞明白。
你想想看,GPU有几千个核心,你总得告诉它「谁干什么活」吧?CUDA用三个层次来组织:
- 线程(Thread):最小的执行单元,每个线程执行相同的代码,但处理不同的数据
- 线程块(Block):一组线程,可以协作、共享数据
- 网格(Grid):一组线程块,覆盖整个计算任务
我画了一张图,帮你理解这个层次关系:
为什么要有这三个层次?原因很简单:
- 线程:处理一个数据元素,比如一个像素、一个矩阵元素
- 线程块:一组线程可以共享数据(通过共享内存),减少全局内存访问
- 网格:覆盖整个数据集,比如一张图片、一个矩阵
写一个最简单的CUDA程序,只需要三步:
// 1. 定义核函数
__global__ void vec_add(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// 2. 配置执行参数
int threads_per_block = 256;
int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
// 3. 启动核函数
vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
注意:核函数不能返回值,不能使用C++标准库(如iostream),不能调用虚函数。这些限制都是因为GPU的执行环境跟CPU完全不同。
1.3 CUDA开发环境搭建——别让环境卡住你
说实话,我见过太多人卡在环境搭建这一步。我自己也踩过不少坑,今天一次性帮你避开。
你需要装三样东西:
- NVIDIA驱动:让操作系统认识你的GPU
- CUDA Toolkit:包含编译器(nvcc)、库文件、调试工具
- Nsight工具链:性能分析和调试工具
我的建议是:先装驱动,再装Toolkit,最后装Nsight。顺序错了容易出问题。
第一步:检查你的GPU
打开终端,运行:
nvidia-smi
如果能看到GPU信息,说明驱动已经装好了。如果报错,去NVIDIA官网下载对应驱动。
个人经验:我曾经在Ubuntu上装驱动,结果黑屏了。后来发现是Secure Boot没关。如果你也遇到类似问题,进BIOS关掉Secure Boot,或者给驱动签名。
第二步:安装CUDA Toolkit
去NVIDIA官网下载对应版本。我个人推荐用runfile安装,比deb包更可控。
# 下载runfile
wget https://developer.download.nvidia.com/compute/cuda/12.3.0/local_installers/cuda_12.3.0_xxx_linux.run
# 安装(注意不要装驱动,只装Toolkit)
sudo sh cuda_12.3.0_xxx_linux.run --toolkit --silent
装完后配置环境变量:
export PATH=/usr/local/cuda-12.3/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-12.3/lib64:$LD_LIBRARY_PATH
第三步:验证安装
写一个简单的程序测试:
#include <cuda_runtime.h>
#include <stdio.h>
int main() {
int device_count;
cudaGetDeviceCount(&device_count);
printf("找到 %d 个GPU设备\n", device_count);
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("设备名称: %s\n", prop.name);
printf("计算能力: %d.%d\n", prop.major, prop.minor);
printf("SM数量: %d\n", prop.multiProcessorCount);
return 0;
}
编译运行:
nvcc -o device_query device_query.cu
./device_query
如果能看到GPU信息,恭喜你,环境搭好了。
第四步:安装Nsight工具链
Nsight包括:
- Nsight Systems:系统级性能分析,看CPU/GPU时间线
- Nsight Compute:核函数级性能分析,看指令、带宽、延迟
- Nsight Eclipse Edition:IDE集成(我基本不用,习惯命令行)
安装很简单:
# 通过CUDA Toolkit自带的安装器
sudo apt install nsight-systems nsight-compute
避坑指南:我曾经在Windows上装Nsight,结果VS插件死活不识别。后来发现是VS版本太新,Nsight还没适配。建议用VS2019或VS2022的特定版本。
本章小结
这一章我们聊了:
- GPU架构从G80到Hopper的演进,核心是SM的进化
- CUDA编程模型的三个层次:线程、线程块、网格
- 开发环境的搭建步骤和常见坑
说实话,这些内容看起来简单,但它们是后面所有优化的基础。我见过太多人跳过这些直接写核函数,结果遇到性能问题完全不知道从哪下手。
下一章,我们会深入内存层次结构——为什么全局内存那么慢?共享内存怎么用?寄存器不够怎么办?这些才是真正拉开性能差距的地方。
课后练习:
- 运行
deviceQuery程序,记录你的GPU的SM数量、最大线程数、共享内存大小 - 修改上面的向量加法代码,改成矩阵加法(二维索引)
- 用
nvidia-smi监控GPU利用率,看看你的程序占用了多少资源