2. CUDA内存管理基础:cudaMalloc/cudaFree机制、CUDA虚拟内存与物理内存、页表与映射
大家好,我是你们的老朋友。今天咱们来聊聊CUDA内存管理。说实话,这玩意儿是算子优化的地基。地基没打好,上面盖多高的楼都得塌。我见过太多同学,算子写得很漂亮,一跑就爆显存,或者慢得像蜗牛。说白了,就是没搞懂GPU是怎么管内存的。
咱们这一章,就专门把CUDA内存管理的底裤扒干净。从最常用的cudaMalloc和cudaFree,到背后的虚拟内存、物理内存,再到页表映射。嗯,内容有点干,但都是硬货。
核心观点:显存碎片问题的根源,在于虚拟地址与物理地址的映射关系被频繁破坏。理解这个映射,你就掌握了解决碎片问题的钥匙。
2.1 cudaMalloc / cudaFree:看似简单,实则暗藏玄机
先说说最基础的cudaMalloc和cudaFree。很多同学觉得这俩函数就是“申请一块显存”和“释放一块显存”,跟C语言的malloc/free差不多。其实不然。
cudaMalloc干了三件事:
- 在GPU的虚拟地址空间中,找一块连续的虚拟地址区间。
- 在GPU的物理内存(显存颗粒)上,分配一块物理内存。
- 建立虚拟地址到物理地址的映射关系,并更新页表。
你看,它不只是“分配内存”,它还在背后默默搭了一座桥——虚拟地址到物理地址的桥。这座桥就是页表。
cudaFree呢?
- 解除虚拟地址到物理地址的映射。
- 释放物理内存。
- 标记虚拟地址区间为“空闲”。
听起来很简单对吧?但问题就出在这里。频繁的cudaMalloc和cudaFree,会导致虚拟地址空间和物理内存都变得支离破碎。我曾在项目中遇到过,一个模型训练到一半,突然报out of memory,但用nvidia-smi一看,显存明明还有不少空闲。这就是典型的显存碎片问题。
避坑指南:我曾经在一个推理服务里,每来一个请求就cudaMalloc一次,处理完就cudaFree。结果跑了不到1000个请求,显存就炸了。后来改用内存池,问题才解决。所以,千万别在热路径上频繁调用cudaMalloc/cudaFree。
2.2 CUDA虚拟内存与物理内存:两个世界的故事
咱们得把“虚拟内存”和“物理内存”这两个概念彻底分开。你想想看,GPU的显存颗粒(比如HBM2e、GDDR6)是物理内存,容量是固定的。但CUDA程序看到的地址,是虚拟地址。
虚拟内存的好处是什么?
- 隔离性:每个CUDA context(或者说每个进程)都有自己的虚拟地址空间,互不干扰。
- 连续性:程序看到的虚拟地址是连续的,但背后的物理内存可能是不连续的。这大大简化了编程。
- 灵活性:虚拟内存可以映射到物理内存,也可以映射到磁盘(swap),虽然GPU一般不这么做。
但虚拟内存也带来了一个问题:碎片化。虚拟地址空间是有限的(比如40位地址空间,也就是1TB),物理内存也是有限的。当虚拟地址空间被切割成很多小块,即使物理内存还有空闲,也可能找不到一块足够大的连续虚拟地址区间来满足cudaMalloc的请求。
我个人的习惯是,在项目初期就用cudaMemGetInfo监控显存使用情况,同时用cuMemGetAddressRange之类的API检查虚拟地址的分配情况。这样能提前发现碎片化的苗头。
2.3 页表与映射:GPU的“地址翻译官”
页表,就是虚拟地址到物理地址的翻译官。GPU硬件在执行内存访问指令时,会拿着虚拟地址去查页表,找到对应的物理地址,然后才真正读写显存。
页表的结构(简化版):
| 虚拟页号 | 物理页框号 | 有效位 | 权限位 |
|---|---|---|---|
| 0x0000 | 0x1000 | 1 | 读写 |
| 0x0001 | 0x2000 | 1 | 读写 |
| 0x0002 | 无效 | 0 | - |
你看,虚拟页0x0000映射到物理页0x1000,虚拟页0x0001映射到物理页0x2000。虚拟页0x0002没有映射,访问它就会触发段错误。
在CUDA中,页表是由GPU驱动和硬件共同维护的。当我们调用cudaMalloc时,驱动会在页表中插入新的映射条目。调用cudaFree时,会删除对应的条目。
页表大小对性能的影响:
- 页表越大,TLB(快表)的命中率越低,内存访问延迟越高。
- 页表碎片化严重时,TLB miss会频繁发生,导致性能下降。
我记得有一次优化一个矩阵乘法的算子,发现性能始终上不去。用nvprof一分析,发现TLB miss率高达30%。后来排查发现,是因为频繁的cudaMalloc/cudaFree导致页表碎片化严重。改用cudaMallocManaged(统一内存)后,TLB miss率降到了5%以下,性能提升了20%。
小技巧:如果你发现算子性能受限于内存访问,不妨用nvidia-smi -q -d MEMORY查看显存碎片情况,或者用cuCtxGetLimit查看当前context的页表大小限制。
2.4 知识体系总览:一张图看懂CUDA内存管理
说了这么多,咱们用一张图来总结一下CUDA内存管理的核心逻辑。这张图是我自己画的,希望能帮你把零散的知识点串起来。
这张图展示了从用户程序到GPU硬件的完整链路。你看,用户程序调用cudaMalloc,CUDA驱动负责管理页表,GPU硬件负责执行地址翻译和实际的内存访问。虚拟地址空间和物理内存通过页表映射起来。碎片问题就出在这个映射关系上。
2.5 小结:记住这三句话
好了,这一章的内容就到这里。最后,我帮你总结三句最核心的话:
- cudaMalloc不是简单的分配,它是在搭桥——建立虚拟地址到物理地址的映射。
- 虚拟地址空间和物理内存是两个独立的世界,页表是连接它们的桥梁。
- 频繁的搭桥和拆桥(cudaMalloc/cudaFree)会导致桥梁结构混乱,也就是碎片化。
下一章,咱们会深入探讨显存碎片的分类和检测方法。到时候,我会分享一些我实际用过的检测工具和脚本。嗯,敬请期待。
课后思考:如果你现在有一个算子,需要频繁分配和释放大小不一的临时缓冲区,你会怎么设计内存管理策略?是直接用cudaMalloc/cudaFree,还是自己实现一个内存池?为什么?