cuModuleGetFunction 返回未找到答案

【问题标题】：cuModuleGetFunction returns not foundcuModuleGetFunction 返回未找到
【发布时间】：2017-02-05 01:28:23
【问题描述】：

我想用 nvrtc JIT 编译器编译 CUDA 内核以提高我的应用程序的性能（因此我的指令获取量增加了，但我节省了多个数组访问）。

函数看起来例如像这样，由我的函数生成器生成（不是那么重要）：

extern "C" __device__ void GetSumOfBranches(double* branches, double* outSum)
{
    double sum = (branches[38])+(-branches[334])+(-branches[398])+(-branches[411]);
    *outSum = sum;
}

我正在使用以下函数编译上面的代码：

CUfunction* FunctionGenerator::CreateFunction(const char* programText)
{
        // When I comment this statement out the output of the PTX file is changing
        // what is the reson?!
        // Bug?
        std::string savedString = std::string(programText);


        nvrtcProgram prog;
        nvrtcCreateProgram(&prog, programText, "GetSumOfBranches.cu", 0, NULL, NULL);

        const char *opts[] = {"--gpu-architecture=compute_52", "--fmad=false"};
        nvrtcCompileProgram(prog, 2, opts);

        // Obtain compilation log from the program.
        size_t logSize;
        nvrtcGetProgramLogSize(prog, &logSize);
        char *log = new char[logSize];
        nvrtcGetProgramLog(prog, log);
        // Obtain PTX from the program.
        size_t ptxSize;
        nvrtcGetPTXSize(prog, &ptxSize);
        char *ptx = new char[ptxSize];
        nvrtcGetPTX(prog, ptx);

        printf("%s", ptx);

        CUdevice cuDevice;
        CUcontext context;
        CUmodule module;
        CUfunction* kernel;
        kernel = (CUfunction*)malloc(sizeof(CUfunction));
        cuInit(0);
        cuDeviceGet(&cuDevice, 0);
        cuCtxCreate(&context, 0, cuDevice);
        auto resultLoad = cuModuleLoadDataEx(&module, ptx, 0, 0, 0);
        auto resultGetF = cuModuleGetFunction(kernel, module, "GetSumOfBranches");
        return kernel;
}

除了cuModuleGetFunction 正在返回CUDA_ERROR_NOT_FOUND 之外，一切正常。出现该错误是因为在 PTX 文件中找不到 GetSumOfBranches。

但是printf("%s", ptx); 的输出是这样的：

// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//

.version 4.3
.target sm_52
.address_size 64

    // .globl   GetSumOfBranches

.visible .func GetSumOfBranches(
    .param .b64 GetSumOfBranches_param_0,
    .param .b64 GetSumOfBranches_param_1
)
{
    .reg .f64   %fd<8>;
    .reg .b64   %rd<3>;


    ld.param.u64    %rd1, [GetSumOfBranches_param_0];
    ld.param.u64    %rd2, [GetSumOfBranches_param_1];
    ld.f64  %fd1, [%rd1+304];
    ld.f64  %fd2, [%rd1+2672];
    sub.rn.f64  %fd3, %fd1, %fd2;
    ld.f64  %fd4, [%rd1+3184];
    sub.rn.f64  %fd5, %fd3, %fd4;
    ld.f64  %fd6, [%rd1+3288];
    sub.rn.f64  %fd7, %fd5, %fd6;
    st.f64  [%rd2], %fd7;
    ret;
}

在我看来，一切都很好，cuModuleGetFunction 应该可以找到 GetSumOfBranches。你能解释一下为什么吗？

第二个问题

当我注释 std::string savedString = std::string(programText); 时，PTX 的输出就是：

// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19856038
// Cuda compilation tools, release 7.5, V7.5.17
// Based on LLVM 3.4svn
//

.version 4.3
.target sm_52
.address_size 64

这很奇怪，因为根本没有使用savedString...

【问题讨论】：

为什么你有问题的功能标记为__device__？那不是 CUDA 内核。您不能从主机代码中调用它。无论如何，对于“为什么这不起作用？”类型的问题。所以你应该提供一个minimal reproducible example
我没想到我只是在编译内核。我认为我也可以编译设备函数。您刚刚解决了我的问题（也许您想回答我的问题以获得积分？-否则谢谢）。在下一个问题中，我将提供一个最小、完整且可验证的示例。

标签： c++ cuda nvrtc

【解决方案1】：

不支持您尝试执行的操作。主机端模块管理 API 和设备 ELF 格式不公开 __device__ 函数，仅公开可通过内核启动 API 调用的 __global__ 函数。

您可以先验或在运行时编译设备函数，并以 JIT 方式将它们与内核链接，然后您可以检索这些内核并调用它们。但这就是你所能做的。

【讨论】：