在 JCuda 中加载多个模块不起作用

Question

在 jCuda 中，可以加载 PTX 或 CUBIN 格式的 cuda 文件并从 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'

即使我将第 3 行替换为第 4 行以在 kernel 之前加载 userFunc，我也会收到此错误。我卡在了这个阶段。这是加载需要在 JCuda 中链接在一起的多个模块的正确方法吗？或者有可能吗？

编辑：

问题的第二部分是

Answer 1

真正简短的回答是：不，您不能在运行时间内将多个模块加载到一个上下文中 API。

您可以随心所欲，但需要显式设置和执行 JIT linking 调用。我不知道在 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 link 会话，将每个 PTX 添加到 linker 会话，然后完成 linker 会话。这将为您提供编译后的 cubin 图像的句柄，它可以像往常一样作为模块加载。将它们放在一起的最简单的驱动程序 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 %d\n", 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;
}

您应该能够编译运行发布的内容并验证它是否正常工作。如果他们实施了 JIT linking 支持，它应该作为 JCUDA 实施的模板。

在 JCuda 中加载多个模块不起作用

Loading multiple modules in JCuda is not working

jit

cuda

jcuda