我的程序中有一个名为 float valueAt(float3 v) 的函数。它应该在给定点返回函数的值。该功能是用户指定的。我目前有这个函数的解释器，但其他人建议我在线编译这个函数，这样它就可以使用机器代码并且速度更快。
我该怎么做？我相信我知道在生成 PTX 时如何加载该函数，但我不知道如何生成 PTX。
【问题讨论】：
nvcc -ptx mycode.cu 从 cuda C/C++ 源代码生成 PTX。显然，用户定义的 CUDA 函数的运行时编译和导入比这更复杂，但我认为没有标准的编程方式可以做到这一点。您可以尝试在程序的适当位置发出系统调用并检查结果，假设您可以将所需的函数放入表示语法正确、可编译的 CUDA 函数的文件中。你的问题要么很简单，要么很复杂。同样，我不确定是否有“标准”方式来做到这一点。
标签： cuda compilation
CUDA 不提供非 PTX 代码的运行时编译方式。
您可以完成您想要的，但不能使用标准的 CUDA API。 PyCUDA 为 CUDA C 代码提供了一种优雅的即时编译方法，其中包括工具链的幕后分叉以编译为设备代码并使用运行时 API 加载。 （可能的）缺点是您需要将 Python 用于您的应用程序的顶层，如果您要将代码发送给第三方，您可能还需要发送一个可以工作的 Python 发行版。
我能想到的唯一其他选择是 OpenCL，它确实支持运行时编译（直到最近才支持它）。 C99 语言库比 CUDA 提供的限制要多得多，而且我发现 API 非常冗长，但运行时编译模型运行良好。
【讨论】：
我已经考虑了这个问题一段时间，虽然我不认为这是一个“很好”的解决方案，但它似乎确实有效，所以我想我会分享它。
基本思路是使用linux产生进程编译然后运行编译后的代码。我认为这很简单，但是由于我将各个部分放在一起，因此我将在此处发布说明，以防对其他人有用。
问题中的问题陈述是能够取一个包含用户定义函数的文件，我们假设它是单个变量f(x)的函数，即y = f(x)，并且x和y可以由float数量表示。
用户将编辑一个名为fx.txt 的文件，其中包含所需的功能。此文件必须符合 C 语法规则。
fx.txt：
y=1/x
然后这个文件被包含在__device__ 函数中：
user_testfunc.cuh：
__device__ float fx(float x){
float y;
#include "fx.txt"
;
return y;
}
它被包含在通过包装器调用的内核中。
cudalib.cu:
#include <math.h>
#include "cudalib.h"
#include "user_testfunc.cuh"
__global__ void my_kernel(float x, float *y){
*y = fx(x);
}
float cudalib_compute_fx(float x){
float *d, *h_d;
h_d = (float *)malloc(sizeof(float));
cudaMalloc(&d, sizeof(float));
my_kernel<<<1,1>>>(x, d);
cudaMemcpy(h_d, d, sizeof(float), cudaMemcpyDeviceToHost);
return *h_d;
}
cudalib.h：
float cudalib_compute_fx(float x);
以上文件被构建到一个共享库中：
nvcc -arch=sm_20 -Xcompiler -fPIC -shared cudalib.cu -o libmycudalib.so
我们需要一个主应用程序来使用这个共享库。
t452.cu：
#include <stdio.h>
#include <stdlib.h>
#include "cudalib.h"
int main(int argc, char* argv[]){
if (argc == 1){
// recompile lib, and spawn new process
int retval = system("nvcc -arch=sm_20 -Xcompiler -fPIC -shared cudalib.cu -o libmycudalib.so");
char scmd[128];
sprintf(scmd, "%s skip", argv[0]);
retval = system(scmd);}
else { // compute f(x) at x = 2.0
printf("Result is: %f\n", cudalib_compute_fx(2.0));
}
return 0;
}
编译如下：
nvcc -arch=sm_20 -o t452 t452.cu -L. -lmycudalib
此时，主应用程序 (t452) 可以执行，它会产生 f(2.0) 的结果，在这种情况下为 0.5：
$ LD_LIBRARY_PATH=.:$LD_LIBRARY_PATH ./t452
Result is: 0.500000
然后用户可以修改fx.txt文件：
$ vi fx.txt
$ cat fx.txt
y = 5/x
只需重新运行应用程序，就会使用新的功能行为：
$ LD_LIBRARY_PATH=.:$LD_LIBRARY_PATH ./t452
Result is: 2.500000
这种方法利用了这样一个事实，即在重新编译/替换共享库时，一个新的 linux 进程将获取新的共享库。另请注意，为了清楚起见，我省略了几种错误检查。至少我会检查 CUDA 错误，并且我还可能会在重新编译之前删除共享对象 (.so) 库，然后在编译后测试它是否存在，以进行编译是否成功进行的基本测试。
此方法完全使用运行时 API 来实现此目标，因此用户必须在其计算机上安装 CUDA 工具包并进行适当设置，以便nvcc 在 PATH 中可用。使用带有 PTX 代码的驱动程序 API 将使这个过程更加清晰（并且不需要用户机器上的工具包），但是 AFAIK 没有使用nvcc 或用户创建的工具链，就无法从 CUDA C 生成 PTX。 nvidia llvm 编译器工具。将来，“标准”CUDA C 工具链中可能会有更“集成”的方法，甚至可能由驱动程序提供。
可以使用设备代码的单独编译和链接来安排类似的方法，这样需要向用户公开的唯一源代码位于user_testfunc.cu（和fx.txt）中。
编辑：现在有一个 CUDA runtime compilation facility，应该用来代替上述内容。
【讨论】：