7

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

让我们从我所拥有的开始：

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

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

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

我的问题：你会怎么做？您能否发布一个示例或链接到具有类似主题的网站？泰

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

4

1 回答 1

6

在他的评论中，罗杰达尔链接了以下帖子

将 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);

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

于 2013-11-07T22:08:41.403 回答