NVRTC编译什么时候应该产生一个CUBIN?

Question

如果我正确理解NVRTC文档中的工作流描述，那么它是如何工作的:

• 从源文本创建NVRTC程序。
• 编译NVRTC程序获取PTX代码
• 设备-使用NVIDIA的驱动程序API (cuLinkCreate,cuLinkAddData,cuLinkComplete)链接PTX代码以获得舱室。

然而……从CUDA 11.3开始，NVRTC有以下API调用:

nvrtcResult nvrtcGetCUBIN ( nvrtcProgram prog, char* cubin );

那么我如何在编译后才有一个小屋呢?

Answer 1

嗯，在主机端，编译后就可以得到正确的机器码，那么在设备端为什么不能呢?

小屋的可用性似乎取决于你的编译目标:

• 如果您的目标是"虚拟体系结构"，即某种计算能力(例如compute_60)，那么您唯一可以得到的是PTX，它还没有特定于任何微体系结构。

• 如果你的目标是一个具体的(微)架构(例如sm_70)，那么编译可以一直进行到SASS组装。

现在，当你使用CUDA驱动程序链接时，你有一个上下文在起作用，并且它总是与物理GPU相关联-一个具体的微架构。这就给了你一个小隔间。

PS:

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