第一章: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里就是:
- 数据准备: CPU在内存里准备好数据。
- 数据搬运: 通过PCIe总线,把数据从CPU内存拷贝到GPU显存(全局内存)。
- 启动Kernel: CPU告诉GPU:“启动这个函数,用N个线程去执行。”
- GPU执行: GPU的SM调度Warp,并行执行Kernel。
- 结果回传: GPU计算完成后,再把结果从显存拷贝回CPU内存。
避坑指南: 我曾经犯过一个错误——在Kernel启动后立刻去读结果。GPU是异步执行的,CPU发完指令后不会等它。如果你立刻去读,读到的可能是旧数据。正确的做法是调用 cudaDeviceSynchronize() 等待GPU完成。记住:CPU和GPU是异步的,同步点要自己控制。
另外,PCIe带宽是个硬伤。我记得有一次做视频处理,数据量很大,每次都要在CPU和GPU之间来回拷贝,结果大部分时间都花在传输上了。后来我改用“流水线”技术——把数据分块,一边传输一边计算,才把性能提上来。这个我们后面会详细讲。
小结
这一章我们聊了三个核心概念:
- SM与CUDA核心: 理解Warp的执行方式,避免分支分化。
- 内存层次: 善用共享内存,避免寄存器溢出。
- CPU-GPU协同: 注意异步执行和PCIe传输开销。
嗯,这些都是基础中的基础。但你别小看它们,我见过太多人因为不理解这些,写出来的程序跑得比CPU还慢。下一章,我们会深入聊聊“内存分配与数据传输”的那些坑。到时候见。