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) etc 的约束。请注意，(addr) 之前的最后一个约束 - "l"- 表示 64 位寻址，如果您使用 32 位寻址，它应该是 "r"。