Error: External calls are not supported Functions defined inside other source files

Hi,

In the following cuda program, a (device))function is called which is defined in another source file. During compilation it gives following error:

Error: External calls are not supported (found non-inlined call to _Z5func2P4RootPi)

Following is the complete program source files:

# cat func.h

typedef struct Root

{

int a;

float f;

};

__device__ void func2(struct Root *root_d, int *x);

# cat func2.cu

#include "func.h"

__device__ void func2(struct Root *root_d, int *x)

{

root_d->a=root_d->a + 10;

root_d->f = root_d->f + 10.0;

*x=root_d->a;

*x = *x +10;

}

# cat func_multi.cu

#include<stdio.h>

#include<cuda.h>

#include "func.h"

__global__ void func1(struct Root *root_d)

{

int *x;

        root_d->a=root_d->a + 10;

        root_d->f = root_d->f + 10.0;

func2(root_d, x);

}

void func3( struct Root *root_d)

{

        func1<<<1,1>>>(root_d);

}

int main()

{

        struct Root *root, *root_d;

        root= (Root *) malloc(1*sizeof(Root));

        root->a=10;

        root->f=11.11;

cudaMalloc((void **)&root_d, 1*sizeof(Root));

        cudaMemcpy(root_d, root, 1*sizeof(Root), cudaMemcpyHostToDevice);

func1<<<1,1>>>(root_d);

func3(root_d);

cudaMemcpy(root, root_d, 1*sizeof(Root), cudaMemcpyDeviceToHost);

printf("\n AFTER COMPLETE ITERATIONS : \n a = %d \t f = %f \n ", root->a, root->f);

return 0;

}

The compilation it throws following error:

# nvcc -arch sm_20  func_multi.cu func2.cu  -o func2_multi

func.h(7): warning: declaration requires a typedef name

func_multi.cu(30): warning: variable "x" is used before its value is set

func_multi.cu(30): warning: variable "__cuda_local_var_39372_6_non_const_x" is used before its value is set

func.h(7): warning: declaration requires a typedef name

func_multi.cu(30): warning: variable "x" is used before its value is set

./func_multi.cu(30): Error: External calls are not supported (found non-inlined call to _Z5func2P4RootPi)

The GPU device is: Tesla C2050 and cuda version is:

nvcc --version

nvcc: NVIDIA ® Cuda compiler driver

Copyright © 2005-2011 NVIDIA Corporation

Built on Thu_May_12_11:09:45_PDT_2011

Cuda compilation tools, release 4.0, V0.2.1221

Is there anything wrong in compilation? How to compile it?

CUDA does not have a linker for device code, therefore all called device functions must be visible from the calling function. This means the called function needs to be either in the same file, or in a file included by the file from which the function is called.

Thanks. Its working now