问题标签 [nvrtc]

概括

我正在将一个基于Scratchapixel 版本的简单光线追踪应用程序移植到一堆 GPU 库中。我使用运行时 API 和驱动程序 API 成功地将它移植到 CUDA，但是Segmentation fault (core dumped)当我尝试使用在运行时编译的 PTX 和 NVRTC 时，它会抛出一个错误。如果我取消注释#include <math.h>内核文件开头的指令（见下文），它仍然可以使用 NVCC（生成的 PTX 完全相同）但使用 NVRTC 编译失败。

我想知道如何使 NVRTC 的行为像 NVCC 一样（甚至可能吗？），或者至少了解这个问题背后的原因。

详细说明

文件kernel.cu（内核源）：

我可以使用以下代码成功编译它：（nvcc --ptx kernel.cu -o kernel.ptx此处为完整的 PTX）并在驱动程序 API 中使用该 PTX，并cuModuleLoadDataEx使用以下代码段。它按预期工作。

即使我取消注释该#include <math.h>行，它也能正常工作（实际上，生成的 PTX 完全相同）。

但是，每当我尝试使用 NVRTC（此处为完整的 PTX）编译这个确切的内核时，它都会成功编译，但会给我一个Segmentation fault (core dumped)调用cuModuleLoadDataEx（当尝试使用生成的 PTX 时）。

如果我取消注释该#include <math.h>行，它会在nvrtcCompileProgram调用中失败并显示以下输出：

我用来用 NVRTC 编译它的代码是：

然后我只需ptxSource使用前面的代码片段加载（注意：该代码块与驱动程序 API 版本和 NVRTC 版本相同）。

到目前为止我注意到/尝试过的其他事情

1. NVCC 生成的PTX和NVRTC生成的PTX完全不同，但我无法理解它们以识别可能的问题。
2. 尝试向编译器指定特定的 GPU 架构（在我的情况下为 CC 6.1），没有区别。
3. 试图禁用任何编译器优化（选项--ftz=false --prec-sqrt=true --prec-div=true --fmad=false）nvrtcCompileProgram。PTX 文件变大了，但仍然存在Segfaulting。
4. 尝试添加--std=c++11或添加--std=c++14到 NVRTC 编译器选项。对于它们中的任何一个，NVRTC 都会生成一个几乎为空的（4 行）PTX，但在我尝试使用它之前不会发出警告或错误。

环境

• 所以：Ubuntu 18.04.4 LTS 64 位
• nvcc --version：Cuda 编译工具，10.1 版，V10.1.168。建于 Wed_Apr_24_19:10:27_PDT_2019
• gcc --version: gcc (Ubuntu 7.5.0-3ubuntu1~18.04) 7.5.0
• 硬件：英特尔 I7-7700HQ、GeForce GTX 1050 Ti

在 OP+1 日编辑

我忘了添加我的环境。请参阅上一节。

你也可以用ptxas编译nvrtc输出吗？——@talonmies 的评论

-生成的nvccPTX 编译时带有警告：

这是由于递归内核函数（更多内容）。可以放心地忽略它。

-生成的nvrtcPTX无法编译并发出错误：

基于这个问题，我添加__device__了Sphere类构造函数并删除了--device-as-default-execution-space编译器选项。它现在生成的 PTX 略有不同，但仍然显示相同的错误。

用now编译会#include <math.h>产生很多“没有执行空间注解的函数被认为是宿主函数，在JIT模式下不允许宿主函数”。除了先前的错误之外的警告。

如果我尝试使用该问题的公认解决方案，它会给我带来一堆语法错误并且无法编译。NVCC 仍然完美无缺。

下载optix6.0.0 sdk并编译包含的示例项目后，我在运行任何示例项目（例如optixHello，optixWhitted）时遇到运行时异常，并显示错误消息：

我的环境是：vs2015, cuda11.0.

代码对于（cmake、compile、runtime）来说工作得很好ooptix sdk 7.0.0，并且大多数预编译的示例项目工作得很好。.exesdk 6.0.0

我写了一个很棒的内核，它会给我带来名利——如果我只能让它用 NVRTC 编译的话：

我希望系统头文件应该被（运行时）编译器识别，就像常规编译器一样，并且这将“正常工作”（以任何特定于 printf 的机器为模）。或者，如果它不起作用，我会期待一条关于“程序stdio.h创建”API 调用（nvrtcCreateProgram()NULLNULL

但是，我得到的是以下内容：

这对我来说似乎很奇怪。这意味着运行时编译器能够查看系统头文件，但无法找到stddef.h，就像 nvcc 或主机端编译器一样。

为什么会发生这种情况，惯用/推荐的解决方法是什么？

注意：我想要一个跨平台的解决方法，而不仅仅是在我的个人机器上工作。

我正在分析使用 nvrtc 库编译的内核（带有调试和 lineinfo）。在分析结果中，许多样本被列为在__nv_nvrtc_builtin_header.h. 但是 - 磁盘上显然没有这样的文件，自然（？）NVIDIA Compute 源视图无法找到它。

我的问题：

• 里面到底有什么__nv_nvrtc_builtin_header.h？
• 我可以查看这个神秘标题的内容吗？（如果有帮助，假设我用来执行编译的代码可以修改/添加到。）

我正在使用 NVRTC 编译内核。相关的 API 调用是：

如您所见，源是原始字符串，与文件无关。这意味着当你--generate-line-info，你得到行号，但没有相关的文件名。这意味着如果您随后使用 NSight Compute - 您将无法看到您的内核源代码。

显然，无论是 NSight Compute 本身还是 NVRTC 本身，都无法确定原始源镜像在某个文件中。但是必须有一些方法来解决这个问题：

