generate c code

In programming guide
“nvcc’s basic workflow consists in separating device code from host code and
compiling the device code into a binary form or cubin object. The generated host
code is output either as C code that is left to be compiled using another tool or as
object code directly by invoking the host compiler during the last compilation stage.”

How to output my c++ host code to c code?


Just compile your c++ code with g++ or whatever compiler you normally use and link everything together at the object level. There is an example of this in the SDK.

Could you make it more clearly? How to separate c++ host code from device code and then compile c++ to c code? Thanks!

Well, you can do it however you want to do it! It doesn’t really matter how you break things up. Just compile all your normal C++ code with g++ and the GPU kernel code with nvcc.

If you really want an example, I can show you how I do it.

Each class in my code that does something on the GPU ends up being spread across 4 files:
<- C code only. Declares any structs used on the GPU and kernel drivers to be used by SomeClass
<- C code only. Defines any kernels and kernel driver functions to be used by SomeClass
<- declaration of the C++ class
<- definition of the C++ class with all host code. Includes SomeClass.cuh and makes calls to the kernel driver function when it needs to run something on the GPU. Class methods provide data structure allocation and management

The *.cu files are compiled with nvcc. The *.cc files are compiled with g++/visual studio. All the resulting object files are linked together at the end to produce a python module. (well, my app produces a python module, yours doesn’t have to).

Thanks for your example. I understand what you said. I use the same method in my application. What I want to know is how to change the c++ host code to c code. Can g++ do this? Moreover, if c++ host code and device code is in a same .cu file how to do it?


I’m not aware of any automatic tool to convert C++ host code to C. You need to do it by hand. I’m not even going to ask why you would want to do such a thing.

You cannot. nvcc does not support this.

I’m confusing about

"…The generated host

code is output either as C code that is left to be compiled using another tool…" in Programming guide2.0

__global__ void kernel()


   .... some code .....

   // this is compiled by nvcc to ptx and then compiled to a cubin


void host_function()


   ... a bunch of C code ....

   // this C code is passed through to the generated file.

   kernel<<< grid, threads >>>();

   // The kernel call is turned into a bunch of calls to the driver API (the generated C-code you keep asking about). The final resulting C file is compiled by gcc/visual c++ to produce a final object file (which includes the cubin from the nvcc compiled kernel).


If you want to see the generated C file, just compile with nvcc -keep and you can look at all the intermediate files.