CUDA禁用L1高速缓存只为一个变量 [英] CUDA disable L1 cache only for one variable

问题描述

是否有关于CUDA 2.0设备的方法来禁用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组装code。

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（浮点）或.s32（INT）等，RETURN_VALUE的约束= D为= F（浮动）或= R（INT这个变化）等。需要注意的是之前的最后一个约束（地址） - 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屋！