CUDA __device__ Unresolved external function

Question

I am trying to figure out how to separate CUDA __device__ codes in separate header files.

I have three files.

File: 1: int2.cuh

 #ifndef INT2_H_ #define INT2_H_ #include "cuda.h" #include "cuda_runtime.h" #include "device_launch_parameters.h" __global__ void kernel(); __device__ int k2(int k); int launchKernel(int dim); #endif /* INT2_H_ */ 

File 2: int2.cu

 #include "int2.cuh" #include "cstdio" __global__ void kernel() { int tid = threadIdx.x; printf("%d\n", k2(tid)); } __device__ int k2(int i) { return i * i; } int launchKernel(int dim) { kernel<<<1, dim>>>(); cudaDeviceReset(); return 0; } 

File 3: CUDASample.cu

 include <stdio.h> #include <stdlib.h> #include "int2.cuh" #include "iostream" using namespace std; static const int WORK_SIZE = 256; __global__ void sampleCuda() { int tid = threadIdx.x; // printf("%d\n", k2(tid)); //Can not call k2 printf("%d\n", tid * tid); } int main(void) { int var; var = launchKernel(16); kernel<<<1, 16>>>(); cudaDeviceReset(); sampleCuda<<<1, 16>>>(); cudaDeviceReset(); return 0; } 

The file is working. I can call the kernel sampleCuda() (in the same file), call the C function launchKernel() (in another file) and call kernel() directly (in another file).

The problem I am having is calling the __device__ function from the sampleCuda() kernel. then it shows the following error. However, the same function can be called in kernel() .

 10:58:11 **** Incremental Build of configuration Debug for project CUDASample **** make all Building file: ../src/CUDASample.cu Invoking: NVCC Compiler /Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 -gencode arch=compute_20,code=sm_20 -odir "src" -M -o "src/CUDASample.d" "../src/CUDASample.cu" /Developer/NVIDIA/CUDA-6.5/bin/nvcc -G -g -O0 --compile --relocatable-device-code=false -gencode arch=compute_20,code=compute_20 -gencode arch=compute_20,code=sm_20 -x cu -o "src/CUDASample.o" "../src/CUDASample.cu" ../src/CUDASample.cu(18): warning: variable "var" was set but never used ../src/CUDASample.cu(8): warning: variable "WORK_SIZE" was declared but never referenced ../src/CUDASample.cu(18): warning: variable "var" was set but never used ../src/CUDASample.cu(8): warning: variable "WORK_SIZE" was declared but never referenced ptxas fatal : Unresolved extern function '_Z2k2i' make: *** [src/CUDASample.o] Error 255 10:58:14 Build Finished (took 2s.388ms) 

Answer 1

The problem is that you defined the __device__ function in a separate compilation unit from __global__ that calls it. You need to either explicitly enable the switch mode of the device by adding the -dc flag, or by moving the definition to the same unit.

From nvcc documentation:

--device-c|-dc Compile each .c / .cc / .cpp / .cxx / .cu file into an object file that contains the code of the roaming device. This is equivalent to --relocatable-device-code = true --compile .

For more information, see Separate Compilation and Linking of CUDA C ++ Device Code .