CUDA: LNK2005 Fehler bei __device__ Funktion, die in Header-Datei verwendet wird

Question

Ich habe eine Geräte Funktion, die in einer Header-Datei definiert ist. Der Grund, warum sie in einer Header-Datei steht, ist, dass sie von einem globalen Kernel verwendet wird, der in einer Header-Datei stehen muss, da es sich um einen Template-Kernel handelt.

Wenn diese Header-Datei in 2 oder mehr .cu-Dateien eingebunden wird, erhalte ich einen LNK2005-Fehler während des Linkens:

FooDevice.cu.obj : Fehler LNK2005: "int __cdecl getCurThreadIdx(void)" (?getCurThreadIdx@@YAHXZ) bereits definiert in Main.cu.obj

Warum wird dieser Fehler verursacht? Wie kann er behoben werden?

Hier ist ein Beispielcode, der den oben genannten Fehler erzeugt:

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"

// Ein weiterer Kernel, der getCurThreadIdx() verwendet
__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 );

    // Verwendung des Template-Kernels
    fooKernel<<< 10, 10 >>>( dInArr, num, dOutArr );

    return 0;
}

Answer 1

Warum tritt dieser Fehler auf?

Weil du deinen Header in FooDevice.cu und Main.cu eingefügt hast, wo er definiert wird. Jetzt hast du zwei Kopien der gleichen Funktion und der Linker erkennt dies.

Wie behebt man das?

Wenn du das Folgende in foo.h definiert hast

template __device__ T foo(T x)
{
    return x;
}

Und zwei .cu-Dateien, die beide foo.h einschließen und auch einen Aufruf enthalten, z.B.

int x = foo(1);

Dann kannst du foo() erzwingen, inline zu sein:

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

und aufrufen:

int x = foo(1);

Dadurch wird verhindert, dass es mehrmals deklariert wird.

Funktionstemplates sind von der One Definition Rule ausgenommen und es kann mehr als eine Definition von ihnen in verschiedenen Übersetzungseinheiten geben. Eine vollständige Funktionstemplatespezialisierung ist kein Template, sondern eine gewöhnliche Funktion. Deshalb musst du das inline-Schlüsselwort verwenden, um die ODR nicht zu verletzen, wenn du sie in eine Headerdatei einfügst, die in mehrere Übersetzungseinheiten eingebunden ist.

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

Siehe auch: http://en.wikipedia.org/wiki/One_Definition_Rule

Ich habe deinen Code so geändert:

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;
}

Und jetzt kompiliert es. Deine Deklaration ohne das Inline von getCurThreadIdx() hat die One Definition Rule verletzt.

Answer 2

Es sollte inliniert werden. Du könntest versuchen, das inline Schlüsselwort hinzuzufügen.

Vielleicht könntest du den unnötigen Code entfernen und ein einfaches Textbeispiel erstellen, damit wir es sehen können? Normalerweise liegt das Problem in den Details...