NVRTC 编译何时应生成 CUBIN？

Posted 2023-03-23

【中文标题】NVRTC 编译何时应生成 CUBIN？

【英文标题】：When should NVRTC compilation produce a CUBIN?

【发布时间】：2021-12-15 07:36:11

【问题描述】：

如果我正确理解 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呢？

【参考方案1】：

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

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

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

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

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

PS：

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