Passare una classe C++/CUDA a SourceModule di PyCUDA

Question

Ho una classe scritta in C++ che utilizza anche alcune definizioni da cuda_runtime.h, questa è una parte del progetto opensource denominata ADOL-C, puoi dare un'occhiata here!

Questo funziona quando sto usando CUDA-C, ma voglio in qualche modo importare questa classe in PyCUDA, se c'è una possibilità di farlo. Quindi, userò questa classe all'interno dei kernel (non in 'main') per definire variabili specifiche che vengono utilizzate per calcolare le derivate di una funzione. C'è un modo per passare questa classe al SourceModule di PyCUDA?

Ho fatto una domanda simile, ma qui vorrei spiegarlo un po 'di più. Quindi, c'è una soluzione che compila il mio codice C usando nvcc -cubin (grazie a talonmies) e poi importandolo con driver.module_from_file(), ma, mi piacerebbe usare SourceModule e scrivere quei kernel all'interno di un file .py, quindi potrebbe essere più user-friendly. Il mio esempio sarebbe simile a questa:

from pycuda import driver, gpuarray 
from pycuda.compiler import SourceModule 
import pycuda.autoinit 
kernel_code_template=""" 
__global__ void myfunction(float* inx, float* outy, float* outderiv) 
{ 
    //defining thread index 
    ... 
    //declare dependent and independet variables as adoubles 
    //this is a part of my question 
    adtl::adouble y[3]; 
    adtl::adouble x[3]; 
    // ... 
} 
""" 

... questo è solo un'idea, ma SourceModule non sa quali sono "di adouble", perché sono definite nella definizione della classe adoublecuda.h, così ho spero che tu capisca meglio la mia domanda. Qualcuno ha la minima idea da dove dovrei cominciare? In caso contrario, scriverò questi kernel in CUDA-C e useremo l'opzione nvcc -cubin.

Grazie per l'aiuto!

Answer 1

Il sistema PyCUDA SourceModule è in realtà solo un modo per ottenere il codice passato in un file, compilare tale file con nvcc in un file cubin e (facoltativamente) caricare il file di cubin nel contesto CUDA corrente. Il modulo del compilatore PyCUDA non conosce assolutamente nulla sulla sintassi o il codice del kernel CUDA e non ha (quasi) alcuna influenza sul codice che viene compilato [il quasi qualificatore è perché può parentesi il codice inviato dall'utente con una dichiarazione extern "C" { } per arrestare il simbolo del linguaggio C++].

Quindi, per fare quello che penso ti stai chiedendo circa, si dovrebbe richiedere solo un #include dichiarazione per qualunque intestazioni il codice dispositivo ha bisogno nella stringa presentato, e di un opportuno insieme di percorsi di ricerca in una lista python passato attraverso il include_dirs opzione parola chiave. Se fai qualcosa di simile:

from pycuda import driver, gpuarray 
from pycuda.compiler import SourceModule 
import pycuda.autoinit 
kernel_code_template=""" 

#include "adoublecuda.h" 
__global__ void myfunction(float* inx, float* outy, float* outderiv) 
{ 
    //defining thread index 
    ... 
    //declare dependent and independet variables as adoubles 
    //this is a part of my question 
    adtl::adouble y[3]; 
    adtl::adouble x[3]; 
    // ... 
} 

""" 

module = SourceModule(kernel_code_template, include_dirs=['path/to/adoublecuda']) 

e dovrebbe funzionare automaticamente (nota non testata, utilizzare a proprio rischio).