第一章:GPU架构基础:CUDA核心与SM概念、内存层次结构概览、GPU与CPU协同工作原理

各位同学,欢迎来到《英伟达GPU内存管理与性能调优实战》的第一课。

说实话,我刚开始接触GPU编程那会儿,也踩过不少坑。很多人上来就写kernel,结果性能一塌糊涂,根本不知道问题出在哪。说白了,你不了解GPU的“脾气”,它就不会好好给你干活。这一章,我们就先把地基打牢。

1.1 CUDA核心与SM:GPU的“工人”与“车间”

先问大家一个问题:GPU里那么多所谓的“CUDA核心”,它们到底是怎么工作的?

我个人习惯把CUDA核心想象成一个“流水线工人”。它很单纯,只会做算术(加减乘除)。但一个工人成不了事,得一群人聚在一起干活。这个“一群人”的集合,就是SM(Streaming Multiprocessor,流式多处理器)。

SM才是真正的执行单元。 每个SM里包含几十到上百个CUDA核心,还有共享内存、寄存器、调度器等等。你可以把SM理解成一个“车间”,车间里有几十个工人(CUDA核心),大家共用一套工具(共享内存)和一位工头(Warp调度器)。

核心概念: 程序员写的是线程(Thread),硬件执行的是Warp(线程束)。一个Warp包含32个线程,它们在同一时刻执行同一条指令。这是理解GPU性能的起点。

我在项目中遇到过一个问题:一个kernel里有很多if-else分支。结果发现,同一个Warp里的32个线程,有的走if,有的走else。GPU没办法,只能先执行if的线程,再执行else的线程。这就是所谓的“Warp divergence”(线程束分化),性能直接腰斩。嗯,这里要注意,写代码时尽量让同一个Warp里的线程走相同的路径。

1.2 内存层次结构:从“手边”到“云端”

GPU的内存,说白了就是一个“金字塔”结构。越靠近计算单元,速度越快,但容量越小。我经常跟团队说:性能调优,本质上就是数据在哪一层的问题。

内存类型 位置 速度 容量 作用域
寄存器 (Register) SM内部 最快 (~1 cycle) 极小 (每个线程几十个) 单个线程
共享内存 (Shared Memory) SM内部 很快 (~5-30 cycles) 小 (几十KB) 单个Block内所有线程
L1/L2缓存 芯片上 中等 几MB 全局
全局内存 (Global Memory) 显存 (VRAM) 慢 (~400-800 cycles) 大 (几GB到几十GB) 所有线程 + Host
本地内存 (Local Memory) 显存 (VRAM) 自动分配 单个线程(寄存器溢出时)

你想想看,如果每个线程都去全局内存里读数据,那GPU大部分时间都在“等数据”,而不是“算数据”。这就是所谓的“内存带宽瓶颈”。

我的经验: 我曾经优化过一个矩阵乘法的例子。一开始直接读写全局内存,带宽利用率只有20%。后来我把数据分块,先加载到共享内存里,让同一个Block的线程反复使用,带宽利用率直接飙到80%以上。说白了,共享内存就是你的“缓存加速器”

还有一个坑:本地内存。当你的kernel里定义的局部变量太多,寄存器放不下时,编译器会自动把它们“溢出”到本地内存。本地内存物理上在显存里,速度很慢。我曾经因为一个kernel里定义了一个大数组,结果性能惨不忍睹。查了半天才发现是寄存器溢出导致的。

1.3 GPU与CPU协同:老板与工人的关系

CPU和GPU怎么配合?我打个比方:CPU是老板,GPU是工人。

老板(CPU)负责逻辑判断、任务调度、串行计算。工人(GPU)负责大规模、重复性的并行计算。老板把任务分好,告诉工人:“你去把这100万个数的平方算出来。”然后工人就闷头干活,干完了通知老板:“活干完了,结果在这。”

这个流程在CUDA里就是:

  1. 数据准备: CPU在内存里准备好数据。
  2. 数据搬运: 通过PCIe总线,把数据从CPU内存拷贝到GPU显存(全局内存)。
  3. 启动Kernel: CPU告诉GPU:“启动这个函数,用N个线程去执行。”
  4. GPU执行: GPU的SM调度Warp,并行执行Kernel。
  5. 结果回传: GPU计算完成后,再把结果从显存拷贝回CPU内存。

避坑指南: 我曾经犯过一个错误——在Kernel启动后立刻去读结果。GPU是异步执行的,CPU发完指令后不会等它。如果你立刻去读,读到的可能是旧数据。正确的做法是调用 cudaDeviceSynchronize() 等待GPU完成。记住:CPU和GPU是异步的,同步点要自己控制

另外,PCIe带宽是个硬伤。我记得有一次做视频处理,数据量很大,每次都要在CPU和GPU之间来回拷贝,结果大部分时间都花在传输上了。后来我改用“流水线”技术——把数据分块,一边传输一边计算,才把性能提上来。这个我们后面会详细讲。

小结

这一章我们聊了三个核心概念:

  • SM与CUDA核心: 理解Warp的执行方式,避免分支分化。
  • 内存层次: 善用共享内存,避免寄存器溢出。
  • CPU-GPU协同: 注意异步执行和PCIe传输开销。

嗯,这些都是基础中的基础。但你别小看它们,我见过太多人因为不理解这些,写出来的程序跑得比CPU还慢。下一章,我们会深入聊聊“内存分配与数据传输”的那些坑。到时候见。