Metal显存不足怎么办？分块计算与数据交换实战方案

Metal应用跑到一半突然被系统中断，控制台里打印出Resource creation failure或者程序因内存压力被杀掉，这类问题的根源几乎都指向同一件事：GPU显存被吃光了。尤其在处理超大纹理、高分辨率视频帧或大规模科学计算时，数据量轻松超过几个GB，而Mac或iPhone的统一内存虽然可观，但留给GPU的配额是有限的。这篇文章不谈玄学优化，只讲两条真正能落地的路线：分块计算和数据交换，并配合代码说明怎么在Metal里实现。

先搞清楚Metal的内存模型：统一内存不等于随便用

Apple Silicon采用统一内存架构（UMA），CPU和GPU共享同一块物理内存，这让数据传输的概念变得模糊，但并不意味着没有限制。Metal内部维护着每个进程的GPU资源上限，当你创建的MTLBuffer和MTLTexture总量逼近上限时，新的资源分配会失败，已有的资源也可能被系统在内存压力下直接回收。

判断显存是否吃紧，可以借助MTLDevice的recommendedMaxWorkingSetSize属性，它返回当前设备建议的工作集大小。注意这只是建议值，不是硬性墙。真正的硬性信号是currentAllocatedSize持续增长且伴随帧率骤降，这时离崩溃就不远了。建议在开发阶段定期打印这两个数值，做到心中有数：

id<MTLDevice> device = MTLCreateSystemDefaultDevice();
NSLog(@"建议工作集: %llu MB", device.recommendedMaxWorkingSetSize / 1024 / 1024);
// 统计当前已分配的缓冲区总大小
uint64_t allocated = 0;
for (id<MTLBuffer> buf in trackedBuffers) {
    allocated += buf.length;
}
NSLog(@"已分配: %llu MB", allocated / 1024 / 1024);

另一个容易被忽视的坑是setBytes和makeBuffer的选择。小数据量（比如不超过4KB的常量）用setBytes直接内联到命令编码器里更高效，完全不走显存分配；而大数组用setBytes反而会把数据拷贝进命令缓冲，占用额外内存还拖慢提交速度。搞清楚这一点，能从源头省下不少显存。

分块计算：把大象切成小块逐块搬运

分块计算的核心思想很朴素：既然一次性装不下整个数据，那就把任务切成N块，每次只把当前需要的块送进GPU，算完立刻释放或复用缓冲区。以矩阵乘法为例，两个8192x8192的float32矩阵，光输入输出就要约800MB。如果切成256x256的tile，每块输入只需约0.5MB，GPU完全可以反复复用同一块缓冲区。

具体实现上，推荐维护一个缓冲区池：预先创建若干个固定大小的MTLBuffer，通过setBuffer:offset:atIndex:配合偏移量复用。注意Metal的缓冲区偏移对齐要求，通常要保证offset是256的倍数（不同设备可能要求不同，查询alignment相关文档确认），否则会出现数据错位这种极难排查的bug。下面是一个简化的分块矩阵乘法骨架：

// Metal着色器：每个线程组处理一个输出tile
kernel void tiledMatMul(device const float* A [[buffer(0)]],
                        device const float* B [[buffer(1)]],
                        device float* C [[buffer(2)]],
                        constant TileParams &params [[buffer(3)]],
                        uint2 gid [[thread_position_in_grid]])
{
    if (gid.x >= params.N || gid.y >= params.N) return;
    float sum = 0.0;
    // 只遍历当前分块范围内的k，配合CPU端分批上传A和B
    for (uint k = params.kStart; k < params.kEnd; ++k) {
        sum += A[uint64_t(gid.y) * params.K + k] * B[uint64_t(k) * params.N + gid.x];
    }
    atomic_fetch_add_explicit((device atomic_float*)&C[uint64_t(gid.y) * params.N + gid.x],
                              sum, memory_order_relaxed);
}

// CPU端：按k维度分块，多次提交命令
for (uint kStart = 0; kStart < K; kStart += kTileSize) {
    uint kEnd = MIN(kStart + kTileSize, K);
    // 只上传当前k范围内的A子块和B子块到预分配的缓冲区
    [sharedAContents replaceBytesInRange:NSMakeRange(0, sliceBytes)
                              withBytes:aSlicePtr];
    [sharedBContents replaceBytesInRange:NSMakeRange(0, sliceBytes)
                              withBytes:bSlicePtr];
    TileParams params = { N, K, kStart, kEnd };
    [encoder setComputePipelineState:pipeline];
    [encoder setBuffer:bufA offset:0 atIndex:0];
    [encoder setBuffer:bufB offset:0 atIndex:1];
    [encoder setBuffer:bufC offset:0 atIndex:2];
    [encoder setBytes:&params length:sizeof(params) atIndex:3];
    [encoder dispatchThreads:grid threadsPerThreadGroup:group];
}
[encoder endEncoding];
[commandBuffer commit];

