CUDA __constant__ 对全局内存的尊重。哪个缓存？

Question

我没有向内核传递大量参数，而是使用了一个 __constant__ 变量。这个变量是一个结构数组，其中包含许多指向全局数据的指针(这些指针将是一个参数列表)；用于调用内核的多个不同数据集的数组。然后内核访问这个数组并取消对全局适当数据的引用。我的问题是，这些数据是通过 L2 缓存还是通过常量缓存缓存？此外，如果后者通过 __ldg() 加载，它会通过 L1 还是仍然是常量缓存？

更具体地说，数据本身位于全局中，但是内核取消引用 __constant__ 变量来访问它。这会对缓存产生不利影响吗？

Answer 1

由立即常量(操作码中的常量)或索引常量(通过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 指令的生命周期之后不会被缓存。