如何使用Metal kernel写入指定Mipmap层级？

Metal 没有提供和GLSL imageStore 直接指定mip层级写入完全对等的单接口能力，这是API设计差异导致的，不是使用方式的问题，可落地的替代方案一共有三种，根据业务场景选择即可：

方案1：CPU端绑定目标层级纹理视图（性能最优，首选方案） #

这是官方原生支持的零开销方案，不需要修改着色器逻辑：

• 不需要绑定完整的带mip链的原始纹理，在CPU端编码渲染/计算命令时，基于原始纹理创建指向单个目标mip层级的纹理视图，视图和原纹理共享同一块显存，写入视图的内容会直接同步到原纹理的对应mip层级，没有额外拷贝开销。
• 核心创建代码示例：

// 创建指向原始纹理第target_mip层的专属视图
id<MTLTexture> mipLevelView = [originTexture newTextureViewWithPixelFormat:originTexture.pixelFormat
    textureType:MTLTextureType2D
    levels:NSMakeRange(target_mip, 1) // 仅取目标mip层级，范围长度为1
    slices:NSMakeRange(0, 1)];

• 将这个视图作为普通2D可写纹理绑定到管线，着色器中直接调用纹理的write方法写入即可，和操作普通无mip的2D纹理完全一致。如果需要切换写入的mip层级，重新绑定对应层级的视图即可。

方案2：单次调度内动态写多mip层级场景：用2D数组纹理替代原生mip链 #

如果需要在一次计算内核调度中，根据运行时逻辑动态选择写入不同mip层级，不想频繁切换绑定资源，可以调整纹理的创建逻辑：

• 不创建带mip链的普通2D纹理，改为创建切片数等于总mip层级数的2D数组纹理，每个数组切片对应一个mip层级，每个切片自身不需要开启mipmap。
• 着色器中写入时直接指定数组切片索引为目标mip层级即可，核心代码示例：

kernel void dynamicWriteMip(
    texture2d_array<float, access::write> mipArray [[texture(0)]],
    constant uint& targetMip [[buffer(0)]],
    uint2 gid [[thread_position_in_grid]]
) {
    float4 output = calcRenderData(gid);
    // 直接写入对应切片，即目标mip层级
    mipArray.write(output, gid, targetMip);
}

• 注意：该方案的显存布局和原生mip链不一致，如果后续需要用到硬件三线性采样等依赖原生mip结构的能力，需要在写入完成后通过Blit命令将各切片内容拷贝到原生mip纹理的对应层级，存在额外拷贝开销，更适合仅需要手动读写各层数据、不依赖硬件自动mip采样的场景。

方案3：临时少量写入：Blit命令中转 #

如果仅需要偶尔写入指定mip层级，不想改动现有纹理创建和绑定逻辑，可以先把待写入的数据渲染/计算到一块和目标mip层级尺寸一致的临时纹理中，写入完成后调用Blit命令编码器的纹理拷贝接口，将临时纹理的内容直接拷贝到目标纹理的指定mip层级即可。

• 优点是代码改动量极小，适配成本低；缺点是多了一次全量显存拷贝，不适合每帧高频写入的场景。

注意：不要尝试在着色器中通过手动计算内存偏移的方式写入非绑定的mip层级，Metal的纹理内存布局对开发者不透明，不同GPU世代、不同厂商的mip数据排列规则存在差异，硬算地址会触发未定义行为，轻则数据写入错误，重则触发GPU重置。

内容的提问来源于stack exchange，提问作者iaomw