• 也许我在 NVRTC API 中遗漏了一些可以使源 <-> 文件关联的东西？
• 也许我们可以操纵生成的编译程序（合理地，而不是手动，或编写我自己的新 API）来建立关联？
• 也许我们可以以某种方式将源代码推入编译的程序中？

我有这个库，其中一部分依赖于 nvrtc，我正在为此构建示例程序。当我在我自己的机器（Devuan GNU/Linux Beowulf）上构建它们时，一切都很好，使用 CUDA 版本 9.x、10.x 和 11.x（无论如何都是大多数）。但是，当我在 Travis CI 构建机器上执行此操作时（每个人都可能无法访问链接），我收到以下错误：

这特别奇怪，因为检测到 CUDA 工具包：

我正在通过 .deb 文件安装 nvrtc 库：

当 NVCC 将内核编译为 PTX 代码时，内核的参数具有以下形式：mangled_name_of_kernel_param_0等mangled_name_of_kernel_param_1。

有没有办法让 NVCC 使用实际的参数名称？或者至少是带有一些后缀/前缀的参数名称以避免与保留名称冲突？

NVRTC 同样的问题。

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

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

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

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

CUDA 运行时 API具有用于从主机端代码处理设备端全局变量的功能cudaGetSymbolAddress()，使用它们的名称（源代码标识符）作为句柄。cudaGetSymbolSize()

在驱动程序 API 中，我们有cuModuleGetGlobal()，它可以让我们做同样的事情......除了它需要一个全局符号所在的 CUmodule。如果您正在使用动态编译并加载/添加到模块中的代码那么你就准备好了。但是，如果这些全局变量是您程序的一部分，使用 NVCC 静态编译而不是动态加载呢？

我会假设每个编译的程序都有某种“主模块”或“默认模块”，其中包含内置的全局变量和函数。我能得到它的句柄吗？

我想知道是否可以使用选项 CU_JIT_LTO 在即时 (JIT) 链接期间改进链接时间优化 (LTO)。如果是这样，我该如何指定这个选项？

我在 NVIDIA 开发者博客中找到了以下代码，但我不明白为什么将 walltime 赋予 CU_JIT_LTO。博客中没有定义 walltime 变量。当我尝试类似的事情时，它对我的​​内核性能没有影响。

来源：https ://developer.nvidia.com/blog/discovering-new-features-in-cuda-11-4/

我的示例案例使用输入选项CU_JIT_INPUT_NVVM来链接使用 LTO 标志（-dlto或-code=lto_80）创建的对象。似乎链接器已经做了一些 LTO，因为内核“执行得更好”，而不是链接没有 LTO 的目标文件，但不如使用 NVCC 链接 LTO。（有关详细结果和讨论，请参见示例案例）

示例案例

为了检查链接时间优化 (LTO) 的有效性，我使用 4 种不同的方法创建了一个简单的程序，并使用每个线程的寄存器数量作为指标。这在我的系统上给出了以下结果（操作系统：ubuntu 20.04，CUDA 工具包：11.5.1，NVIDIA 驱动程序：495.44，GPU：NVIDIA RTX 3080）。

结果解读：

从单个翻译文件创建程序应该可以得到最好的结果。编译器可以看到所有的函数实现并使用它来优化内核。这导致 30 个寄存器/线程。

使用 NVCC 与 LTO 链接绝对有效。它使用与从单个 .cu 文件编译的程序相同数量的寄存器 (30)，而没有 LTO 的情况则不同（使用 44 个寄存器）。

使用 NVRTC/JIT 链接文件比在没有 LTO 的情况下链接 NVCC 做得“更好”（当我们只关注寄存器使用时），但不如使用 LTO 链接 NVCC。内核使用 38 个寄存器/线程。

注意：我的目标不是减少寄存器的使用，我只是把它作为一个指标。因为来自单个翻译文件的程序使用 30 个寄存器/线程，所以我假设一个完全优化的链接程序将具有相同的“最终可执行代码”，因此使用相同数量的寄存器。因为情况并非如此，所以我开始研究 JIT 选项。

CU_JIT_LTO 选项：

我尝试使用 JIT_option CU_JIT_LTO 进一步优化 NVRTC/JIT 案例中的链接。但是，我不确定如何使用此选项。我尝试了以下两种方式（有关更多上下文，请参见下面的文件cuda 代码。链接代码从第 41 行开始）：

方法 1：将选项 CU_JIT_LTO 添加到cuLinkCreate(...). 这似乎没有任何效果。int lto = 0该代码在和时使用相同数量的寄存器int lto = 1。

方法 2：将选项 CU_JIT_LTO 添加到cuLinkAddFile(...)和cuLinkAddData(...)。这会立即给出错误 CUDA_ERROR_INVALID_VALUE。

所以现在我的问题是：应该如何使用选项 CU_JIT_LTO？

文件：

下面是两个文件。按照以下步骤运行示例（在 linux 操作系统上）：

1. 将 cuda 代码保存在扩展名为 .cu 的文件中（例如：code.cu）
2. 将 bash 脚本保存在文件中（例如：run.sh）
3. 从终端运行此命令：bash run.sh code.cu

cuda代码：

bash脚本：

子问题：生成 NVVM IR 文件

要生成func_link_nvrtc_lto.o与命令一起使用的文件cuLinkAddFile(linker, CU_JIT_INPUT_NVVM, "func_link_nvrtc_lto.o", ...)，我必须添加-ptx标志，如下面的命令所示。我没有在任何文档中找到这一点，而是通过尝试和错误。我想知道是否有更好的方法来生成这样的文件。