为什么 CUDA 中的常量内存大小受到限制？答案

【问题标题】：Why is the constant memory size limited in CUDA?为什么 CUDA 中的常量内存大小受到限制？
【发布时间】：2012-04-21 05:10:07
【问题描述】：

根据"CUDA C Programming Guide"，只有在命中多处理器常量缓存时，常量内存访问才会受益（第 5.3.2.4 节）¹。否则，半扭曲的内存请求可能比合并全局内存读取的情况更多。那么为什么将恒定内存大小限制为 64 KB？

为了不问两次，再问一个问题。据我了解，在 Fermi 架构中，纹理缓存与 L2 缓存相结合。纹理使用是否仍然有意义，或者全局内存读取以相同的方式缓存？

¹恒定内存（第 5.3.2.4 节）

常量内存空间驻留在设备内存中，缓存在F.3.1和F.4.1节中提到的常量缓存中。

对于计算能力为 1.x 的设备，对 warp 的常量内存请求首先被拆分为两个请求，每个请求一个半warp，它们是独立发出的。

然后，一个请求会被拆分为与初始请求中不同的内存地址一样多的单独请求，从而将吞吐量降低与单独请求数相等的因子。

然后在缓存命中的情况下以常量缓存的吞吐量为结果请求提供服务，否则以设备内存的吞吐量提供服务。

【问题讨论】：

标签： cuda gpgpu gpu-constant-memory

【解决方案1】：

对于计算能力为 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 级缓存。所有一级缓存都通过二级缓存访问设备内存。

64 KB 常量限制是每个 CU 模块，它是一个 CUDA 编译单元。 CUmodule 的概念隐藏在 CUDA 运行时，但可以通过 CUDA Driver API 访问。

【讨论】：

格雷格，我很抱歉可能不够清楚。我知道如何使用常量内存。在这个问题中，我想知道如果仅缓存 8 KB 实际上提供比全局内存更好的性能，为什么常量内存的大小限制为 64 KB。以同样的方式，可以通过 L1 缓存访问全局内存。那么从程序员的角度来看，使用全局内存和常量内存有什么区别，因为它们都以相同的方式缓存？
常量缓存的大小针对图形和计算着色器进行了优化。 CUDA API 公开了常量缓存，以便开发人员可以利用额外的缓存。当所有线程访问同一地址时，常量缓存具有最佳性能。命中速度非常快。未命中可以具有与 L1 未命中相同的性能。常量缓存可用于减少 L1 的抖动。计算能力 1.x 设备没有 L1 或 L2 缓存，因此使用常量缓存来优化某些访问。