CUDA 仅对一个变量禁用 L1 缓存 [英] CUDA disable L1 cache only for one variable

问题描述

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

Is there any way on CUDA 2.0 devices to disable L1 cache only for one specific variable? I know that one can disable L1 cache at compile time adding the flag -Xptxas -dlcm=cg to nvcc for all memory operations. However, I want to disable cache only for memory reads upon a specific global variable so that all of the rest of the memory reads to go through the L1 cache.

根据我在网上所做的搜索，可能的解决方案是通过 PTX 汇编代码.

Based on a search I have done in the web, a possible solution is through PTX assembly code.

推荐答案

如上所述你可以使用内联 PTX，这里是一个例子:

As mentioned above you can use inline PTX, here is an example:

__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".

You can easily vary this by swapping .f64 for .f32 (float) or .s32 (int) etc., the constraint of return_value "=d" for "=f" (float) or "=r" (int) etc. Note that the last constraint before (addr) - "l" - denotes 64 bit addressing, if you are using 32 bit addressing, it should be "r".

这篇关于CUDA 仅对一个变量禁用 L1 缓存的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！