Passing a C++/CUDA class to PyCUDA's SourceModule

Question

I have a class written in C++ that uses also some definitions from cuda_runtime.h, this is a part from opensource project named ADOL-C, you can have a look here!

This works when I'm using CUDA-C, but I want somehow to import this class in PyCUDA, if there is a possibility of doing that. So, I will use this class inside of kernels (not in 'main') to define specific variables that are used for computing the derivatives of a function. Is there any way for of passing this class to PyCUDA's SourceModule?

I asked a similar question, but here I would like to explain a bit more that that. So, there is a solution compiling my C code using nvcc -cubin (thanks to talonmies) and then importing it with driver.module_from_file(), but, I would like to use SourceModule and write those kernels inside of a .py file, so it could be more user-friendly. My example would look something like this:

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

... this is just an idea, but SourceModule will not know what are "adouble's", because they are defined in class definition adoublecuda.h, so I hope you understand my question now better. Does anyone have a clue where should I begin? If not, I will write this kernels in CUDA-C, and use nvcc -cubin option.

Thanks for help!

Answer 1

The PyCUDA SourceModule system is really only a way of getting the code you pass into a file, compiling that file with nvcc into a cubin file, and (optionally) loading that cubin file into the current CUDA context. The PyCUDA compiler module knows absolutely nothing about CUDA kernel syntax or code, and has (almost) no influence on the code that is compiled [the almost qualifier is because it can bracket user submitted code with an extern "C" { } declaration to stop C++ symbol mangling].

So to do what I think you are asking about, you should only require an #include statement for whatever headers your device code needs in the submitted string, and a suitable set of search paths in a python list passed via the include_dirs keyword option. If you do something like this:

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

and it should automagically work (note untested, use at own risk).