CUDA：头文件中使用的 __device__ 函数出现 LNK2005 错误

Question

我有一个在头文件中定义的设备函数。它位于头文件中的原因是因为它由全局内核使用，而该内核需要位于头文件中，因为它是模板内核。

当此头文件包含在 2 个或更多 .cu 文件中时，我在链接过程中收到 LNK2005 错误：

FooDevice.cu.obj：错误 LNK2005：“int __cdecl getCurThreadIdx(void)" (?getCurThreadIdx@@YAHXZ) 已定义 在Main.cu.obj中

为什么会出现这个错误？如何修复它？

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

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() inline:

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

并调用：

int x = foo<int>(1);

这将阻止它被声明为多个次。

函数模板不受
一个定义规则，可能更多
比它们的一个定义
不同的翻译单位。满的
函数模板专业化是
不是模板，而是普通的
函数，所以你需要使用内联
如果需要，关键字不要违反 ODR
将它们放入包含的头文件中
分成几个翻译单元。

取自 http://www.velocityreviews.com/forums/t447911-why-does-explicit-specialization-of-function-templates-cause-genesis-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() 违反了单一定义规则。

Answer 2

它应该内联。您可以尝试添加 inline 关键字。

也许您可以删除不必要的代码并创建一个简单的文本示例供我们查看？通常问题出在细节上...