cuda - 如何在运行时生成、编译和运行 CUDA 内核

Question

嗯，我有一个非常微妙的问题:)

让我们从我拥有的开始:

1. 数据，大​​量数据，复制到 GPU
2. 程序，由 CPU(主机)生成，需要针对该数组中的每个数据进行评估
3. 程序更改非常频繁，可以生成为 CUDA 字符串、PTX 字符串或其他内容(？)，并且在每次更改后需要重新评估<

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

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

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

Edit: OK now, PTX kernel can be run from PTX string (char array) directly. Anyways I still wonder, is there some better / faster solution to this? There is still conversion from string to some PTX bytecode, which should be possibly avoided. I also suspect, that some clever way of creating device specific Cuda binary from PTX might exist, which would remove JIT compiler lag (is small, but it can add up if you have huge numbers of kernels to run) :)

Answer 1

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

Passing the PTX program to the CUDA driver directly

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

但是，如果您必须从 CUDA 内核开始即时创建 PTX 代码，会发生什么情况？遵循 CUDA Expression templates 中的方法，您可以定义一个包含 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);

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