根据“CUDA C Programming Guide”，只有在命中多处理器常量缓存时，常量内存访问才会受益（第 5.3.2.4 节）1。否则，半扭曲的内存请求可能比合并全局内存读取的情况更多。那么为什么恒定的内存大小限制为 64 KB？
为了不问两次，再问一个问题。据我了解，在 Fermi 架构中，纹理缓存与 L2 缓存相结合。纹理使用是否仍然有意义或全局内存读取以相同的方式缓存？
1常量内存（第 5.3.2.4 节）
常量内存空间驻留在设备内存中，并缓存在 F.3.1 和 F.4.1 节中提到的常量缓存中。
对于计算能力为 1.x 的设备，对 warp 的恒定内存请求首先被拆分为两个请求，一个用于每个半 warp，它们是独立发出的。
然后，一个请求被拆分为与初始请求中不同的内存地址一样多的单独请求，从而将吞吐量降低等于单独请求数量的因子。
然后在缓存命中的情况下以常量缓存的吞吐量为结果请求提供服务，否则以设备内存的吞吐量提供服务。
计算能力 1.0-3.0 设备的恒定内存大小为 64 KB。缓存工作集只有 8KB（参见 CUDA 编程指南 v4.2 表 F-2）。
驱动程序、编译器和声明的变量使用常量内存__device__ __constant__。驱动程序使用常量内存来传递参数、纹理绑定等。编译器在许多指令中使用常量（参见反汇编）。
可以使用主机运行时函数读取和写入放置在常量内存中的变量cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()（请参阅 CUDA 编程指南 v4.2 部分 B.2.2）。常量内存在设备内存中，但通过常量缓存访问。
在 Fermi 纹理上，常数、L1 和 I-Cache 都是每个 SM 内部或周围的 1 级缓存。所有 1 级缓存都通过 L2 缓存访问设备内存。
每个 CUmodule 的 64 KB 常量限制是一个 CUDA 编译单元。CUmodule 的概念隐藏在 CUDA 运行时下，但可以通过 CUDA 驱动程序 API 访问。