Chargement de plusieurs modules dans JCuda ne fonctionne pas

Question

En jCuda on peut charger des fichiers cuda au format PTX ou CUBIN et appeler (lancer) __global__ fonctions (noyaux) de Java. En gardant cela à l'esprit, je veux développer un framework avec JCuda qui obtient la fonction __device__ de l'utilisateur dans un fichier .cu lors de l'exécution, le charge et l'exécute. Et j'ai déjà mis en œuvre une fonction __global__, dans lequel chaque thread trouve le point de départ de ses données connexes, effectuer un certain calcul, initialisation, puis appeler la fonction __device__ de l'utilisateur.

Voici mon code pseudo-noyau:

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

    // initialize 

    userFunc(args); 

    // rest of the kernel 
} 

et la fonction de l'utilisateur __device__:

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

Et côté Java, voici la partie que je charge les modules (modules sont fabriqués à partir ptx fichiers qui ont été créés avec succès à partir de fichiers cuda avec cette commande: 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 

Lorsque j'essaie de l'exécuter j'ai eu une erreur à la ligne 3: CUDA_ERROR_NO_BINARY_FOR_GPU. Après quelques recherches, j'obtiens que mon fichier ptx a une erreur de syntaxe. Après l'exécution de cette commande a suggéré:

ptxas -arch=sm_30 kernel.ptx

Je suis:

ptxas fatal : Unresolved extern function 'userFunc'

Même quand je Substituer à la ligne 3 avec 4 pour charger userFunc avant noyau je reçois cette erreur. Je suis resté bloqué à cette phase. Est-ce la bonne façon de charger plusieurs modules qui doivent être reliés entre eux dans JCuda? Ou est-ce même possible?

Edit:

Deuxième partie de la question est here

Answer 1

La réponse est très court: Non, vous ne pouvez pas charger plusieurs modules dans un contexte de l'API d'exécution.

Vous pouvez faire ce que vous voulez, mais cela nécessite une configuration explicite et l'exécution d'un appel de liaison JIT. Je n'ai aucune idée de comment (ou même si) cela a été implémenté dans JCUDA, mais je peux vous montrer comment le faire avec l'API de pilote standard. Attendez ...

Si vous avez une fonction de l'appareil dans un fichier, et un noyau dans un autre, par exemple:

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

et

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

Vous pouvez les compiler pour PTX comme d'habitude:

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

Lors de l'exécution, votre code doit créer une session de lien JIT, ajoutez chaque PTX à la session de liaison, puis finalisez le session de liaison.Cela vous donnera un handle pour une image cubin compilée qui peut être chargée comme un module comme d'habitude. Le plus simple possible code API de pilote pour mettre cela ensemble ressemble à ceci:

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

Vous devriez être en mesure de compiler et exécuter ce que diffusée et vérifier que cela fonctionne OK. Il devrait servir de modèle pour une implémentation JCUDA, si le support de liaison JIT est mis en œuvre.