Breaking change for Eigen code in NVCC 11.3 and later vs NVCC 11.2 and earlier

Hello,

starting with CUDA 11.3.0 and up to and including CUDA 11.4.1, nvcc complains about

error: identifier "Eigen::fix<(int)0> " is undefined in device code

while compiling a recent (modified) version of Eigen.
Earlier versions of CUDA up to 11.2.2 are working fine.

Here is a simple way to reproduce the problem:

git clone https://gitlab.com/fwyzard/eigen.git -b extend_CUDA_support_v1

cat > test.cu <<@EOF
#include <Eigen/Core>
#include <Eigen/Eigenvalues>

using Matrix5d = Eigen::Matrix<double, 5, 5>;

__global__
void kernel(Matrix5d const *in, Matrix5d *out)
{
  (*out) = in->inverse();
}
@EOF

/usr/local/cuda-11.4.1/bin/nvcc -c -std=c++17 -O3 --expt-relaxed-constexpr --expt-extended-lambda -Ieigen -DEIGEN_DONT_PARALLELIZE --diag-suppress 20012,20014 test.cu -o test.cu.o

which fails with

eigen/Eigen/src/LU/PartialPivLU.h(412): error: identifier "Eigen::fix<(int)-1> " is undefined in device code
eigen/Eigen/src/LU/PartialPivLU.h(422): error: identifier "Eigen::fix<(int)-1> " is undefined in device code
eigen/Eigen/src/LU/PartialPivLU.h(422): error: identifier "Eigen::fix<(int)-1> " is undefined in device code
eigen/Eigen/src/LU/PartialPivLU.h(422): error: identifier "Eigen::fix<(int)-1> " is undefined in device code
eigen/Eigen/src/LU/PartialPivLU.h(422): error: identifier "Eigen::fix<(int)-1> " is undefined in device code
eigen/Eigen/src/Core/products/GeneralBlockPanelKernel.h(1617): error: identifier "Eigen::fix<(int)0> " is undefined in device code
eigen/Eigen/src/Core/products/GeneralBlockPanelKernel.h(1617): error: identifier "Eigen::fix<(int)0> " is undefined in device code
eigen/Eigen/src/Core/products/GeneralBlockPanelKernel.h(1617): error: identifier "Eigen::fix<(int)0> " is undefined in device code
eigen/Eigen/src/Core/products/GeneralBlockPanelKernel.h(1617): error: identifier "Eigen::fix<(int)1> " is undefined in device code
...
Error limit reached.
100 errors detected in the compilation of "test.cu".
Compilation terminated.

while the same works with CUDA up to version 11.2.2:

/usr/local/cuda-11.2.2/bin/nvcc -c -std=c++17 -O3 --expt-relaxed-constexpr --expt-extended-lambda -Ieigen -DEIGEN_DONT_PARALLELIZE --diag-suppress 20012,20014 test.cu -o test.cu.o
/usr/local/cuda-11.2.2/bin/cuobjdump -lelf test.cu.o
ELF file 1: test.cu.1.sm_52.cubin

Is this a known issue ?

Is it a new bug in NVCC – or an old bug in Eigen that was not caught by NVCC earlier ?

Thanks you,
.Andrea

While trying to write a simpler test case, I may have figured out what the problem is.

In C++14 and later mode, Eigen declares Eigen::fix<N> as

template<int N>
static const internal::FixedInt<N> fix{};

If I make the same declaration available for device code as

template<int N>
#ifdef __CUDA_ARCH__
__device__
#endif
static const internal::FixedInt<N> fix{};

then the code compiles and runs fine.

So, back to my earlier question: is this a bug in CUDA 11.3 and later? Or is it a bug in Eigen, that earlier versions of NVCC failed to notice ?

Thank you,
.Andrea

Hi dear customer ,
Could you please file us a ticket following the instruction here Getting Help with CUDA NVCC Compiler . We will take a look.

Of course: https://developer.nvidia.com/nvidia_bug/3361469 .