📚 课程目录第4章显存复用策略（下）

4. 显存复用策略（下）

好，咱们接着聊显存复用。上一讲我们把池化、Alias、生命周期管理这些基础策略讲透了。这一讲，我要带你深入几个更硬核的方向——Shared Memory、跨进程IPC、数据压缩，还有显存卸载。

说实话，这些技术我在实际项目中踩过的坑，比吃过的盐还多。你想想看，一个吞吐压测跑下来，显存爆了，或者带宽打满了，那种感觉……嗯，咱们今天就把这些痛点一个个解决掉。

核心观点：显存复用不只是“省着用”，更是“调度着用”。Shared Memory让你在芯片内部玩出花，IPC让你跨进程共享，压缩让你用精度换吞吐，卸载让你把显存当缓存用。

4.1 显存共享（Shared Memory）—— 片上的“快车道”

Shared Memory，说白了就是GPU芯片内部的一块高速缓存。它跟L1 cache在物理上是同一块硬件，但你可以手动控制它的内容。我个人的习惯是：只要数据会被反复访问，优先考虑塞进Shared Memory。

为什么？因为Global Memory的延迟是几百个cycle，而Shared Memory只有几十个cycle。你想想看，一个循环里读10次数据，用Shared Memory能省多少时间？

4.1.1 典型用法：矩阵分块（Tiling）

矩阵乘法是Shared Memory最经典的场景。我来给你看一段代码，这是我当年优化一个推荐系统模型时用的套路：

// 每个线程块处理一个子矩阵
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    float sum = 0.0f;
    for (int t = 0; t < N / TILE_SIZE; ++t) {
        // 协作加载到Shared Memory
        As[ty][tx] = A[by * TILE_SIZE + ty][t * TILE_SIZE + tx];
        Bs[ty][tx] = B[t * TILE_SIZE + ty][bx * TILE_SIZE + tx];
        __syncthreads();
        
        // 从Shared Memory计算
        for (int k = 0; k < TILE_SIZE; ++k)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }
    C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}

我的经验：Shared Memory的大小是有限的——每个SM通常只有48KB到164KB。我曾经在一个项目里试图塞一个256x256的矩阵，结果直接编译失败。记住：TILE_SIZE不要超过32，否则bank conflict会让你欲哭无泪。

4.1.2 避坑指南：Bank Conflict

Shared Memory被分成32个bank，每个bank宽度4字节。如果多个线程同时访问同一个bank的不同地址，就会发生bank conflict，导致串行化。

我曾经犯过一个低级错误：用一个二维数组，列方向访问。结果性能直接掉了一半。后来改成行方向访问，或者加padding，才把带宽跑满。

4.2 跨进程显存复用（IPC）—— 进程间的“握手”

多进程场景下，每个进程都有自己的虚拟地址空间。如果你在进程A里分配了一块显存，进程B想直接用，怎么办？

答案就是CUDA IPC。它允许你把一个进程中的设备指针“导出”给另一个进程。我当年做分布式推理时，这个技术救了我的命。

4.2.1 核心API

流程很简单：

进程A调用 cudaIpcGetMemHandle 获取一个句柄
把句柄通过共享内存或socket传给进程B
进程B调用 cudaIpcOpenMemHandle 拿到指针

// 进程A
float *d_ptr;
cudaMalloc(&d_ptr, size);
cudaIpcMemHandle_t handle;
cudaIpcGetMemHandle(&handle, d_ptr);
// 发送handle给进程B...

// 进程B
cudaIpcMemHandle_t handle; // 从进程A接收
float *d_ptr_b;
cudaIpcOpenMemHandle(&d_ptr_b, handle, cudaIpcMemLazyEnablePeerAccess);
// 现在d_ptr_b和d_ptr指向同一块显存

注意：IPC只支持同一台机器上的GPU。跨机器的场景，你得用GPUDirect RDMA或者NCCL。另外，记得在进程退出前调用cudaIpcCloseMemHandle，否则会内存泄漏。我曾经因为这个被运维同事追着骂了三天。

4.3 显存压缩技术（FP16/INT8）—— 用精度换吞吐

显存不够？带宽打满？最简单的办法就是——压缩数据。

FP16和INT8是两种最常见的压缩格式。FP16把32位浮点砍成16位，精度损失大约1-2个数量级，但显存占用和带宽直接减半。INT8更狠，直接降到8位整数，但需要做量化校准。

4.3.1 什么时候用FP16？

我个人习惯：只要模型对精度不敏感，优先上FP16。比如推理场景，尤其是视觉模型，FP16几乎不影响准确率。

但训练场景要小心。我记得有一次训练一个BERT模型，用FP16跑了一周，结果loss死活不降。后来发现是梯度下溢了——FP16的最小正数是6e-8，比FP32的1e-38大得多。解决方案是混合精度训练：前向用FP16，反向用FP32，再配合loss scaling。

4.3.2 INT8量化

INT8需要做校准。你拿一批数据跑一遍，统计出每个tensor的min/max，然后映射到[-128, 127]区间。TensorRT和ONNX Runtime都支持自动量化。

格式	显存节省	精度损失	适用场景
FP32	1x	无	训练、高精度推理
FP16	2x	~0.1%	推理、混合精度训练
INT8	4x	~0.5-1%	推理（需校准）

我的建议：别一上来就上INT8。先跑FP16，如果精度达标，就别折腾了。INT8的校准过程很麻烦，而且对异常值敏感。我曾经在一个语音模型上试INT8，结果因为一个离群点，整个模型输出全是噪声。

4.4 显存卸载（Offloading）—— 把显存当缓存用

显存不够了怎么办？卸载到CPU内存或者SSD上。这个思路其实很简单：把不常用的数据挪走，用的时候再搬回来。

但这里有个关键问题：卸载的粒度。你是一次卸载整个tensor，还是按层卸载？我见过有人把整个模型参数卸载到CPU，结果每次forward都要从CPU搬数据，性能直接崩了。

4.4.1 卸载策略

按层卸载：只卸载那些长时间不用的层。比如Transformer的中间层，在推理时只用到一次，可以卸载。
按batch卸载：把当前不用的batch数据卸载到CPU，等需要时再预取回来。
异步卸载：用CUDA Stream做异步传输，让计算和传输重叠。

// 异步卸载示例
cudaStream_t stream;
cudaStreamCreate(&stream);

// 异步拷贝到CPU
cudaMemcpyAsync(h_ptr, d_ptr, size, cudaMemcpyDeviceToHost, stream);

// 同时做其他计算
kernel<<<grid, block, 0, stream>>>(d_other);

// 等传输完成
cudaStreamSynchronize(stream);

注意：PCIe带宽是有限的。一张A100的PCIe 4.0带宽只有64GB/s，而HBM带宽是2TB/s。你卸载得越多，PCIe瓶颈越严重。我建议：卸载比例不要超过30%，否则得不偿失。

知识体系总览

下面这张图，是我梳理的显存复用策略全景。你可以把它当作一个决策树：

这张图把四个策略串起来了。你从上往下看：先判断你的瓶颈是延迟还是带宽？延迟敏感就用Shared Memory；多进程场景用IPC；带宽打满就上压缩；显存不够就卸载。

好了，这一讲的内容就到这里。记住：没有银弹。每个策略都有适用场景，也有代价。你需要根据实际压测数据，找到那个平衡点。