CUDA __device__ 函数的 JIT 编译
JIT compilation of CUDA __device__ functions
我有一个固定的内核，我希望能够合并用户定义的设备函数来改变输出。用户定义的函数将始终具有相同的输入参数，并且将始终输出一个标量值。如果我在编译时知道用户定义的函数，我可以将它们作为指向内核的指针传递（并且如果没有给定函数，则有一个默认设备函数在输入上运行）。我可以在 运行 时间访问用户定义函数的 PTX 代码，我想知道我是否可以使用 NVIDIA 的 jitify 之类的东西在 运行 时间编译 PTX，获取指向设备函数的指针，然后将此设备函数传递给预编译的内核函数。
我看到一些帖子接近回答这个问题 (How to generate, compile and run CUDA kernels at runtime)，但大多数建议在 运行 时编译整个内核和设备功能。鉴于设备功能具有固定的输入和输出，我看不出有任何理由无法提前编译内核功能。我缺少的部分是如何在 运行 时间只编译设备函数并获取指向它的指针然后传递给内核函数。
I have access to the user defined function's PTX code at runtime and
am wondering if I could use something like NVIDIA's jitify to compile
the PTX at run time, get a pointer to the device function, and then
pass this device function to the precompiled kernel function.
不，你不能那样做。 NVIDIA 的 API 不公开设备功能，只公开完整的内核。所以没有办法获得运行时编译的设备指针。
您可以使用运行时使用 NVRTC. However, you can only do this via the driver module APIs 编译的设备函数执行预编译内核（PTX 或 cubin）的运行时链接。运行时 API 没有公开该功能（并且根据我对运行时 API 工作方式的理解，如果不对注入嵌入式静态编译代码的方式进行一些重大的架构更改，它可能无法公开在运行时）。
您可以执行以下操作：
- 使用--keep生成您的cuda项目，并为您的cuda项目查找生成的ptx或cubin。
- 在运行时，生成您的 ptx（在我们的实验中，我们需要将函数指针存储在设备内存区域中，声明一个全局变量）。
- 在运行时从 cuLinkCreate, adding first the ptx or cubin from the --keep output and then your runtime generated ptx with cuLinkAddData 开始构建一个新模块。
- 最后，调用你的内核。但是您需要使用新生成的模块而不是使用 <<<>>> 符号来调用内核。在后一种情况下，它将在函数指针未知的模块中。最后一个阶段应该使用驱动程序 API 完成（您可能还想尝试运行时 API cudaLaunchKernel）。
主要元素是确保从生成的模块调用内核，而不是从与您的程序神奇链接的模块。
我有一个固定的内核，我希望能够合并用户定义的设备函数来改变输出。用户定义的函数将始终具有相同的输入参数，并且将始终输出一个标量值。如果我在编译时知道用户定义的函数，我可以将它们作为指向内核的指针传递（并且如果没有给定函数，则有一个默认设备函数在输入上运行）。我可以在 运行 时间访问用户定义函数的 PTX 代码，我想知道我是否可以使用 NVIDIA 的 jitify 之类的东西在 运行 时间编译 PTX，获取指向设备函数的指针，然后将此设备函数传递给预编译的内核函数。
我看到一些帖子接近回答这个问题 (How to generate, compile and run CUDA kernels at runtime)，但大多数建议在 运行 时编译整个内核和设备功能。鉴于设备功能具有固定的输入和输出，我看不出有任何理由无法提前编译内核功能。我缺少的部分是如何在 运行 时间只编译设备函数并获取指向它的指针然后传递给内核函数。
I have access to the user defined function's PTX code at runtime and am wondering if I could use something like NVIDIA's jitify to compile the PTX at run time, get a pointer to the device function, and then pass this device function to the precompiled kernel function.
不，你不能那样做。 NVIDIA 的 API 不公开设备功能，只公开完整的内核。所以没有办法获得运行时编译的设备指针。
您可以使用运行时使用 NVRTC. However, you can only do this via the driver module APIs 编译的设备函数执行预编译内核（PTX 或 cubin）的运行时链接。运行时 API 没有公开该功能（并且根据我对运行时 API 工作方式的理解，如果不对注入嵌入式静态编译代码的方式进行一些重大的架构更改，它可能无法公开在运行时）。
您可以执行以下操作：
- 使用--keep生成您的cuda项目，并为您的cuda项目查找生成的ptx或cubin。
- 在运行时，生成您的 ptx（在我们的实验中，我们需要将函数指针存储在设备内存区域中，声明一个全局变量）。
- 在运行时从 cuLinkCreate, adding first the ptx or cubin from the --keep output and then your runtime generated ptx with cuLinkAddData 开始构建一个新模块。
- 最后，调用你的内核。但是您需要使用新生成的模块而不是使用 <<<>>> 符号来调用内核。在后一种情况下，它将在函数指针未知的模块中。最后一个阶段应该使用驱动程序 API 完成（您可能还想尝试运行时 API cudaLaunchKernel）。
主要元素是确保从生成的模块调用内核，而不是从与您的程序神奇链接的模块。