ptxas 文件中的 CUDA 外部类链接和未解析的外部函数

Question

我正在使用 CUDA，我创建了一个 int2_ 类来处理复整数。

ComplexTypes.h文件中的类声明如下：

namespace LibraryNameSpace
{
    class int2_ {

        public:
            int x;
            int y;

            // Constructors
            __host__ __device__ int2_(const int,const int);
            __host__ __device__ int2_();
            // etc.

            // Equalities with other types      
            __host__ __device__ const int2_& operator=(const int);
            __host__ __device__ const int2_& operator=(const float);
            // etc.

    };
}

ComplexTypes.cpp文件中的类实现如下：

#include "ComplexTypes.h"

__host__ __device__         LibraryNameSpace::int2_::int2_(const int x_,const int y_)           { x=x_; y=y_;}
__host__ __device__         LibraryNameSpace::int2_::int2_() {}
// etc.

__host__ __device__ const   LibraryNameSpace::int2_& LibraryNameSpace::int2_::operator=(const int a)                        { x = a;            y = 0.;             return *this; }
__host__ __device__ const   LibraryNameSpace::int2_& LibraryNameSpace::int2_::operator=(const float a)                      { x = (int)a;       y = 0.;             return *this; }
// etc.

一切正常。在main（包括ComplexTypes.h）中，我可以处理int2_ 数字。

在CudaMatrix.cu 文件中，我现在包含ComplexTypes.h 并定义和正确实例化__global__ 函数：

template <class T1, class T2>
__global__ void evaluation_matrix(T1* data_, T2* ob, int NumElements)
{
    const int i = blockDim.x * blockIdx.x + threadIdx.x;
    if(i < NumElements) data_[i] = ob[i];
}

template __global__ void evaluation_matrix(LibraryNameSpace::int2_*,int*,int);

CudaMatrix.cu 文件的情况似乎与main 函数对称。然而，编译器抱怨：

Error   19  error : Unresolved extern function '_ZN16LibraryNameSpace5int2_aSEi'    C:\Users\Documents\Project\Test\Testing_Files\ptxas simpleTest

请考虑：

1. 在将实现移至单独的文件之前，在 main 文件中包含声明和实现时，一切正常。
2. 有问题的指令是data_[i] = ob[i]。

有人知道发生了什么吗？

Answer 1

我在上面的帖子中遵循的程序有两个问题：

1. ComplexTypes.cpp 文件名必须转为ComplexTypes.cu，以便nvcc 可以拦截CUDA 关键字__device__ 和__host__。 Talonmies 在他的评论中指出了这一点。实际上，在发布之前，我已经将文件名从 .cpp 更改为 .cu，但编译器抱怨并显示相同的错误。于是，我就老老实实的往后退了一步；

2. 在 Visual Studio 2010 中，必须使用 View -> Property Pages；配置属性 -> CUDA C/C++ -> 通用 -> 生成可重定位设备代码 -> 是 (-rdc=true)。这是单独编译所必需的。确实，在NVIDIA CUDA Compiler Driver NVCC，据说：

CUDA 的工作原理是将设备代码嵌入到主机对象中。在整个程序编译中，它将可执行设备代码嵌入到宿主对象中。在单独编译中，我们将可重定位设备代码嵌入到宿主对象中，并运行设备链接器（nvlink）将所有设备代码链接在一起。然后通过主机链接器将 nvlink 的输出与所有主机对象链接在一起，形成最终的可执行文件。 relocatable vs 可执行设备代码的生成由 --relocatable-device-code={true,false} 选项控制，可简写为 - rdc={true,false}.