caching - 使部分(但不是全部)(CUDA) 内存访问不被缓存

Question

我刚刚注意到(CUDA 内核)内存访问完全有可能不被缓存(参见 this answer here on SO )。

这可以吗...

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

Answer 1

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