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

Question

根据"CUDA C编程指南",只有在多处理器常量高速缓存被命中时,常量内存访问才会受益(见第5.3.2.4节)¹.否则,对于半翘曲,可能存在比合并的全局存储器读取更多的存储器请求.那么为什么常量内存大小限制为64 KB？

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

---

¹ 恒定存储器(第5.3.2.4节)

常量存储空间驻留在设备存储器中,并缓存在F.3.1和F.4.1节中提到的常量高速缓存中.

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

然后,请求被分成多个单独的请求,因为初始请求中存在不同的内存地址,吞吐量减少的因子等于单​​独请求的数量.

然后,在高速缓存命中的情况下,或者在设备存储器的吞吐量下,以恒定高速缓存的吞吐量来服务所得到的请求.

Answer 1

对于计算能力1.0-3.0设备,常量内存大小为64 KB.高速缓存工作集仅为8KB(参见CUDA编程指南v4.2表F-2).

常量内存由声明的驱动程序,编译器和变量使用__device__ __constant__.驱动程序使用常量内存来传递参数,纹理绑定等.编译器在许多指令中使用常量(请参阅反汇编).

放置在恒定存储器变量可以读取和使用主机运行时函数写入cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()(见CUDA编程指南4.2版部分B.2.2).常量内存位于设备内存中,但可通过常量高速缓存进行访问.

在Fermi纹理上,常量,L1和I-Cache都是每个SM中或周围的1级缓存.所有级别1都通过L2缓存缓存访问设备内存.

64 KB常量限制是每CUmodule,它是一个CUDA编译单元.CUmodule的概念隐藏在CUDA运行时下,但可由CUDA驱动程序API访问.