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

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++ - > Common - >生成可重定位设备代码 - >是(-rdc = true).这对于单独编译是必要的.实际上,在NVIDIA CUDA编译器驱动程序NVCC中,据说:

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