我只是注意到(CUDA 内核)内存访问完全有可能不被缓存(参见例如这个答案在这里).
这可以吗...
- 对于单个内核?
- 在运行时而不是在编译时?
- 仅用于写入而不是读取和写入?
- 仅当您单独编译该内核时,因为这是通过代码生成启用的指令级功能。您还可以使用内联 PTX 汇编器来发出
ld.global.cg
内核中特定加载操作的指令[参见here详情]。
- 不,这是 PTX 的指令级功能。您可以在运行时 JIT 包含非缓存内存加载的代码版本,但这在技术上仍然是编译。您可能可以使用一些模板技巧和单独编译来让运行时保存使用或不使用缓存构建的同一代码的两个版本,并在运行时在这些版本之间进行选择。您还可以使用相同的技巧来获取给定内核的两个版本,无论是否有内联 PTX 用于未缓存加载 [请参阅here实现这一目标的一种可能性]
- 这些非缓存指令以字节级粒度绕过 L1 缓存到达 L2 缓存。因此它们只是加载(所有写入都会使 L1 缓存无效并存储到 L2)。
本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)