第一章: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):一组线程块,覆盖整个计算任务

我画了一张图,帮你理解这个层次关系:

Grid(网格) Block (0,0) Block (1,0) Block (2,0) T(0,0) T(1,0) T(0,1) T(1,1) T(0,2) T(1,2) T(0,3) T(1,3) 共享内存(Shared Memory) 每个Block内部 线程可协作

为什么要有这三个层次?原因很简单:

  • 线程:处理一个数据元素,比如一个像素、一个矩阵元素
  • 线程块:一组线程可以共享数据(通过共享内存),减少全局内存访问
  • 网格:覆盖整个数据集,比如一张图片、一个矩阵

写一个最简单的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开发环境搭建——别让环境卡住你

说实话,我见过太多人卡在环境搭建这一步。我自己也踩过不少坑,今天一次性帮你避开。

你需要装三样东西:

  1. NVIDIA驱动:让操作系统认识你的GPU
  2. CUDA Toolkit:包含编译器(nvcc)、库文件、调试工具
  3. 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编程模型的三个层次:线程、线程块、网格
  • 开发环境的搭建步骤和常见坑

说实话,这些内容看起来简单,但它们是后面所有优化的基础。我见过太多人跳过这些直接写核函数,结果遇到性能问题完全不知道从哪下手。

下一章,我们会深入内存层次结构——为什么全局内存那么慢?共享内存怎么用?寄存器不够怎么办?这些才是真正拉开性能差距的地方。

课后练习:

  1. 运行deviceQuery程序,记录你的GPU的SM数量、最大线程数、共享内存大小
  2. 修改上面的向量加法代码,改成矩阵加法(二维索引)
  3. 用nvidia-smi监控GPU利用率,看看你的程序占用了多少资源

可可精品资料库