3

在 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 以在内核之前加载userFunc ，我也会收到此错误。我被困在了这个阶段。这是加载需要在 JCuda 中链接在一起的多个模块的正确方法吗？或者甚至有可能吗？

编辑：

问题的第二部分在这里

4

1 回答 1

3

真正简短的回答是：不，您不能将多个模块加载到运行时 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 添加到链接器会话，然后完成链接器会话。这将为您提供一个已编译的 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 实现的模板。

于 2015-09-10T15:32:38.797 回答