How to separate the kernel file CUDA with the main .cpp file

Question

When I build the code with kernelAdd() function and main() function in the same file mainFunc.cu, it's ok.

But when I separate the kernelAdd() function in the kernelAdd.cu file and the main file in main.cpp file, it's built with the 2 errors:

"error C2065: 'add' : undeclared identifier"

and "error C2059: syntax error : '<'"

I built them in Visual Studio 2008 and Cuda v5.0.

And how do I correct its errors?

Thanks!

kernelAdd.cu

__global__ void add(int a, int b, int *c) {
*c = a + b;
}

mainFunc.cpp

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <conio.h>

int main(void) {

int c;
int *devC;

cudaMalloc((void**) &devC, sizeof(int));
add<<<1,1>>>(2,7,devC);
cudaMemcpy(&c, devC, sizeof(int), cudaMemcpyDeviceToHost);

printf("2+7=%d\n", c);

cudaFree(devC);

getch();
return 0;
}

Answer 1

error C2065: 'add' : undeclared identifier

This error has nothing to do with CUDA. The add function belongs to one compilation unit (kernelAdd.cu), while the other compilation unit (mainFunc.cpp) does not know anything about it. In order to give it this information you have to create an additional header file kernelAdd.h with the function declaration:

__global__ void add(int a, int b, int *c);

And include it in the mainFunc.cpp:

#include "kernelAdd.h"

Every cu or cpp file is compiled separately and only knows about the functions it sees in the header files it includes.

error C2059: syntax error : '<'

Now I'm guessing here (don't have VS, or even Windows to check), but it seems that VS chooses the compiler for each of the files in the project based on its extension. Thus mainFunc.cpp gets compiled with a common C++ compiler, but the <<<>>> syntax for kernel calls does not belong to the standard C++ — it is from CUDA. All the CUDA-specific syntax must be used only in those files that are going to be compiled with nvcc.

So one way to solve your problem is to rename mainFunc.cpp to mainFunc.cu. You can still keep your main file .cpp, of course, but then you will have to move the kernel call to some normal C++ function inside a cu file and expose it in a standard C++ header file, which your .cpp file will include.

Answer 2

The only addition to C to make CUDA C is the triple angle-bracket syntax to launch a kernel (<<<>>>). Everything else uses existing C features. Specifying a function as __global__ will cause nvcc to compile it for the device and create the symbols etc. so that it can be called from the host.

This means that:

1. The device code (__global__ function etc.) must be in a .cu file.
2. The host code that uses the <<<>>> syntax to launch a kernel must be in a .cu file.

You can still have all your other host code in .cpp files, you just need to put a stub in the .cu file to call the kernel, e.g. void launch_add(...) { add<<<...>>>(...); }.