一、背景：我踩过的Metal性能坑

1.1 坑的具体样子

我去年做一个开放世界手游项目，其中地形块的实时光照计算是用Metal计算着色器写的。本来这个功能是赶进度做的，跑是能跑，但测的时候发现，加载1000个地形块（每个16×16大小）要300多毫秒，玩家在地图上移动时，地形加载和光照渲染卡得厉害，用户反馈很不好。后来我花了一周查问题，从代码逻辑到GPU特性折腾下来，才发现是共享内存的布局坑了我。

1.2 发现问题的过程

一开始我以为是计算量太大，或者线程组大小设错了，改了线程组从8×8到32×32，性能没什么变化。后来用Xcode的Metal GPU计数器看，发现共享内存的内存冲突率高达78%——这才反应过来，是线程对共享内存的访问太集中，导致GPU的内存模块堵车了，并行效率自然上不去。

二、原来的错误写法：我踩的第一个坑

2.1 错误代码的具体内容

// 技术栈：Metal Shading Language (MSL)
kernel void computeTerrainLight(
    const device float* inHeightmap [[ buffer(0) ]],  // 输入的地形高度图
    device float* outLightmap [[ buffer(1) ]],       // 输出的光照图
    threadgroup float sharedBlock[16][16] [[ threadgroup ]], // 线程组共享内存，用来缓存地形块局部数据
    uint2 gridPos [[ thread_position_in_grid ]],     // 当前线程在整个计算网格的位置
    uint2 localPos [[ thread_position_in_threadgroup ]] // 当前线程在所属线程组的位置
) {
    // 错误写法：直接用线程组内坐标当共享内存索引，会触发严重的内存冲突
    sharedBlock[localPos.y][localPos.x] = inHeightmap[gridPos.y * 1024 + gridPos.x];
    threadgroup_barrier(mem_flags::mem_threadgroup); // 等所有线程加载完共享内存

    // 后续光照计算，要反复读共享内存的相邻数据
    float light = 0.0;
    for(int i=-1; i<=1; i++){
        for(int j=-1; j<=1; j++){
            light += sharedBlock[localPos.y+i][localPos.x+j];
        }
    }
    outLightmap[gridPos.y * 1024 + gridPos.x] = light /9.0;
}

2.2 错误的核心原因

Metal的共享内存会分成一个个叫bank的内存模块，每个bank一次只能处理一个线程的内存请求。如果多个线程同时访问同一个bank，就必须排队，这就是常说的内存冲突，相当于大家挤同一个窗口办业务，效率直接被拉低。刚才的代码里，同一行的所有线程（比如固定localPos.y，localPos.x从0到15），访问的共享内存是连续的同一块，全部落入同一个bank，自然要排队堵车。

三、重构后的共享内存访问模式：怎么改的

3.1 优化后的代码

// 技术栈：Metal Shading Language (MSL)
kernel void computeTerrainLightOptimized(
    const device float* inHeightmap [[ buffer(0) ]],
    device float* outLightmap [[ buffer(1) ]],
    threadgroup float sharedBlock[16][16] [[ threadgroup ]],
    uint2 gridPos [[ thread_position_in_grid ]],
    uint2 localPos [[ thread_position_in_threadgroup ]]
) {
    // 优化点：用线程坐标的交错偏移，让同一行线程访问不同bank
    // 偏移值=线程y坐标+线程x坐标，把连续地址打散
    uint linearOffset = localPos.y + localPos.x;
    uint sharedY = linearOffset / 16; // 得到共享内存的行号
    uint sharedX = linearOffset % 16; // 得到共享内存的列号
    sharedBlock[sharedY][sharedX] = inHeightmap[gridPos.y *1024 + gridPos.x];
    
    threadgroup_barrier(mem_flags::mem_threadgroup);

    // 后续光照计算，访问共享内存时也用同样的偏移逻辑，避免再次冲突
    float light =0.0;
    for(int i=-1; i<=1; i++){
        for(int j=-1; j<=1; j++){
            uint sampleLinear = (localPos.y+i) + (localPos.x+j);
            light += sharedBlock[sampleLinear/16][sampleLinear%16];
        }
    }
    outLightmap[gridPos.y *1024 + gridPos.x] = light /9.0;
}

3.2 优化的核心逻辑

简单说就是“把人分散到不同窗口”：用线程坐标的交错偏移，把原本连续的内存地址打散，让同一行的线程访问不同的bank，就像每个人选了不同的窗口，不用排队。这个调整只改了共享内存的索引逻辑，后续计算完全没变，却直接解决了内存冲突的问题。

四、优化的应用场景、优缺点和注意事项

4.1 适用的场景

哪些地方需要这么改？主要是需要大量线程重复访问局部数据的计算任务，比如游戏里的粒子系统计算、物理碰撞预计算、纹理滤波处理、实时光照采样这些，都是共享内存冲突的重灾区，也是这个优化见效最快的地方。

4.2 技术的优缺点

优点很明显：一是改造成本极低，只调整索引公式，不用加额外计算逻辑，也不用换硬件；二是适用范围广，几乎所有用共享内存的Metal计算着色器都能试。缺点也有：一是需要了解目标GPU的bank数量，苹果A系列GPU一般是32个bank，M系列是64个，偏移量要微调；二是只解决内存访问问题，要是任务本身是纯计算密集型的，提升可能不明显。

4.3 必须注意的细节

第一，一定要查目标设备的GPU bank信息，比如M系列用32为偏移，冲突率会更低；第二，别滥用共享内存，如果数据每个线程只用一次，就用全局内存，共享内存大小有限（一般线程组最大16KB），塞太多会溢出；第三，一定要用Metal的GPU计数器测，不能凭感觉，冲突率降到10%以下，性能就基本没问题；第四，要测试不同的线程组大小，16×16和32×32的交错偏移效果可能不一样。

五、总结

这个坑其实很多开发新手或者赶项目的人都会踩——大家一般只关注“功能能不能跑”，不会深入看硬件的内存特性。Metal的共享内存本来是用来加速局部数据访问的，但如果布局不对，反而会拖慢速度。我这次重构只是改了几行索引代码，就把地形光照的计算时间从320ms降到了110ms，玩家的流畅度提升非常明显。这个经验就是想告诉大家，写着色器的时候，不仅要写对，还要写得让硬件“舒服”，这样才能拿到最好的性能。