在内核代码中调用 Eigen 函数时出现 NVCC 编译错误

Question

我在 Ubuntu 18.04.4 上运行的项目使用 Eigen 3.3.9。在我尝试修改我的项目以获得 CUDA 支持之前，一切都运行良好。我尝试过 cuda/10.0 和 cuda/10.2。

当我尝试调用特征矩阵的行列式（）时，出现问题。我的编译器产生以下错误。

voxelize.cu(301): error: calling a __host__ function("Eigen::MatrixBase< ::Eigen::Matrix<float, (int)3, (int)3, (int)0, (int)3, (int)3> > ::determinant const") from a __global__ function("MeshIterKernel") is not allowed

voxelize.cu(301): error: identifier "Eigen::MatrixBase< ::Eigen::Matrix<float, (int)3, (int)3, (int)0, (int)3, (int)3> > ::determinant const" is undefined in device code

voxelize.cu(301): error: calling a __host__ function("Eigen::MatrixBase< ::Eigen::Matrix<float, (int)3, (int)3, (int)0, (int)3, (int)3> > ::determinant const") from a __global__ function("MeshIterKernel") is not allowed

voxelize.cu(301): error: identifier "Eigen::MatrixBase< ::Eigen::Matrix<float, (int)3, (int)3, (int)0, (int)3, (int)3> > ::determinant const" is undefined in device code

voxelize.cu(301): error: calling a __host__ function("Eigen::MatrixBase< ::Eigen::Matrix<float, (int)3, (int)3, (int)0, (int)3, (int)3> > ::determinant const") from a __global__ function("MeshIterKernel") is not allowed

还有同一根日志的几行。

我非常确定支持 Eigen，因为我没有收到其他 Eigen 函数的任何错误。我看到编译时忽略有关 device 和 host 注释的警告，但在其中之一中没有看到决定因素。

./Eigen/src/Core/Block.h(341): warning: __host__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

./Eigen/src/Core/Block.h(341): warning: __device__ annotation is ignored on a function("BlockImpl_dense") that is explicitly defaulted on its first declaration

./Eigen/src/Core/Transpose.h(66): warning: __host__ annotation is ignored on a function("Transpose") that is explicitly defaulted on its first declaration

./Eigen/src/Core/Transpose.h(66): warning: __device__ annotation is ignored on a function("Transpose") that is explicitly defaulted on its first declaration

这是产生错误的部分代码。

__device__
bool Philipp(Vector3f A, Vector3f B, Vector3f C, Vector3f D, Vector3f P)
{
    Vector3f a = A - P, b = B - P, c = C - P, d = D - P;
    Matrix3f Da, Db, Dc, Dd;
    Da << b, c, d;
    Db << a, c, d;
    Dc << a, b, d;
    Dd << a, b, c;

    if((Da.determinant() * Dc.determinant() >= 0) && (Db.determinant() * Dd.determinant() >= 0) && (Da.determinant() * Db.determinant() <= 0))
        return true;
    return false;
}

如有任何帮助，我们将不胜感激。如果这种情况继续发生，我将实现我自己的行列式函数。

编辑： 这是我可以减少以重现该问题的最少代码。

#include <Eigen/Dense>
using namespace Eigen;

__device__
bool Foo()
{
    Vector3f a, b, c;
    Matrix3f Da;
    a << 1,2,3;
    b << 2,5,7;
    c << 6,1,8;
    Da << a, b, c;
    return Da.determinant() >= 0;
}

__global__ 
void Kernel()
{
    Foo();
}

int main(int argc, char** argv)
{
    Kernel<<<1,1>>>();
    return 0;
}

我将 Eigen 文件夹提取到工作目录中，并将其包含在

nvcc add.cu -I .

Answer 1

我最近也遇到同样的问题。我认为这可能是因为 metric() 函数没有在 Eigen 库中定义为 CUDA 设备函数。

在 Eigen/src/LU/Determinant.h 中的 determinant() 函数之前添加 EIGEN_DEVICE_FUNC 应该可以解决此问题。我尝试了这个并且有效。

以下代码显示了Determinant.h的一部分。

/** \lu_module
  *
  * \returns the determinant of this matrix
  */
template<typename Derived>
EIGEN_DEVICE_FUNC   //←← add here
inline typename internal::traits<Derived>::Scalar MatrixBase<Derived>::determinant() const
{
  eigen_assert(rows() == cols());
  typedef typename internal::nested_eval<Derived,Base::RowsAtCompileTime>::type Nested;
  return internal::determinant_impl<typename internal::remove_all<Nested>::type>::run(derived());
}