NVRTC 编译应该什么时候产生 CUBIN？
When should NVRTC compilation produce a CUBIN?
如果我正确理解 NVRTC documentation 中的工作流程描述，那么它的工作原理如下：
- 从源文本创建 NVRTC 程序。
- 编译NVRTC程序获取PTX代码。
- Device-link PTX代码使用NVIDIA的Driver API (
cuLinkCreate, cuLinkAddData, cuLinkComplete)获取cubin.
但是...从 CUDA 11.3 开始，NVRTC 有以下 API 调用：
nvrtcResult nvrtcGetCUBIN ( nvrtcProgram prog, char* cubin );
那么编译后怎么会有cubin呢？
好吧，在主机端编译后就可以得到正确的机器码，那为什么不在设备端呢？
看来 cubin 的可用性取决于你编译的目标：
如果您的目标是“虚拟架构”，即某种计算能力（例如 compute_60 - 那么您唯一可以获得的是 PTX，它尚未特定于任何微架构。
如果您针对的是具体（微）架构（例如 sm_70），则编译可以一直进行到 SASS 放在 cubin 中的程序集。
现在，当您 link 使用 CUDA 驱动程序时，您有一个 context 在起作用，并且它始终与物理 GPU - 一个具体的微架构相关联.所以这必然会给你一个小房间。
PS:
- 其他开关也可能影响 cubin 输出的可用性，例如
--dlink-time-opt.
- 在 CUDA 11.3 之前，我们根本无法
nvrtcGetCUBIN()。这似乎也影响了模块的创建，即您是否可以使用 PTX 与 CUBIN 创建模块。
如果我正确理解 NVRTC documentation 中的工作流程描述，那么它的工作原理如下：
- 从源文本创建 NVRTC 程序。
- 编译NVRTC程序获取PTX代码。
- Device-link PTX代码使用NVIDIA的Driver API (
cuLinkCreate,cuLinkAddData,cuLinkComplete)获取cubin.
但是...从 CUDA 11.3 开始，NVRTC 有以下 API 调用：
nvrtcResult nvrtcGetCUBIN ( nvrtcProgram prog, char* cubin );
那么编译后怎么会有cubin呢？
好吧，在主机端编译后就可以得到正确的机器码，那为什么不在设备端呢？
看来 cubin 的可用性取决于你编译的目标：
如果您的目标是“虚拟架构”，即某种计算能力（例如
compute_60- 那么您唯一可以获得的是 PTX，它尚未特定于任何微架构。如果您针对的是具体（微）架构（例如
sm_70），则编译可以一直进行到 SASS 放在 cubin 中的程序集。
现在，当您 link 使用 CUDA 驱动程序时，您有一个 context 在起作用，并且它始终与物理 GPU - 一个具体的微架构相关联.所以这必然会给你一个小房间。
PS:
- 其他开关也可能影响 cubin 输出的可用性，例如
--dlink-time-opt. - 在 CUDA 11.3 之前，我们根本无法
nvrtcGetCUBIN()。这似乎也影响了模块的创建，即您是否可以使用 PTX 与 CUBIN 创建模块。