NVCC Compile Error when calling Eigen function in kernel code

Question

I am using Eigen 3.3.9 for my project running on Ubuntu 18.04.4. Everything worked well before I tried to modify my project for CUDA support. I've tried cuda/10.0 and cuda/10.2.

My problem occurs when I try to call determinant() for a Eigen Matrix. My compiler produces the following errors.

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

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

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

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

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

and a few more rows of that same log.

I am pretty sure that Eigen is supported since I haven't gotten any errors for other Eigen functions. I see warnings about device and host annotations ignored when compiling, but haven't seen determinant in one of those.

./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

Here is part of the code that produces error.

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

Any help is appreciated. I am going to implement my own determinant function if this keep on happening.

EDIT: This is the minimal code I can reduce to reproduce the issue.

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

I have my Eigen folder extracted right in the working directory, and included it with

nvcc add.cu -I .

NVCC Compile Error when calling Eigen function in kernel code

0 Answers