如何在运行时生成,编译和运行CUDA内核

Question

好吧,我有一个非常微妙的问题:)

让我们从我拥有的东西开始:

1. 数据,大量数据,复制到GPU
2. 由CPU(主机)生成的程序,需要针对该阵列中的每个数据进行评估
3. 该程序的变化很频繁,可以作为CUDA字符串,字符串PTX或别的东西(？)产生并需要重新评估每一个变化之后

我想要的:基本上只是想让它尽可能有效(快速),例如.避免将CUDA编译为PTX.解决方案甚至可以完全针对特定设备,这里不需要大的兼容性:)

我所知道的:我已经知道函数cuLoadModule,它可以从存储在文件中的PTX代码加载和创建内核.但我想,必须有一些其他方法可以直接创建内核,而不必先将其保存到文件中.或者也许可以将其存储为字节码？

我的问题:你会怎么做？您可以发布一个示例或链接到类似主题的网站吗？TY

编辑:好了,PTX内核可以直接从PTX字符串(char数组)运行.无论如何我仍然想知道,有没有更好/更快的解决方案呢？仍然存在从字符串到某些PTX字节码的转换,应该可以避免.我也怀疑,从PTX创建设备特定的Cuda二进制文件的一些聪明的方法可能存在,这将删除JIT编译器滞后(很小,但如果你有大量的内核要运行它可以加起来):)

Answer 1

在他的评论中，Roger Dahl 链接了以下帖子

直接将PTX程序传递给CUDA驱动

其中解决了两个函数的使用，即cuModuleLoad和cuModuleLoadDataEx。前者用于从文件加载 PTX 代码并将其传递给nvcc编译器驱动程序。后者避免了 I/O，并且能够将 PTX 代码作为 C 字符串传递给驱动程序。在任何一种情况下，您都需要已经拥有 PTX 代码，无论是作为 CUDA 内核编译的结果（加载或复制并粘贴到 C 字符串中）还是作为手写源代码。

但是，如果您必须从 CUDA 内核开始即时创建 PTX 代码，会发生什么？按照CUDA 表达式模板中的方法，您可以定义一个包含 CUDA 内核的字符串，例如

ss << "extern \"C\" __global__ void kernel( ";
ss << def_line.str() << ", unsigned int vector_size, unsigned int number_of_used_threads ) { \n";
ss << "\tint idx = blockDim.x * blockIdx.x + threadIdx.x; \n";
ss << "\tfor(unsigned int i = 0; i < ";
ss << "(vector_size + number_of_used_threads - 1) / number_of_used_threads; ++i) {\n";
ss << "\t\tif(idx < vector_size) { \n";
ss << "\t\t\t" << eval_line.str() << "\n";
ss << "\t\t\tidx += number_of_used_threads;\n";
ss << "\t\t}\n";
ss << "\t}\n";
ss << "}\n\n\n\n";

然后使用系统调用将其编译为

int nvcc_exit_status = system(
         (std::string(NVCC) + " -ptx " + NVCC_FLAGS + " " + kernel_filename 
              + " -o " + kernel_comp_filename).c_str()
    );

    if (nvcc_exit_status) {
            std::cerr << "ERROR: nvcc exits with status code: " << nvcc_exit_status << std::endl;
            exit(1);
    }

最后使用cuModuleLoadandcuModuleGetFunction从文件加载 PTX 代码并将其传递给编译器驱动程序，如

    result = cuModuleLoad(&cuModule, kernel_comp_filename.c_str());
    assert(result == CUDA_SUCCESS);
    result =  cuModuleGetFunction(&cuFunction, cuModule, "kernel");
    assert(result == CUDA_SUCCESS);

当然，表达式模板与这个问题无关，我只是引用了我在这个答案中报告的想法的来源。