为什么CUDA中的恒定内存大小受到限制？

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

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

常量内存（节5.3.2.4）

恒定存储器空间驻留在设备存储器，并且在第F.3.1和F.提到的常数的高速缓存被缓存4.1。

对于计算能力1.x的设备，对于一个warp的恒定内存请求首先被分成两个请求，每个half-warp一个，分别发出。

然后，请求被分割为多个独立的请求，因为初始请求中存在不同的内存地址，吞吐量减少了一个等于单独请求数量的因子。

然后，以高速缓存命中情况下的常量高速缓存的吞吐量或其他情况下的设备内存吞吐量来处理结果请求。

来源

2012-04-21 AdelNick

计算能力1.0-3.0设备的常量内存大小为64 KB。缓存工作集只有8KB（请参见CUDA编程指南v4.2表F-2）。

驱动程序，编译器和声明为__device__ __constant__的变量使用常量内存。驱动程序使用常量内存来传递参数，纹理绑定等。编译器在许多指令中使用常量（请参见反汇编）。

使用主机运行时功能cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()（参见CUDA编程指南v4.2第B.2.2节）可以读取和写入置于恒定内存中的变量。常量内存位于设备内存中，但通过常量缓存进行访问。

关于费米纹理，常量，L1和I-Cache都是每个SM中或其周围的1级高速缓存。所有1级缓存都通过L2缓存访问设备内存。

64 KB的常量限制是CUmodule，它是一个CUDA编译单元。 CUmodule的概念隐藏在CUDA运行时环境下，但可通过CUDA驱动程序API访问。

来源

2012-04-21 20:59:25

格雷格，我很抱歉可能不够清楚。我知道如何使用常量内存。在这个问题中，我想知道为什么只有缓存8 KB才能提供比全局内存更好的性能，为什么常量内存的大小限制为64 KB。以同样的方式可以通过L1缓存访问全局内存。因此，从编程人员的角度来看，使用全局内存或常量内存之间的区别是什么，因为它们都以相同的方式被缓存？ – AdelNick 2012-04-22 16:14:33

常量缓存的大小适合图形和计算着色器。 CUDA API公开了常量缓存，以便开发人员可以利用额外的缓存。当所有线程访问相同的地址时，常量缓存具有最佳性能。点击速度非常快。错过可以有一个L1错过的相同表现。常量缓存可用于减少L1的抖动。计算能力1.x设备没有L1或L2缓存，因此常量缓存用于优化某些访问。 – 2012-04-22 19:57:02

为什么CUDA中的恒定内存大小受到限制？

回答

相关问题