不使用NVCC时是否使用了NVIDIA的JIT编译缓存？
Is NVIDIA's JIT compilation cache used when you don't use NVCC?
我们都应该知道（但知道的人还不够多），当您使用 NVCC 构建 CUDA 程序时，运行 它在完全编译 (SASS) 代码的设备上二进制文件中不包含特定设备的 - 中间 PTX 代码是 JITed，结果实际上用于 运行ning 你的内核。在此 JITing 期间，JIT compilation cache 启动，因此，下次您 运行 相同的可执行文件时，可以跳过编译，只加载结果。
现在，假设我正在编写动态编译内核的 C++ 文件，在 运行 时间，而不是使用 NVCC，例如：
- 我使用 NVRTC 的
nvrtcCompileProgram() 编译 CUDA C++ 代码，针对具体的架构（例如 sm_70）。
- 我使用CUDA驱动的
cuModuleLoad()用内核加载一个PTX文件。
编译结果会放在那个缓存中吗？
根据我的经验观察，答案似乎是：
Compilation with
JIT cache used?
NVRTC targeting concrete architecture
No
NVRTC targeting "virtual" architecture
???
Loading a module with the CUDA driver
Yes
但我还没有对此进行广泛的探索；也许有更官方的保证。
您描述的缓存行为与 nvcc 或 nvrtc 无关。 运行time JIT 编译代码的缓存是驱动级别的机制，主要用于实现新硬件与旧代码的兼容性。
当 运行使用 运行时间或驱动程序 API 将 CUDA 代码 运行 内核 运行 时，需要考虑三种情况：
应用程序向驱动程序提供兼容的 SASS（作为 运行 时间 API 应用程序中的静态链接负载，或 SASS从文件加载，或 SASS 使用 nvrtc 以物理架构作为目标发出）。在这种情况下，SASS 被加载并执行。不涉及缓存。
应用程序提供有效的 PTX 代码（在不存在兼容 SASS 的情况下来自 fatbinary 有效负载，或者通过驱动程序 API 加载，无论该有效负载的来源是，在将虚拟体系结构用作目标的情况下，它包括 nvrtc）。在这种情况下，驱动程序触发 PTX 的 JIT 编译并加载结果 SASS 以执行。这是缓存发生的地方。驱动程序将检查 JIT 输出的用户特定私有缓存，如果它存在并且找到与之前编译的 PTX 的匹配项，它会从缓存中检索 SASS 并使用它而不是编译相同的 PTX再次。可以通过将 CUDA_CACHE_DISABLE 设置为 1 来破坏此机制。可以在 here 中找到有关此机制及其控件的更全面的讨论。如果PTX无效，则返回无效（或不兼容）的PTX错误信息给调用者，执行失败
该应用程序既不提供兼容 SASS，也不提供 PTX。在这种情况下，一个 no binary for GPU（或其 运行time API 等价物）错误将返回给调用者并且执行失败。在这种情况下，驱动程序 PTX 缓存不起作用。
那么对于你的两种情况：
I use NVRTC's nvrtcCompileProgram() to compile CUDA C++ code, targeting a concrete architecture (e.g. sm_70).
在这种情况下，您属于上述第一种或第三种情况。如果有效，二进制有效负载将被加载并执行，如果无效则失败并出现错误。没有缓存发生。
I use the CUDA driver's cuModuleLoad() to load a PTX file with the kernel.
在这种情况下，情况 2 适用。驱动程序执行缓存检查并重新使用缓存中先前的 JIT 传递输出，或者尝试执行 JIT 编译并在发生缓存未命中时缓存结果。如果 PTX 有效且兼容，则内核 运行s.
我们都应该知道（但知道的人还不够多），当您使用 NVCC 构建 CUDA 程序时，运行 它在完全编译 (SASS) 代码的设备上二进制文件中不包含特定设备的 - 中间 PTX 代码是 JITed，结果实际上用于 运行ning 你的内核。在此 JITing 期间，JIT compilation cache 启动，因此，下次您 运行 相同的可执行文件时，可以跳过编译，只加载结果。
现在，假设我正在编写动态编译内核的 C++ 文件，在 运行 时间，而不是使用 NVCC，例如：
- 我使用 NVRTC 的
nvrtcCompileProgram()编译 CUDA C++ 代码，针对具体的架构（例如sm_70）。 - 我使用CUDA驱动的
cuModuleLoad()用内核加载一个PTX文件。
编译结果会放在那个缓存中吗？
根据我的经验观察，答案似乎是：
| Compilation with | JIT cache used? |
|---|---|
| NVRTC targeting concrete architecture | No |
| NVRTC targeting "virtual" architecture | ??? |
| Loading a module with the CUDA driver | Yes |
但我还没有对此进行广泛的探索；也许有更官方的保证。
您描述的缓存行为与 nvcc 或 nvrtc 无关。 运行time JIT 编译代码的缓存是驱动级别的机制，主要用于实现新硬件与旧代码的兼容性。
当 运行使用 运行时间或驱动程序 API 将 CUDA 代码 运行 内核 运行 时，需要考虑三种情况：
应用程序向驱动程序提供兼容的 SASS（作为 运行 时间 API 应用程序中的静态链接负载，或 SASS从文件加载，或 SASS 使用 nvrtc 以物理架构作为目标发出）。在这种情况下，SASS 被加载并执行。不涉及缓存。
应用程序提供有效的 PTX 代码（在不存在兼容 SASS 的情况下来自 fatbinary 有效负载，或者通过驱动程序 API 加载，无论该有效负载的来源是，在将虚拟体系结构用作目标的情况下，它包括 nvrtc）。在这种情况下，驱动程序触发 PTX 的 JIT 编译并加载结果 SASS 以执行。这是缓存发生的地方。驱动程序将检查 JIT 输出的用户特定私有缓存，如果它存在并且找到与之前编译的 PTX 的匹配项，它会从缓存中检索 SASS 并使用它而不是编译相同的 PTX再次。可以通过将
CUDA_CACHE_DISABLE设置为 1 来破坏此机制。可以在 here 中找到有关此机制及其控件的更全面的讨论。如果PTX无效，则返回无效（或不兼容）的PTX错误信息给调用者，执行失败该应用程序既不提供兼容 SASS，也不提供 PTX。在这种情况下，一个 no binary for GPU（或其 运行time API 等价物）错误将返回给调用者并且执行失败。在这种情况下，驱动程序 PTX 缓存不起作用。
那么对于你的两种情况：
I use NVRTC's
nvrtcCompileProgram()to compile CUDA C++ code, targeting a concrete architecture (e.g.sm_70).
在这种情况下，您属于上述第一种或第三种情况。如果有效，二进制有效负载将被加载并执行，如果无效则失败并出现错误。没有缓存发生。
I use the CUDA driver's cuModuleLoad() to load a PTX file with the kernel.
在这种情况下，情况 2 适用。驱动程序执行缓存检查并重新使用缓存中先前的 JIT 传递输出，或者尝试执行 JIT 编译并在发生缓存未命中时缓存结果。如果 PTX 有效且兼容，则内核 运行s.