CUDA:头文件中使用的__device__函数的LNK2005错误

Question

我有一个在头文件中定义的设备功能.它在头文件中的原因是因为它由全局内核使用,因为它是模板内核所以需要在头文件中.

当这个头文件包含在2个或更多.cu文件中时,我在链接期间收到LNK2005错误:

FooDevice.cu.obj:错误LNK2005:已在Main.cu.obj中定义的"int __cdecl getCurThreadIdx(void)"(？getCurThreadIdx @@ YAHXZ)

为什么会出现此错误？怎么解决？

以下是产生上述错误的示例代码:

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

为什么会出现此错误？

因为你已经在FooDevice.cu和Main.cu中包含了你的头文件,所以你现在有两个相同功能的副本,链接器检测到这个.

怎么解决？

如果在foo.h中定义了以下内容

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

两个.cu文件都包含foo.h并且还包含对它的调用,例如

int x = foo<int>(1);

然后你可以强制foo()内联:

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

并致电:

int x = foo<int>(1);

这将阻止它被多次声明.

函数模板是One Defintion Rule的免除,并且可以在不同的翻译单元中对它们进行多个定义.完整功能模板专业化不是模板,而是普通函数,因此如果要将它们放在包含在多个翻译单元中的头文件中,则需要使用内联关键字来违反ODR.

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

另见:http://en.wikipedia.org/wiki/One_Definition_Rule

我改变了你的代码:

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

它现在编译.没有getCurThreadIdx()内联的声明违反了一个定义规则.