1

如果我正确理解了NVRTC 文档中的工作流程描述，那么它是如何工作的（假设是 CUDA 源）：

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

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

nvrtcResult nvrtcGetCUBIN ( nvrtcProgram prog, char* cubin );

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

4

1 回答 1

1

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

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

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

  • 如果您的目标是一个具体的（微）架构（例如sm_70），那么编译可以一直进行到放置在 cubin 中的 SASS 程序集。

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

PS：

  1. 其他开关也可能影响 cubin 输出的可用性，例如--dlink-time-opt.
  2. 在 CUDA 11.3 之前，我们根本做不到nvrtcGetCUBIN()。这似乎也影响了模块的创建，即您是否可以使用 PTX 与 CUBIN 创建模块。
于 2021-10-30T21:17:18.803 回答