如何在运行时生成、编译和运行CUDA核函数

Question

如何在运行时生成、编译和运行CUDA核函数

cudacompilationgpgpuptx

7

嗯，我有一个相当微妙的问题 :)

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

数据，大量数据数组，已复制到GPU
程序，由CPU（主机）生成，需要对该数组中的每个数据进行评估
程序经常更改，可以生成为CUDA字符串、PTX字符串或其他什么东西 (?)，并且每次更改后需要重新评估

我想要的：基本上只是想尽可能地使其有效（快速），例如避免将CUDA编译为PTX。解决方案甚至可以完全针对设备特定，这里不需要太大的兼容性 :)

我知道的：我已经知道函数cuLoadModule，它可以从文件中加载和创建PTX代码的内核。但是我认为，必须有其他方法可以直接创建内核，而无需首先将其保存到文件中。或者也许可以将其存储为字节码？

我的问题：你会怎样做？你能发布一个示例或链接到类似主题的网站吗？TY

编辑：好的，现在，PTX内核可以直接从PTX字符串（char数组）中运行。无论如何，我仍然想知道，是否有更好/更快的解决方案？仍然存在从字符串转换为某些PTX字节码的问题，这应该可以避免。我还怀疑，可能存在一种聪明的方式，可以从PTX创建设备特定的Cuda二进制文件，从而消除JIT编译器的滞后（虽然很小，但如果您需要运行大量内核，则可能会累加）：）

- teejay

2

https://dev59.com/nXDYa4cB1Zd3GeqPB5Z8 - Roger Dahl

1

在CUDA表达式模板中，作者使用表达式模板技术为每个表达式类型在运行时生成一个CUDA内核。请查看代码。您对这样的东西感兴趣吗？ - Vitality

+1 指向有趣且相关的文章，最终它们会从文件中加载 PTX，但我喜欢它们如何完全抽象出CUDA的向量运算 :) 很高兴看到其他人是如何做的，谢谢您链接代码Jack！ - teejay

我认为上述论文是使用CUDA表达式模板的第一个例子。我想到的批评是，也许需要在运行时即时生成CUDA代码，并编译和加载PTX代码，这可能会挫败使用表达式模板的优势。如果您对CUDA中的表达式模板感兴趣，那么自那以后，其他库已经被开发出来：Newton使用thrust，J.M. Cohen，“Processing Device Arrays with C++ Metaprogramming”，GPU Computing Gems - Jade Edition等等。 - Vitality

最近，我们开发了一个BlueBird表达式模板库，可在主机和设备上运行，并旨在接近Matlab的语法。目前它是一个beta版本。 - Vitality

就像你所说的那样 - 对于某些应用程序来说，动态PTX代码生成和JIT编译成设备二进制文件的开销可能太大了。这就是为什么我仍然希望，其他一些CUDA大师能够来到这里展示一些避免这些问题的神奇方法 :) - teejay

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

在他的评论中，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);
    }

最后使用cuModuleLoad和cuModuleGetFunction从文件中加载PTX代码，并将其传递给编译器驱动程序，例如：

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

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