Pasando una clase C++/CUDA al SourceModule de PyCUDA

Question

Tengo una clase escrita en C++ que también usa algunas definiciones de cuda_runtime.h, esta es una parte del proyecto de código abierto llamado ADOL-C, ¡puedes mirar here!

Esto funciona cuando estoy usando CUDA-C, pero quiero de alguna manera importar esta clase en PyCUDA, si existe la posibilidad de hacerlo. Entonces, usaré esta clase dentro de kernels (no en 'main') para definir variables específicas que se usan para calcular las derivadas de una función. ¿Hay alguna forma de pasar esta clase al SourceModule de PyCUDA?

Hice una pregunta similar, pero aquí me gustaría explicar un poco más que eso. Entonces, hay una solución compilando mi código C usando nvcc -cubin (gracias a talonmies) y luego importándolo con driver.module_from_file(), pero me gustaría usar SourceModule y escribir esos núcleos dentro de un archivo .py, entonces podría ser más fácil de usar. Mi ejemplo sería algo como esto:

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

... esto es sólo una idea, pero SourceModule no sabrá cuáles son "de adouble", ya que se definen en adoublecuda.h definición de clase, por lo que Espero que entiendas mi pregunta ahora mejor. ¿Alguien tiene una pista de dónde debería comenzar? Si no, escribiré estos núcleos en CUDA-C y usaré la opción nvcc -cubin.

¡Gracias por la ayuda!

Answer 1

El sistema PyCUDA SourceModule es realmente solo una forma de obtener el código que pasa a un archivo, compilar ese archivo con nvcc en un archivo cubin y (opcionalmente) cargar ese archivo cubin en el contexto actual de CUDA. El módulo compilador PyCUDA no sabe absolutamente nada sobre la sintaxis ni el código del kernel CUDA, y tiene (casi) ninguna influencia sobre el código compilado [el calificador casi es porque puede encerrar el código enviado por el usuario con una declaración extern "C" { } para detener el crecimiento del símbolo C++].

Así que para hacer lo que creo que usted está planteando, sólo se debe requerir una declaración #include por cualquier cabeceras código de su dispositivo necesita en la cadena presentada, y un conjunto adecuado de las rutas de búsqueda en una lista pitón pasado a través de la include_dirs opción de palabra clave. Si hace algo como esto:

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']) 

y debe funcionar automágicamente (obsérvese no probado, use bajo su propia responsabilidad).