CUDA __constant__尊重全球记忆.哪个缓存？

我没有将大量参数传递给内核,而是使用__constant__变量.这个变量是一个结构数组,它包含许多指向全局数据的指针(这些指针是一个参数列表); 一个数组,用于调用内核的多个不同数据集.然后,内核访问此数组并取消引用全局适当的数据.我的问题是,这些数据是通过L2还是常量缓存缓存的？而且,如果后者和(如果加载的__ldg()话)通过L1或仍然是常量高速缓存？

更具体地说,数据本身位于全局中,但内核取消引用__constant__变量以获取它.这会对缓存产生负面影响吗？

由立即常量(操作码中的常量)或索引常量(通过ldc指令访问)访问的常量变量由(bank,offset)对访问,而不是由地址访问.这些读取通过立即常量和索引常量缓存.在某些芯片上,这些是相同的缓存.常量访问的示例是:

// immediate constant
ADD r0, r1, c[bank][offset]

// r1 has packed version of bank, offset
LDC r0, r1

传递cc2.0及更高版本的参数,以便您可以看到立即的常量访问.

常量访问通过常量内存层次结构,最终产生一个全局地址,该地址可以在系统内存或设备内存中.

如果将常量变量设置为指向全局的指针,则将通过数据层次结构读取数据.

如果定义const变量,编译器可以选择将只读数据放在存储体/偏移量或地址中.

如果您查看SASS(nvdisasm或工具),您将看到LD说明.根据芯片的不同,这些数据可以缓存在L1/Tex缓存中,然后缓存在L2缓存中.

SHARED
LDS/STS/ATOMS             -> shared memory

GENERIC
LD/ST (generic to shared) -> shared memory
LD/ST (generic to global) -> L1/TEX -> L2
LD/ST (generic to local)  -> L1/TEX -> L2

LOCAL
LDL/STL (local)           -> L1/TEX -> L2

GLOBAL
LDG/STG (global)          -> TEX    -> L2

INDEXED CONSTANT
LDC -> indexed constant cache -> ...-> L2

L2未命中可以转到设备存储器或固定系统存储器.

在你提到的情况下,很可能通过立即常量访问常量变量(假设合理的常量大小,最佳性能),并且取消引用的指针将导致全局内存访问.

在GK110上,LDG指令缓存在纹理缓存中.

在Maxwell LDG.CI指令缓存在纹理缓存中.LDG.CA操作缓存在纹理缓存(GM20x)中.所有其他LDG访问都通过纹理缓存,但不会在warp指令的生命周期内缓存.