为什么CUDA中常量内存大小受限？

Question

为什么CUDA中常量内存大小受限？

14

根据"CUDA C编程指南", 只有当多处理器常量缓存被命中时，常量内存访问才会受益（第5.3.2.4节）¹。否则，半warp可能比聚合全局内存读取的情况下产生更多的内存请求。那么为什么常量内存大小限制为64 KB呢？

还有一个问题，为了不重复提问。据我所了解，在Fermi架构中，纹理缓存与L2缓存合并。纹理使用是否仍然有意义，或者全局内存读取以相同的方式进行缓存？

¹常量内存（第5.3.2.4节）

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

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

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

如果缓存命中，则以常量缓存的吞吐量服务于结果请求；否则以设备内存的吞吐量服务于结果请求。

- AdelNick

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Greg Smith · Accepted Answer

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

64 KB常量限制是每个CUmodule（CUDA编译单元）的限制。CUmodule的概念被隐藏在CUDA运行时中，但可通过CUDA驱动程序API访问。