1. 显存溢出初探：什么是显存溢出（OOM）

做GPU开发的朋友，十有八九都遇到过这么个场景：程序跑着跑着，突然崩了，终端里甩出一行血红的报错——CUDA out of memory。嗯，这就是我们今天要聊的主角：显存溢出，简称OOM。

说白了，显存溢出就是你的GPU核函数向显卡申请显存，但显卡说「我没那么多空间了」。就像你往一个杯子里倒水，水满了还继续倒，水就溢出来了。GPU也是一样，显存就那么大，你申请的量超过了可用容量，CUDA runtime就会直接罢工。

OOM的典型表现

我刚开始做CUDA开发那会儿，第一次遇到OOM还以为是显卡坏了。后来见得多了，发现OOM其实有几种固定的「症状」：

程序直接崩溃：最常见的表现。运行到某个kernel launch或者cudaMalloc时，程序直接退出，终端打印OOM错误。
系统卡死或黑屏：如果你把显存吃到超过显卡的物理上限，驱动可能会触发保护机制，导致整个桌面卡住。我在项目里遇到过，当时还以为死机了，结果等了几十秒系统自己恢复了。
性能突然断崖式下降：有时候不是直接崩，而是显存快满的时候，驱动开始做内存交换（如果支持），性能会掉到让你怀疑人生。
cudaMalloc返回错误码：如果你在代码里检查了CUDA API的返回值，会发现cudaErrorMemoryAllocation这个错误。很多新手不检查返回值，所以程序崩了都不知道原因。

⚠️ 注意： OOM不一定是显存真的用完了。有时候是显存碎片化导致的——明明总空闲显存够，但连续的大块内存分配不了。这个问题我在优化一个大规模矩阵运算时踩过坑，后面会细讲。

为什么GPU核函数容易引发OOM？

这个问题其实很有意思。你想想看，CPU程序内存溢出虽然也有，但远没有GPU这么频繁。为什么会这样？

我总结了几点核心原因：

显存容量天生受限：消费级显卡的显存通常只有8GB、12GB、24GB，和CPU动辄64GB、128GB的系统内存比起来，简直是「小水管」。你一个不小心，分配几个大张量就满了。
GPU核函数的数据「胃口」大：GPU擅长并行计算，所以核函数处理的数据量通常很大。比如一个简单的向量加法，如果你处理1000万个元素，光数据就要占几十MB。要是做深度学习训练，一个batch的数据加上中间激活值，轻松吃掉几个GB。
中间结果需要驻留显存：CPU程序可以把中间结果写回磁盘或者堆外内存，但GPU核函数不行。核函数执行期间，所有数据都必须待在显存里。你想想看，一个复杂的核函数可能有几十个中间变量，每个都是大块头。
内存管理不够精细：很多开发者写CUDA代码时，习惯一次性把所有数据都加载到显存，用完也不及时释放。我见过一个项目，核函数里分配了临时缓冲区，但忘记在函数末尾调用cudaFree，结果每次调用都泄漏一点显存，跑几次就OOM了。

核心观点： GPU核函数的OOM问题，本质上是「有限显存」和「海量数据」之间的矛盾。解决思路无非两条：要么减少显存占用，要么优化数据流动。

一个简单的OOM示例

光说不练假把式。我们来看一段代码，它几乎肯定会触发OOM：

// 一个注定OOM的核函数示例
__global__ void oom_kernel(float* d_data, int N) {
    // 每个线程分配一个巨大的局部数组
    float local_buffer[1024 * 1024]; // 4MB per thread
    
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        for (int i = 0; i < 1024 * 1024; i++) {
            local_buffer[i] = d_data[tid] * i;
        }
        d_data[tid] = local_buffer[0];
    }
}

int main() {
    float* d_data;
    cudaMalloc(&d_data, 1024 * sizeof(float));
    
    // 启动1024个线程，每个线程4MB局部变量
    // 总共需要 1024 * 4MB = 4GB 显存！
    oom_kernel<<<1, 1024>>>(d_data, 1024);
    
    cudaDeviceSynchronize();
    return 0;
}

这段代码的问题很明显：每个线程都声明了一个4MB的局部数组。如果你启动1024个线程，光这个局部变量就要占4GB显存。而局部变量是分配在寄存器或local memory里的，这些资源都是有限的。跑起来不出意外的话，马上就会OOM。

💡 小技巧： 我个人习惯在写核函数之前，先估算一下每个线程需要多少寄存器、多少共享内存、多少局部变量。用cudaOccupancyMaxPotentialBlockSize这个API可以帮你算出来。别等到程序崩了才去排查。

知识体系总览

为了让你对本章内容有个整体印象，我画了一张图，把显存溢出的核心知识点串了起来：

这张图把OOM的四个核心维度串在了一起：定义、表现、原因、解决思路。你可以把它当作一个思维导图，后续章节我们会逐一深入每个分支。

一点个人经验

我记得刚入行那会儿，接手过一个图像处理项目。核函数里有个临时数组，我图省事直接声明了固定大小 float temp[1024*1024]。当时觉得1024x1024也就4MB，没啥大不了的。结果核函数一启动，每个线程都来一份，几千个线程同时跑，显存直接爆了。

后来我养成了一个习惯：写核函数之前，先算一笔账。每个线程用多少寄存器？共享内存分配了多少？局部变量有多大？把这些加起来，乘以线程总数，心里就有数了。这个习惯帮我避免了很多OOM的坑。

💡 避坑指南： 我曾经在调试一个OOM问题时，花了整整两天才发现是某个cudaMalloc分配后忘了free。后来我写了个简单的显存跟踪工具，每次分配和释放都打日志，问题一目了然。建议你也试试这个方法。

好了，关于显存溢出的基本概念就聊到这里。记住一句话：OOM不是bug，是资源管理的问题。理解了这一点，后面的优化技巧你学起来会轻松很多。

📚 课程目录

1. 显存溢出初探：什么是显存溢出（OOM）

OOM的典型表现

为什么GPU核函数容易引发OOM？

一个简单的OOM示例

知识体系总览

一点个人经验