这种做法的代价是多次commit带来的同步开销，以及原子操作的性能损耗。如果任务本身允许，更好的方式是在GPU端用threadgroup_memory缓存tile数据，减少全局内存访问。分块策略没有银弹，块太小会让调度开销占比过高，块太大又回到显存不足的老路，一般从让单个tile的数据量落在几百KB到几MB之间开始试验比较稳妥。

纹理类任务同理。超大图片处理可以拆成多通道渲染，每次只绑定一个子区域的纹理。利用blitEncoder的copyFromTexture:toTexture:把已经算完的子区域搬走，再复用同一段纹理内存处理下一块。针对支持Apple GPU的设备，还可以考虑memoryless渲染目标，这种纹理只存在于tile内存中，不占用主存，对临时中间结果非常友好。

数据交换：把冷数据挪出GPU，用的时候再搬回来

分块解决的是“算不过来”，交换解决的是“存不下”。有些场景必须保留全部中间结果，比如多轮迭代的机器学习推理、需要回读的历史帧缓存。这时可以采用分级存储策略：热数据留在MTLBuffer，温数据放到CPU侧的普通内存，冷数据落到磁盘，通过队列管理换入换出。

实现交换时，MTLStorageModeShared的缓冲区是最省事的选择，CPU和GPU都能直接访问，天然适合做交换的中转站。如果性能敏感，优先用MTLStorageModeManaged并严格配对didModifyRange:和synchronizeResource:，避免整块同步带来的隐性开销。一个简单的淘汰策略是最近最少使用（LRU），维护一个双向链表即可：

// 简化的LRU交换管理器
@interface BufferSwapper : NSObject
- (void)touch:(NSUInteger)bufferId;          // 标记最近使用
- (id<MTLBuffer>)loadIn:(NSUInteger)bufferId; // 换入GPU
- (void)evictIfNeeded:(NSUInteger)requiredBytes; // 换出冷数据
@end

@implementation BufferSwapper
- (void)evictIfNeeded:(NSUInteger)requiredBytes {
    while (self.gpuUsage + requiredBytes > self.budget && !self.lruList.isEmpty) {
        NSUInteger victim = [self.lruList popLeastRecentlyUsed];
        // 把victim对应的GPU缓冲区内容拷回CPU侧暂存内存，然后释放GPU资源
        NSData *cold = [self snapshotGpuBuffer:victim];
        [self.coldStore setObject:cold forKey:@(victim)];
        [self.gpuBuffers removeObjectForKey:@(victim)];
        self.gpuUsage -= [self.sizes objectForKey:@(victim)].unsignedIntegerValue;
    }
}
@end

要注意的是，换入换出本身有成本。一次几GB的memcpy在统一内存上虽然不涉及PCIe传输，但页表更新和缓存污染仍然明显。所以交换粒度要选好：太细会导致管理开销爆炸，太粗又浪费带宽。经验值是每次交换的单位控制在1MB到64MB之间。

还有一个进阶技巧是利用MTLHeap来做显存的统一调度。堆允许你预占一块大内存，然后在里面灵活分配和释放子缓冲区，避免频繁调用makeBuffer造成的碎片。配合makeAliasable，同一块堆内存还能在时间上不重叠的资源之间复用，这在多pass渲染管线里能省下可观的显存。

监控与调优：让问题在上线前暴露

做了分块和交换之后，不等于一劳永逸，必须建立监控手段。开发阶段用Instruments的Metal System Trace模板，可以看到每一帧的显存占用曲线和命令缓冲执行时间；生产环境则建议在代码里埋点，周期性上报currentAllocatedSize，配合服务端告警提前发现内存泄漏式的增长。

常见的泄漏来源包括忘记释放中间MTLTexture、命令编码器持有强引用导致资源延迟释放、以及自动释放池在长循环中没有及时drain。尤其是在循环里创建大量临时Metal对象的代码，务必用@autoreleasepool包裹每一次迭代，否则对象会堆积到runloop结束才释放，短时间内显存就会被打爆。这些细节配合分块与交换策略，才能真正让大规模Metal任务稳定跑起来。

Metal显存管理分块计算修改时间：2026-09-09 20:57:11