Cargando múltiples módulos en JCuda no está funcionando

jit

Question

En jCuda uno puede cargar archivos cuda como formato PTX o CUBIN y llamar ( __global__ ) __global__ funciones (kernels) desde Java.

Teniendo esto en cuenta, quiero desarrollar un framework con JCuda que obtenga la función __device__ del usuario en un archivo .cu en tiempo de ejecución, lo cargue y lo ejecute. Y ya he implementado una función __global__ , en la cual cada hilo descubre el punto de inicio de sus datos relacionados, realiza algunos cálculos, inicializa y luego llama a la función __device__ del usuario.

Aquí está mi pseudo código de kernel:

extern "C" __device__ void userFunc(args);
extern "C" __global__ void kernel(){

    // initialize

    userFunc(args);

    // rest of the kernel
}

Y la función __device__ del usuario:

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

Y en el lado de Java, aquí está la parte que cargué los módulos (los módulos están hechos de archivos ptx que se crean con éxito desde archivos cuda con este comando: 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

Cuando intento ejecutarlo CUDA_ERROR_NO_BINARY_FOR_GPU error en la línea 3: CUDA_ERROR_NO_BINARY_FOR_GPU . Después de buscar, entiendo que mi archivo ptx tiene algún error de sintaxis. Después de ejecutar este comando sugerido:

ptxas -arch=sm_30 kernel.ptx

Tengo:

ptxas fatal : Unresolved extern function ''userFunc''

Incluso cuando reemplazo la línea 3 por 4 para cargar UserFunc antes del núcleo , obtengo este error. Me quedé atrapado en esta fase. ¿Es esta la forma correcta de cargar múltiples módulos que deben vincularse entre sí en JCuda? ¿O es posible?

Editar:

La segunda parte de la pregunta está aquí

Answer 1

La respuesta realmente corta es: No, no puede cargar múltiples módulos en un contexto en la API de tiempo de ejecución.

Puede hacer lo que quiera, pero requiere una configuración y ejecución explícita de una llamada de enlace JIT. No tengo idea de cómo (o incluso si) se ha implementado en JCUDA, pero puedo mostrarte cómo hacerlo con la API de controlador estándar. Espere...

Si tiene una función de dispositivo en un archivo y un kernel en otro, por ejemplo:

// 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;
}

y

// 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]);
}

Puede compilarlos en PTX como siempre:

$ 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_

En tiempo de ejecución, su código debe crear una sesión de enlace JIT, agregar cada PTX a la sesión del enlazador y luego finalizar la sesión del enlazador. Esto le dará un control sobre una imagen de cubos compilados que puede cargarse como un módulo como de costumbre. El código de API de controlador más simple posible para armar esto se ve así:

#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;
}

Debería poder compilar y ejecutar esto como se publicó y verificar que funciona bien. Debería servir como plantilla para una implementación JCUDA, si tienen implementado el soporte de enlace JIT.