CUDA External calls not supported

Question

I am developing a CUDA 4.0 application running on a Fermi card. According to the specs, Fermi has Compute Capability 2.0 and therefore should support non-inlined function calls.

I compile every class I have with nvcc 4.0 in a distinct obj file. Then, I link them all with g++-4.4.

Consider the following code :

[File A.cuh]

#include <cuda_runtime.h>

struct A
{
    __device__ __host__ void functionA();
};

[File B.cuh]

#include <cuda_runtime.h>

struct B
{
    __device__ __host__ void functionB();
};

[File A.cu]

#include "A.cuh"
#include "B.cuh"

void A::functionA()
{
    B b;
    b.functionB();
}

Attempting to compile A.cu with nvcc -o A.o -c A.cu -arch=sm_20 outputs Error: External calls are not supported (found non-inlined call to _ZN1B9functionBEv).

I must be doing something wrong, but what ?

Answer 1

As explained on this thread on the NVidia forums, it appears that even though Fermi supports non-inlined functions, nvcc still needs to have all the functions available during compilation, i.e. in the same source file: there is no linker (yep, that's a pity...).

Answer 2

functionB is not declared and therefore considered external call. As the error said external calls are not supported. Implement functionB and it will work.