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

Question

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

考虑到这一点，我想使用 JCuda 开发一个框架来获取用户的 __device__功能在 .cu文件在运行时加载并运行它。我已经实现了 __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 中链接在一起的多个模块的正确方法吗？或者说有可能吗？

编辑:

问题的第二部分是here

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 链接 session ，将每个 PTX 添加到链接器 session ，然后完成链接器 session 。这将为您提供已编译的 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 链接支持，它应该充当 JCUDA 实现的模板。