Metal中constant与const device地址空间：只读缓冲区选用场景咨询

我一直以为可写缓冲区要声明为device，只读缓冲区用constant就好，但看了包括官方资源在内的很多代码示例后发现并非如此——device其实能覆盖constant的所有适用场景，但很多开发者会刻意在只读缓冲区用device，甚至在同一个kernel里同时用constant和const device，比如下面这段代码：

kernel void test
(
 constant float4 *buffer0 [[buffer(0)]],
 const device float4 *buffer1 [[buffer(1)]],
 device float4 *out_buffer [[buffer(2)]],
 )
{
    // both buffer0 and buffer1 are read-only, why they possibly would need different address spaces???
}

相关文档说明不够清晰，想请教：只读缓冲区在哪些场景下更适合用device/const device而非constant地址空间？

适用场景说明 #

• 缓冲区大小超出constant空间限制：Metal中constant地址空间的总容量有严格上限（不同GPU架构通常在64KB到256KB之间），如果只读缓冲区体积超过这个阈值，只能选择const device——device空间没有这类大小限制，可容纳更大的数据集。

• 需要动态随机或非连续寻址：constant空间的缓存是按固定块大小优化的，更适合连续、固定模式的读取（比如每个线程按相同偏移访问）。若kernel需要对缓冲区做大量随机访问、非连续寻址，或不同线程访问偏移差异大，const device的缓存机制会更适配这类不规则访问，效率更高。

• 缓冲区需频繁动态更新：虽然constant缓冲区也能更新，但更新开销远大于device缓冲区——constant空间的缓存是全局共享的，更新时需要同步所有计算单元的缓存状态。如果只读缓冲区需要频繁在CPU和GPU间更新，用const device能避免频繁刷新全局缓存带来的性能损耗。

• 统一代码逻辑或跨平台兼容：部分开发者会统一用const device处理所有只读缓冲区，不用区分不同地址空间的限制，代码逻辑更一致；另外，跨平台GPU代码框架（如适配Metal和CUDA）中，const device的语义和CUDA的__device__ const更接近，便于代码移植。

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