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(浮点数)或.s32(int)等来轻松改变这一点.对于"= f"(浮点数)或"= r"(int)等,return_value"= d"的约束.请注意,之前的最后一个约束(addr) - "l" - 表示64位寻址,如果使用32位寻址,则应为"r".

Answer 2

内联PTX可用于加载和存储变量.ld.cg和st.cg指令仅缓存L2中的数据.缓存操作符在PTX ISA 2.3文档的8.7.8.1缓存操作符一节中描述.说明或兴趣是ld和st.在CUDA中使用Inline PTX Assembly中描述了内联PTX.