1. CUDA入门:GPU计算的历史与演进

说起GPU计算,我最早接触是在2008年。那时候CUDA刚发布不久,我还在做CPU上的数值模拟。一个偶然的机会,我发现用GPU算矩阵乘法,速度能快几十倍。说实话,当时我整个人都震惊了。

GPU最初是为了图形渲染设计的。你想啊,屏幕上几百万个像素点,每个都要计算颜色、光照、纹理,这天然就是并行计算的活儿。后来大家发现,这种大规模并行能力,用在科学计算上也很香。于是就有了GPGPU(通用GPU计算)这个概念。

NVIDIA在2006年推出了CUDA,算是正式把GPU计算的大门打开了。我个人觉得,CUDA最大的贡献不是硬件,而是让程序员能用C语言写GPU程序。你想想看,以前搞GPU编程得用图形API,那叫一个痛苦。

CUDA生态全景

现在的CUDA生态,已经远不止一个编程模型了。我把它分成几个层次:

  • 底层:CUDA Toolkit,包含编译器nvcc、运行时库、驱动API
  • 中间层:cuBLAS(线性代数)、cuFFT(傅里叶变换)、cuRAND(随机数)等库
  • 上层:深度学习框架(TensorFlow、PyTorch)、科学计算库(OpenMM、GROMACS)

我在项目中遇到过一个问题:直接用CUDA写卷积运算,性能还不如调cuDNN库。后来我学乖了,能用库就别自己造轮子。除非你确实需要定制化操作。

核心观点:CUDA生态的成熟度很高。90%的常见计算任务,都有现成的优化库。新手入门,先学会调库,再学写核函数。

GPU架构概览

要理解CUDA,得先搞明白GPU长什么样。我画了一张图,帮你理清思路。

GPU架构概览 GPU(图形处理器) GPC(图形处理簇) TPC(纹理处理簇) SM(流多处理器) 包含多个CUDA核心 TPC(纹理处理簇) SM(流多处理器) 包含多个CUDA核心 SM内部结构 Warp Warp Warp Warp GPC(图形处理簇) TPC SM TPC SM 显存(Global Memory) 所有SM共享,容量大但延迟高 通常为GDDR6/HBM2e SM Warp(32线程) 显存

这张图展示了GPU的核心层次结构。最底层是SM(流多处理器),它是GPU执行计算的基本单元。每个SM里有多个CUDA核心,还有共享内存、寄存器等资源。

Warp是NVIDIA引入的一个概念。一个Warp包含32个线程,它们在同一时刻执行同一条指令。说白了,这就是GPU的"线程束"。我刚开始学的时候,总觉得Warp这个概念很绕。后来我把它理解成"一个小组,32个人一起干活",就好懂了。

SP(流处理器)其实就是CUDA核心。每个SP能执行一个浮点运算。现在的安培架构,一个SM里可以有128个SP。你算算,一个GPU几十个SM,那就是几千个核心同时干活。

小技巧:写CUDA代码时,要时刻想着Warp。如果一个Warp里的线程分支了(if-else),性能会下降。我踩过这个坑,后来写代码都尽量让Warp内线程走同一路径。

CUDA安装与开发环境搭建

搭建环境这事儿,说简单也简单,说坑也多。我按步骤来:

  1. 装NVIDIA驱动:去NVIDIA官网下载对应显卡的驱动。别用Windows自动更新的,那个版本太老。
  2. 装CUDA Toolkit:我建议用11.x版本,兼容性好。别追最新版,除非你爱折腾。
  3. 装Nsight:这是NVIDIA的IDE,基于Eclipse。调试GPU代码必备。

注意:我曾经因为驱动版本和Toolkit版本不匹配,折腾了一整天。后来我养成了习惯:装之前先查兼容性矩阵。NVIDIA官网有张表,告诉你哪个驱动配哪个Toolkit。

验证安装是否成功,打开命令行,输入:

nvcc --version
nvidia-smi

如果能看到版本信息,说明环境搭好了。nvidia-smi还能看到GPU温度、显存使用情况,调试时很有用。

第一个CUDA程序:Hello World

终于到写代码的环节了。CUDA的Hello World,不是打印一句话那么简单。我们要在GPU上执行一个核函数。

#include <stdio.h>

// 核函数:在GPU上执行
__global__ void helloFromGPU() {
    printf("Hello World from GPU!\n");
}

int main() {
    // 调用核函数,启动1个block,每个block里1个线程
    helloFromGPU<<<1, 1>>>();
    
    // 等待GPU执行完毕
    cudaDeviceSynchronize();
    
    printf("Hello World from CPU!\n");
    return 0;
}

编译命令:

nvcc hello.cu -o hello
./hello

输出结果:

Hello World from GPU!
Hello World from CPU!

这里有几个关键点:

  • __global__ 修饰符告诉编译器,这个函数在GPU上执行,但从CPU调用
  • <<<1, 1>>> 是执行配置,第一个参数是block数,第二个是每个block的线程数
  • cudaDeviceSynchronize() 确保GPU执行完,不然CPU可能先退出

我第一次跑通这个程序时,心里还挺激动的。虽然只是打印一句话,但背后是CPU和GPU协同工作的完整流程。

总结一下:CUDA编程的核心就是写核函数,然后在CPU端调用它。GPU负责并行计算,CPU负责控制流程。这个分工,贯穿了整个CUDA编程模型。

嗯,第一章就到这里。代码不多,但概念很重要。后面我们会一步步深入,从线程层次结构到共享内存优化,再到多GPU编程。路还长,慢慢来。

可可精品资料库