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