How to use preprocessor symbols in PyCUDA?

Question

There is a program written in C with some CUDA kernels and I need to adapt this program to run in python using PyCUDA. Now, in this C code there are some global variables defined and some of those kernels are accessing those variables via #ifdef. What I would like to do is to define those global variables in python and then just copy my CUDA C kernels code and run them using SourceModule scheme (let's say I'm just lazy). I imagine something like this:

my_global_var=True
mod=SourceModule(""" __global__ void func() {
...
#ifdef my_global_var
do something
#endif }
""")

Is there any way to do that? Clearly, that naive attempt doesn't work at all. Should I use some kind of special pointer or something like that? In other words, what exacly happens when CUDA C kernel compiled with PyCUDA have #ifdef in it?

Edit: Well, it is possible that I'm misusing the "global variable" term here. The C code looks more or less like this:

#define X
__global__ void func(...) {
...
#ifdef X
do something
#endif }

What I'm trying to achieve is to use that C code in PyCUDA without messing with it too much. So I would like to just copy "global_ void func...", put it in PyCUDA SourceModule and define that X (no matters if it is global variable or something different) in python (like in the first pseudocode). Is it possible?

Answer 1

As with compiling any C or C++ code, external manipulation of preprocessor symbols (by which I mean outside of the code itself) requires passing additional arguments to the compiler. If you had this code:

__global__ void func(...) {
...
#ifdef X
do something
#endif 

}

then the way to set X during compilation would be to pass -DX to the compiler, which defines the symbol X during the current preprocessor pass.

You can do exactly the same thing in PyCUDA by passing options to the build using the options keyword in the SourceModule constructor. Compiler options are passed as a list, so your PyCUDA example could be implemented something like this:

my_global_var=True

....

build_options = []
if my_global_var:
    build_options.append('-Dmy_global_var')

mod=SourceModule(""" __global__ void func() {
...
#ifdef my_global_var
do something
#endif }
""", options=build_options)

[Disclaimer: Totally untested, use at own risk - I don't have a working PyCUDA installation ATM]

Here, we simply build a list containing as many entries as there are options you need to pass to the compiler, and then instantiate a SourceModule instance with those options. After that, it should just work...