NVRTC 编译何时应生成 CUBIN？

Question

如果我正确理解 NVRTC documentation 中的工作流程描述，以下是它的工作原理（假设是 CUDA 源）：

• 从源文本创建 NVRTC 程序。
• 编译 NVRTC 程序以获取 PTX 代码。
• 使用 NVIDIA 的驱动程序 API（cuLinkCreate、cuLinkAddData、cuLinkComplete）设备链接 PTX 代码以获取 cubin。

但是...从 CUDA 11.3 开始，NVRTC 具有以下 API 调用：

nvrtcResult nvrtcGetCUBIN ( nvrtcProgram prog, char* cubin );

那么我怎样才能只编译后有一个cubin呢？

Answer 1

好吧，在主机端，您只需编译后就可以得到正确的机器代码，那么为什么不在设备端呢？

似乎 cubin 的可用性取决于您编译的目标：

• 如果您的目标是“虚拟架构”，即某种计算能力（例如 compute_60 - 那么您唯一可以获得的就是 PTX，它还不是特定于任何微架构的。

• 如果您针对具体（微）架构（例如sm_70），则编译可以一直进行到放置在 cubin 中的 SASS 程序集。

现在，当您使用 CUDA 驱动程序进行链接时，您有一个上下文在起作用，它总是与物理 GPU 相关联 - 一个具体的微架构。所以这必然会给你一个cubin。

PS：

1. 其他开关也可能影响立方体输出的可用性，例如--dlink-time-opt。
2. 在 CUDA 11.3 之前，我们根本不能nvrtcGetCUBIN()。这似乎也影响了模块的创建，即您是否可以使用 PTX 与 CUBIN 创建模块。