在JCuda中加载多个模块不工作

Question

在jCuda中，可以将cuda文件加载为PTX或CUBIN格式，并从Java中调用(启动)__global__函数(内核)

考虑到这一点，我想用JCuda开发一个框架，在运行时在.cu文件中获取用户的__device__函数，加载并运行它。我已经实现了一个__global__函数，在这个函数中，每个线程找到它的相关数据的起点，进行一些计算和初始化，然后调用用户的__device__函数。

下面是我的内核伪代码:

extern "C" __device__ void userFunc(args);
extern "C" __global__ void kernel(){
    // initialize
    userFunc(args);
    // rest of the kernel
}

和用户的__device__功能:

extern "C" __device__ void userFunc(args){
    // do something
}

在Java方面，这是我加载模块的部分(模块是由ptx文件制成的，这些文件是用这个命令从cuda文件成功创建的:nvcc -m64 -ptx path/to/cudaFile -o cudaFile.ptx)

CUmodule kernelModule = new CUmodule(); // 1 
CUmodule userFuncModule = new CUmodule(); // 2
cuModuleLoad(kernelModule, ptxKernelFileName); // 3 
cuModuleLoad(userFuncModule, ptxUserFuncFileName); // 4

当我尝试运行它时，我在第3行得到错误:CUDA_ERROR_NO_BINARY_FOR_GPU。经过一些搜索，我得到我的ptx文件有一些语法错误。执行以下建议命令后:

ptxas -arch=sm_30 kernel.ptx

我:

ptxas fatal : Unresolved extern function 'userFunc'

甚至当我用4替换第3行来加载userFunc在内核之前，我得到这个错误。我被困在了这个阶段。这是正确的方式来加载多个模块，需要在JCuda链接在一起吗?或者这有可能吗?

编辑:

问题的第二部分在这里

Answer 1

真正简短的答案是:不，你不能在运行时API中将多个模块加载到一个上下文中。

你可以做你想做的，但它需要显式地设置和执行JIT链接调用。我不知道如何(甚至是否)在JCUDA中实现，但我可以向您展示如何使用标准驱动程序API来实现它。抓住…

如果在一个文件中有一个设备函数，而在另一个文件中有一个内核，例如:

// test_function.cu
#include <math.h>
__device__ float mathop(float &x, float &y, float &z)
{
        float res = sin(x) + cos(y) + sqrt(z);
        return res;
}

和

// test_kernel.cu
extern __device__ float mathop(float & x, float & y, float & z);
__global__ void kernel(float *xvals, float * yvals, float * zvals, float *res)
{
        int tid = threadIdx.x + blockIdx.x * blockDim.x;
        res[tid] = mathop(xvals[tid], yvals[tid], zvals[tid]);
}

你可以像往常一样编译成PTX:

$ nvcc -arch=sm_30 -ptx test_function.cu
$ nvcc -arch=sm_30 -ptx test_kernel.cu
$ head -14 test_kernel.ptx
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-19324607
// Cuda compilation tools, release 7.0, V7.0.27
// Based on LLVM 3.4svn
//
.version 4.2
.target sm_30
.address_size 64
        // .globl       _Z6kernelPfS_S_S_
.extern .func  (.param .b32 func_retval0) _Z6mathopRfS_S_

在运行时，您的代码必须创建一个JIT链接会话，将每个PTX添加到链接会话，然后完成链接会话。这将给你一个编译后的镜像的句柄，它可以像往常一样作为一个模块加载。最简单的驱动程序API代码如下所示:

#include <cstdio>
#include <cuda.h>
#define drvErrChk(ans) { drvAssert(ans, __FILE__, __LINE__); }
inline void drvAssert(CUresult code, const char *file, int line, bool abort=true)
{
    if (code != CUDA_SUCCESS) {
        fprintf(stderr, "Driver API Error %04d at %s %dn", int(code), file, line);
        exit(-1);
    }
}
int main()
{
    cuInit(0);
    CUdevice device;
    drvErrChk( cuDeviceGet(&device, 0) );
    CUcontext context;
    drvErrChk( cuCtxCreate(&context, 0, device) );
    CUlinkState state;
    drvErrChk( cuLinkCreate(0, 0, 0, &state) );
    drvErrChk( cuLinkAddFile(state, CU_JIT_INPUT_PTX, "test_function.ptx", 0, 0, 0) );
    drvErrChk( cuLinkAddFile(state, CU_JIT_INPUT_PTX, "test_kernel.ptx" , 0, 0, 0) );
    size_t sz;
    char * image;
    drvErrChk( cuLinkComplete(state, (void **)&image, &sz) );
    CUmodule module;
    drvErrChk( cuModuleLoadData(&module, image) );
    drvErrChk( cuLinkDestroy(state) );
    CUfunction function;
    drvErrChk( cuModuleGetFunction(&function, module, "_Z6kernelPfS_S_S_") );
    return 0;
}

你应该能够编译和运行这个张贴，并验证它的工作正常。它应该作为JCUDA实现的模板，如果他们已经实现了JIT链接支持。