CUDA: LNK2005 error on __device__ function used in header file

Question

I have a device function that is defined in a header file. The reason it is in a header file is because it is used by a global kernel, which needs to be in a header file since it is a template kernel.

When this header file is included across 2 or more .cu files, I get a LNK2005 error during linking:

FooDevice.cu.obj : error LNK2005: "int __cdecl getCurThreadIdx(void)" (?getCurThreadIdx@@YAHXZ) already defined in Main.cu.obj

Why is this error caused? How to fix it?

Here is sample code to produces the above error:

FooDevice.h:

#ifndef FOO_DEVICE_H
#define FOO_DEVICE_H

__device__ int getCurThreadIdx()
{
    return ( ( blockIdx.x * blockDim.x ) + threadIdx.x );
}

template< typename T >
__global__ void fooKernel( const T* inArr, int num, T* outArr )
{
    const int threadNum = ( gridDim.x * blockDim.x );

    for ( int idx = getCurThreadIdx(); idx < num; idx += threadNum )
        outArr[ idx ] = inArr[ idx ];

    return;
}

__global__ void fooKernel2( const int* inArr, int num, int* outArr );

#endif // FOO_DEVICE_H

FooDevice.cu:

#include "FooDevice.h"

// One other kernel that uses getCurThreadIdx()
__global__ void fooKernel2( const int* inArr, int num, int* outArr )
{
    const int threadNum = ( gridDim.x * blockDim.x );

    for ( int idx = getCurThreadIdx(); idx < num; idx += threadNum )
        outArr[ idx ] = inArr[ idx ];

    return;
}

Main.cu:

#include "FooDevice.h"

int main()
{
    int num             = 10;
    int* dInArr         = NULL;
    int* dOutArr        = NULL;
    const int arrSize   = num * sizeof( *dInArr );

    cudaMalloc( &dInArr, arrSize );
    cudaMalloc( &dOutArr, arrSize );

    // Using template kernel
    fooKernel<<< 10, 10 >>>( dInArr, num, dOutArr );

    return 0;
}

Answer 1

Why is this error caused?

Because you have included your header in FooDevice.cu and Main.cu where it gets defined so you now have two copies of the same function and the linker detects this.

How to fix it?

If you have the following defined in foo.h

template<typename T> __device__ T foo(T x)
{
    return x;
}

And two .cu files that both include foo.h and also contain a call to it, e.g.

int x = foo<int>(1);

Then you can force foo() inline:

template<typename T>
inline __device__ T foo(T x)
{
    return x;
}

and call:

int x = foo<int>(1);

This will stop it from being declared multiple times.

Function templates are an exempt of One Defintion Rule and may be more than one definition of them in different translation unit's. Full function template specialization is not a template, rather an ordinary function, so you need to use inline keyword not to violate ODR if you want to put them in a header file included into several translation unit's.

Taken from http://www.velocityreviews.com/forums/t447911-why-does-explicit-specialization-of-function-templates-cause-generation-of-code.html

See also: http://en.wikipedia.org/wiki/One_Definition_Rule

I changed your code like this:

inline __device__ int getCurThreadIdx()
{
    return ( ( blockIdx.x * blockDim.x ) + threadIdx.x );
}

template< typename T >
__global__ void fooKernel( const T* inArr, int num, T* outArr )
{
    const int threadNum = ( gridDim.x * blockDim.x );

    for ( int idx = getCurThreadIdx(); idx < num; idx += threadNum )
        outArr[ idx ] = inArr[ idx ];

    return;
}

And it now compiles. Your declaration without the inline of getCurThreadIdx() was violating the one definition rule.