我可以将“静态”CUDA 内核启动与 PTX 代码结合起来并获得一个工作二进制文件吗？

Posted 2023-03-23

【中文标题】我可以将“静态”CUDA 内核启动与 PTX 代码结合起来并获得一个工作二进制文件吗？

【英文标题】：Can I combine a "static" CUDA kernel launch with PTX code and get a working binary?

【发布时间】：2021-08-17 16:42:34

【问题描述】：

假设我使用一个 CUDA 程序 - 例如 CUDA vectorAdd 示例，并剪掉内核的实现，但仍然有启动命令：

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);

假设我自己写了 PTX，因为我是一个 DIY 类型的人，所以现在我有 vectorAdd.cu 没有内核的 CUDA 代码和 vectorAdd.ptx。

我现在可以生成一个可执行文件，它可以像未修改的 vectorAdd 那样工作，但在 PTX 中运行代码吗？

（假设 PTX 没有尝试任何有趣的事情或做错任何事情。）

注意事项：

这个问题是一个变体：

How can I create an executable to run a kernel in a given PTX file?

除了在那个问题中，发帖人愿意使用驱动程序 API 来动态加载和编译使用驱动程序 API 的 PTX 文件。在这里，这不是一个选项：C++ 代码使用三人字形 CUDA 运行时启动，这不能改变。

我不介意创建涉及生成其他文件的可执行文件的过程，例如一个立方体。

【问题讨论】：

我完全不明白你的问题（或它的理由）。两件事：另一个问题 isn't 根本不使用 nvrtc（即实时设备 C++ 编译器），您似乎在描述的是基本的运行时 API PTX JIT 功能，它已经存在自第一个 CUDA beta 版本以来。据我所知，您所问的问题与您所链接的问题完全相同。你期待什么不同的答案？

@talonmies：见最后一段的编辑。您对 NVRTC 的看法是正确的，其他问题没有使用它。

所以您问是否可以使用运行时 API 进行 JIT？为什么不实际问这个问题？无论如何，答案是一个字——不。

@talonmies：不，我不想做任何 JIT。我想提前编译——但内核代码不是 C++/CUDA，而是 PTX。

@talonmies: 1. 这不是无关紧要的，因为 PTX 是用例的输入。 2. 不是“还是不行”，因为之前你没看懂我问的意思，就关闭了。如果现在的答案是“否”，请重新打开/重新投票，然后回答“否”。

【参考方案1】：

在调用者可见的标头中将您的函数定义为

__ global __ void vectorAdd(void* d_A, void* d_B, void* d_C, int numElements);

使用空声明创建文件vectorAdd.cu

__ global __ void vectorAdd(void* d_A, void* d_B, void* d_C, int numElements) 

使用合适的选项调用

nvcc --keep vectorAdd.cu

将 vectorAdd.ptx 替换为您的版本 调用

nvcc -fatbin -dlink

来创建 fatbin 和 cubin 文件 调用 nvcc -link 链接 .cubin 文件和 .cudafe1.cpp 或 cudafe1.c（取决于语言）文件。它们还依次包含 .cudafe1.stub.c 和 .fatbin.c 文件 在您的项目中使用生成的 .obj 或 .o 文件 (Windows/Linux) 以 CUDA 运行时方式调用 vectorAdd>>

（作为高级 DIY 人，您将来会想编写 SASS 代码，这是特定于设备的低级汇编语言。）

【讨论】：

在步骤 (5.) 中，我为什么文件调用 nvcc？

在ptx文件上（加了），也可以指定Cuda架构

好的，nvcc -link 怎么样？它将采取哪些论据？我猜不是.c 或.cpp 文件。

nvcc -link 编译和链接主机和设备代码。它可以为主机代码调用指定的 C++ 编译器。 .cpp 和 .c 文件在步骤 3 中生成，.fatbin.c（包含在 .cudafe1.cpp 中）在步骤 5 中生成。您可以使用与通常的 makefile/IDE 相同的参数进行编译.cu 文件，作为输入，您指定提到的 .cpp/.c 文件。内核位于 fatbin.c 中，由运行时 API 自动加载。