使一些但不是全部 (CUDA) 内存访问未缓存

Question

我刚刚注意到（CUDA 内核）内存访问完全可以不缓存（参见 this answer here on SO）。

这能做到吗...

对于单独的单个内核？
在运行时而不是在编译时？
仅用于写入而不用于读取和写入？

Answer 1

仅当您单独编译该内核时，因为这是由代码生成启用的指令级功能。您还可以使用内联 PTX 汇编程序为内核中的特定加载操作发出 ld.global.cg 指令 [有关详细信息，请参阅 here]。
不，它是 PTX 的指令级功能。您可以在运行时 JIT 包含非缓存内存加载的代码版本，但这在技术上仍然是编译。您可能会使用一些模板技巧和单独的编译来让运行时保存相同代码的两个版本，这些版本使用或不使用缓存构建，并在运行时在这些版本之间进行选择。您还可以使用相同的技巧来获得给定内核的两个版本，有或没有用于未缓存加载的内联 PTX [有关实现此目的的一种可能性，请参见 here]
这些非缓存指令绕过具有字节级粒度的L1缓存到L2缓存。所以它们只是加载（所有写入都使 L1 缓存无效并存储到 L2）。

Answer 2

我不知道以前是否可行，但 CUDA 8.0 为您提供了针对特定 reads/writes 微调缓存的可能性。有关详细信息，请参阅 PTX manual。

例如，要使此代码在读取时始终进入主内存：

const float4 val = input[i];

你可以这样写：

float4 val;
const float4* myinput = input+i;
asm("ld.global.cv.v4.f32 {%0, %1, %2, %3}, [%4];" : "=f"(val.x), "=f"(val.y), "=f"(val.z), "=f"(val.w) : "l"(myinput));

我设法将我的一个缓存密集型内核的速度提高了大约 20%，使用非缓存读取和写入设计只访问一次的数据

使一些但不是全部 (CUDA) 内存访问未缓存

Making some, but not all, (CUDA) memory accesses uncached

caching

cuda

gpgpu