caching - CUDA 仅对一个变量禁用 L1 缓存

Question

CUDA 2.0 设备上是否有任何方法可以仅为一个特定变量禁用 L1 缓存？我知道可以在编译时禁用 L1 缓存，为所有内存操作添加-Xptxas -dlcm=cg标志nvcc。但是，我只想对特定全局变量的内存读取禁用缓存，以便所有其余的内存读取都通过 L1 缓存。

根据我在网络上进行的搜索，可能的解决方案是通过 PTX 汇编代码。

Answer 1

如上所述，您可以使用内联 PTX，这是一个示例：

__device__ __inline__ double ld_gbl_cg(const double *addr) {
  double return_value;
  asm("ld.global.cg.f64 %0, [%1];" : "=d"(return_value) : "l"(addr));
  return return_value;
}

您可以通过将 .f64 替换为 .f32 (float) 或 .s32 (int) 等来轻松改变这一点，return_value "=d" 的约束为 "=f" (float) 或 "=r" (int)等。注意 (addr) 之前的最后一个约束 - “l” - 表示 64 位寻址，如果您使用 32 位寻址，它应该是“r”。

Answer 2

内联 PTX 可用于加载和存储变量。ld.cg 和 st.cg 指令只缓存 L2 中的数据。缓存操作符在PTX ISA 2.3文档的 8.7.8.1 缓存操作符部分进行了描述。指令或兴趣是ld和st。在 CUDA中使用内联 PTX 程序集中描述了内联 PTX 。

Answer 3

如果您将变量声明为volatile，那么它将仅缓存在 Fermi GPU 的 L2 缓存中。请注意，某些编译器优化（例如删除重复加载）不会对 volatile 变量执行，因为编译器假定它们可能由另一个线程编写。