9

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 ?

talonmies
  • 70,661
  • 34
  • 192
  • 269
user703016
  • 37,307
  • 8
  • 87
  • 112

3 Answers3

10

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...).

Luc Touraille
  • 79,925
  • 15
  • 92
  • 137
1

True, CUDA 5.0 does it. I can't get it to expose external device variables but device methods work just fine. Not by default.

The nvcc option is "-rdc=true". In Visual Studio and Nsight it is an option in the project properties under Configuration Properties -> CUDA C/C++ -> Common -> Generate Relocatable Device Code.

Boyko Karadzhov
  • 566
  • 4
  • 9
1

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

Daniel
  • 30,896
  • 18
  • 85
  • 139