NVCC使用ldg而不是普通负载的充分条件(经验性)是什么？

Question

“数据自动化系统C编程指南”他说 G.4.2节

在内核的整个生命周期中只读的数据也可以通过使用__ldg()函数读取而缓存在上一节中描述的只读数据缓存中.当编译器检测到某些数据满足只读条件时，它将使用__ldg()读取它。编译器可能并不总是能够检测到某些数据是否满足只读条件。用const和__restrict__限定符标记用于加载此类数据的指针会增加编译器检测只读条件的可能性。

我的问题：

• “增加可能性”--为什么const __restrict__指针的可能性不是100%？毕竟，您确实不允许使用它们进行写作(除非您使用const_cast或其他类似的邪恶工具)。
• 根据您的经验，编译器在const __restrict__指针上使用const __restrict__本身的充分条件是什么？
• 相反，在哪些情况下，我应该显式地使用__ldg()从const __restrict__指针加载。

Answer 1

只读(纹理)缓存与全局内存不一致.因此，除了验证您的函数没有更改指向的数据之外，编译器还需要检查内核中的任何其他地方都没有修改内存。否则，该函数可能会忽略对纹理缓存仍然包含陈旧数据的数据的更改。

Answer 2

基于其他注释和回答，我将冒昧地陈述用于从指针__ldg()读取p的以下条件--每个条件都足够(但不是必要的)：

条件A

1. p是__global__函数的参数。
2. p是const __restrict__。

条件B

1. p是__device__函数f的参数。
2. p是const __restrict__。
3. 函数f只与指针算法的结果一起调用，该指针本身满足条件A。