为何只读缓存更适合分散读取？兼析与常量缓存的逻辑结构差异

一、性能差异的底层原因 #

1. 常量缓存为啥适合统一读取 #

• 常量缓存靠广播机制优化：当一个warp里所有线程都读取同一地址时，硬件仅需从全局内存读取一次数据，直接广播给整个warp的32个线程，完全避免重复内存访问，大幅节省带宽与延迟。
• 但它对分散访问兼容性极差——如果warp内线程访问不同地址，会触发多次独立内存请求，性能直接暴跌。

2. 只读缓存为啥适配分散读取 #

• 只读缓存（比如通过__ldg()指令调用的缓存）就是为随机/分散访问设计的：哪怕warp内线程访问不同地址，只要这些地址处于同一内存块，硬件能自动合并请求，不会像常量缓存那样发起大量独立请求。
• 它的缓存调度逻辑更灵活，专门针对线程间非统一内存访问做了优化，即便地址分散，也能高效利用内存带宽，性能不会出现断崖式下跌。

二、逻辑结构的核心区别 #

• 缓存范围与可见性
  • 常量缓存是全设备共享的，对应CUDA中的__constant__内存空间，所有线程块共用同一套缓存内容，且只能在kernel启动前由主机端加载，运行时无法修改。
  • 只读缓存是每个SM（流式多处理器）私有的，对应带__restrict__修饰的只读全局内存，或通过__ldg()读取的内存，缓存内容按需加载，不同SM的只读缓存可存储不同数据。
• 缓存大小与组织方式
  • 常量缓存容量通常较小（早期CUDA架构为64KB），内部结构围绕同地址广播优化，仅聚焦批量同地址访问的效率。
  • 只读缓存容量更大（比如Volta及以后架构中，L1的只读部分或独立只读缓存可达128KB以上），采用组相联结构，与普通数据缓存类似，专门优化随机访问的命中率和请求合并能力。
• 写入权限
  • 常量缓存对应的内存是纯只读的，kernel运行时绝对无法写入，只能由主机端提前配置。
  • 只读缓存对应的内存是主机可写、kernel只读的——主机端可在kernel运行前或运行中（通过统一内存）修改内容，但kernel内仅能读取，无法写入。

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