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!

score 7 · Accepted Answer · answered Jul 02 '12 at 11:22

7

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).

answered Jul 02 '12 at 11:22

talonmies

70,661
34
192
269

Wow, that's the solution I was looking for! I just wanted to include this header file, so that my kernels know where is definition of adouble class, but I didn't know how. I will not use this adouble class inside of a "main", but I will need to figure out how to get this adouble array from gpu. As you can see, adouble class has only two private members: `double val` `double ADVAL` Maybe I will need to create a struct in python similar to this. Thank you very much for helping me! – Banana Jul 02 '12 at 11:41
When I try to include this class I get too much errors saying: "this declaration may not have extern "C" linkage". Do I need to change adoublecuda.h or there is something else? – Banana Jul 02 '12 at 14:19
As I noted in my answer, SourceModule can bracket code strings with an `extern "C" {}` declaration. With pure C++ definitions in your code, you don't want that. You can disable that behaviour with the `no_extern_c=True` keyword argument. There will be symbol mangling in the output, you might need to take that into accounnt in your Python code. I don't have a working PyCUDA installation at the moment to test. – talonmies Jul 02 '12 at 14:26
Yeah, if I use statement like this: `mod = SourceModule(kernel_code_template, include_dirs=['path/to/adoublecuda'],no_extern_c=True)` then I get an error when trying to access `myfunction` like: "cuModuleGetFunction failed: not found " – Banana Jul 03 '12 at 08:04
3

Here is the solution for that: [Using C++ Function Templates in PyCUDA](http://wiki.tiker.net/PyCuda/Examples/C%2B%2BFunctionTemplates) – Banana Jul 03 '12 at 09:45

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

1 Answers1