Bug 1266

POSSIBLE FIX:

After taking a look into the problematic file, changing in Eigen/src/Core/arch/CUDA/Half.h lines 526

from

return Eigen::internal::raw_uint16_to_half(

to

return Eigen::half_impl::raw_uint16_to_half(

seems to work since the function is defined some lines above in this namespace. 


--------------------------------------------------------------------------------


REMAINING COMPILATION PROBLEM:

However, I still have problems to compile Eigen. When calling

#include <Eigen/Core>
#include <Eigen/LU>

Eigen::Matrix<float , 6 , 6> A = <initialize this>;
Eigen::Matrix<float , 6 , 1> b = <initialize this>;
Eigen::Matrix<float , 6 , 1> result = A.fullPivLu().solve(b);

I get the error:

<path to lib folder in project>/Eigen-3.3-beta2/Eigen/src/Core/MatrixBase.h(332): warning: function "Eigen::MatrixBase<Derived>::fullPivLu [with Derived=Eigen::Matrix<float, 6, 6, 0, 6, 6>]" was referenced but not defined

Linking CXX static library <myLib>.a
[ 60%] Built target <myLib>
Linking CXX executable <myAppUsingLib>
<myLib>(<fileWithEigen>.cu.o): undefined reference to `Eigen::MatrixBase<Eigen::Matrix<float, 6, 6, 0, 6, 6> >::fullPivLu() const'
collect2: error: ld returned 1 exit status

The solver seems to have problems now.


Thanks in advance.

I'm sorry to post so frequently, but now I had some time again to test a bit more.

Interesstingly, it does not matter which decomposition/solver I use. Invoking the solver functions
- partialPivLu()
- fullPivLu()
- householderQr()
- colPivHouseholderQr()
- fullPivHouseholderQr()
- llt()
- ldlt()
- jacobiSvd() 
(listed at http://eigen.tuxfamily.org/dox-devel/group__TutorialLinearAlgebra.html), I'm getting the exact same error:

The invoked function was referenced but not defined.

I guess some defines prevent nvcc from reading the definition so its linker gets into trouble. But this is all speculation since I haven't figured out where Eigen actually defines these functions.

Again: Everything works perfectly fine with Eigen 3.2.9 and the code example above should run on the CPU, no kernel or device stuff. It's just defined in a .cu file and must be there.

I can't reproduce your errors, but apparently I only have CUDA 2.1 available.
However, the error and fix regarding raw_uint16_to_half makes sense to me, I'll fixed that as suggested:
https://bitbucket.org/eigen/eigen/commits/cfdf8c8e364faa37a23c58c5

Are you able to build the cuda_basic unit test? For me this compiles and runs also without problems.

I following your suggestion and got the following result:


stotko@computer:<path-to-lib>/Eigen-3.3-beta2/build$ make cuda_basic
Building NVCC (Device) object test/CMakeFiles/cuda_basic.dir//./cuda_basic_generated_cuda_basic.cu.o
<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(60): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(61): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(62): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(63): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(64): warning: invalid error number

<path-to-lib>/Eigen-3.3-beta2/Eigen/src/Core/util/DisableStupidWarnings.h(65): warning: invalid error number

Scanning dependencies of target cuda_basic
Linking CXX executable cuda_basic
Built target cuda_basic
stotko@computer:<path-to-lib>/Eigen-3.3-beta2/build$ ./test/cuda_basic 
Initializing random number generator with seed 1469812041
Repeating each test 10 times
CUDA device info:
  name:                        GeForce GTX 780
  capability:                  3.5
  multiProcessorCount:         12
  maxThreadsPerMultiProcessor: 2048
  warpSize:                    32
  regsPerBlock:                65536
  concurrentKernels:           1
  clockRate:                   901500
  canMapHostMemory:            1
  computeMode:                 0
stotko@computer:<path-to-lib>/Eigen-3.3-beta2/build$ 


Besides the warnings everything looks normal. I also got these warnings in my project, but didn't mention them here.

Ok, the corresponding #pragmas should better be dis/enabled depending on the NVCC-version (you did not need to copy the entire list of warnings ...).
I'm not sure which one is supported by which NVCC version, so I'm cc-ing Benoit (who added them).

If the cuda_basic test compiles and links, then maybe you need to adapt your linking. Or maybe you did not actually put the correct #include inside you .cu files.

Could you attach a (complete) failing example? (I may not be able to tests this before next week, though).

Actually, our cuda_basic test does only test eigenvalues but no other decomposition at the moment. --> test-needed

Created attachment 718 [details]
A failing example

Before running the example:

Please fill the empty Eigen folder with the current (newest?) Eigen version from the defaulkt branch. Bugzilla does not allow me to include this.


What the example contains:
 - CMakeLists.txt : Similar but much simplified structure like in my project
 - main.cpp and main.cu: The failing examples
 - Eigen dummy folder representing my local copy of the library

Expected behavior:
 - Everything compiles and runs fine with trivial output.

Obseved behavior:
 - main.cpp compiles and runs fine with expected output.
 - main.cu fails to compile with the same error like already posted.

Oops, forgot to mention if this might cause problems: I compiled the example the following way.

mkdir build
cd build
cmake ..
make

Pretty standard pipeline, so no special or strange things here.

Thanks again and have a nice weekend.

Actually, I found out how to fix/disable all of my warnings:

--------------------------------------------------------------------------------

First of all, pragma with numbers >=2647 cause the invalid error number warning on NVCC 6.5, so they should only be used for newer versions (maybe >= 7.0 but not sure which one exactly). I also don't know and haven't found how many numbers each version of NVCC supports. Is this documented somewhere?

--------------------------------------------------------------------------------

Suppressing the "calling a __host__ function from a __host__ __device__ function is not allowed" warnings like in DisableStupidWarnings.h works for me with

#pragma diag_suppress 2527
#pragma diag_suppress 2529

I have no idea why exactly these two numbers did the trick since I haven't found a complete list explaining them but only a partial one until number 1786:
http://www.ssl.berkeley.edu/~jimm/grizzly_docs/SSL/opt/intel/cc/9.0/lib/locale/en_US/mcpcom.msg

--------------------------------------------------------------------------------

I also got this warning in my project and the failing example:

<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/util/Memory.h(276): warning: list initialization syntax is a C++11 feature
          detected during:
            instantiation of "Eigen::internal::aligned_stack_memory_handler<T>::aligned_stack_memory_handler(T *, size_t, __nv_bool) [with T=float]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/SolveTriangular.h(66): here
            instantiation of "void Eigen::internal::triangular_solver_selector<Lhs, Rhs, Side, Mode, 0, 1>::run(const Lhs &, Rhs &) [with Lhs=const Eigen::Block<const Eigen::Matrix<float, 6, 6, 0, 6, 6>, -1, -1, false>, Rhs=Eigen::Block<Eigen::Matrix<float, 6, 1, 0, 6, 1>, -1, 1, false>, Side=1, Mode=5]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/SolveTriangular.h(177): here
            instantiation of "void Eigen::TriangularViewImpl<_MatrixType, _Mode, Eigen::Dense>::solveInPlace<Side,OtherDerived>(const Eigen::MatrixBase<OtherDerived> &) const [with _MatrixType=const Eigen::Block<const Eigen::Matrix<float, 6, 6, 0, 6, 6>, -1, -1, false>, _Mode=5U, Side=1, OtherDerived=Eigen::Block<Eigen::Matrix<float, 6, 1, 0, 6, 1>, -1, 1, false>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/TriangularMatrix.h(507): here
            instantiation of "void Eigen::TriangularViewImpl<_MatrixType, _Mode, Eigen::Dense>::solveInPlace(const Eigen::MatrixBase<OtherDerived> &) const [with _MatrixType=const Eigen::Block<const Eigen::Matrix<float, 6, 6, 0, 6, 6>, -1, -1, false>, _Mode=5U, OtherDerived=Eigen::Block<Eigen::Matrix<float, 6, 1, 0, 6, 1>, -1, 1, false>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/LU/FullPivLU.h(777): here
            instantiation of "void Eigen::FullPivLU<_MatrixType>::_solve_impl(const RhsType &, DstType &) const [with _MatrixType=Eigen::Matrix<float, 6, 6, 0, 6, 6>, RhsType=Eigen::Matrix<float, 6, 1, 0, 6, 1>, DstType=Eigen::Matrix<float, 6, 1, 0, 6, 1>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/Solve.h(143): here
            instantiation of "void Eigen::internal::Assignment<DstXprType, Eigen::Solve<DecType, RhsType>, Eigen::internal::assign_op<Scalar, Scalar>, Eigen::internal::Dense2Dense, void>::run(DstXprType &, const Eigen::internal::Assignment<DstXprType, Eigen::Solve<DecType, RhsType>, Eigen::internal::assign_op<Scalar, Scalar>, Eigen::internal::Dense2Dense, void>::SrcXprType &, const Eigen::internal::assign_op<Scalar, Scalar> &) [with DstXprType=Eigen::Matrix<float, 6, 1, 0, 6, 1>, DecType=Eigen::FullPivLU<Eigen::Matrix<float, 6, 6, 0, 6, 6>>, RhsType=Eigen::Matrix<float, 6, 1, 0, 6, 1>, Scalar=float]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/AssignEvaluator.h(813): here
            instantiation of "void Eigen::internal::call_assignment_no_alias(Dst &, const Src &, const Func &) [with Dst=Eigen::Matrix<float, 6, 1, 0, 6, 1>, Src=Eigen::Solve<Eigen::FullPivLU<Eigen::Matrix<float, 6, 6, 0, 6, 6>>, Eigen::Matrix<float, 6, 1, 0, 6, 1>>, Func=Eigen::internal::assign_op<float, float>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/PlainObjectBase.h(721): here
            instantiation of "Derived &Eigen::PlainObjectBase<Derived>::_set_noalias(const Eigen::DenseBase<OtherDerived> &) [with Derived=Eigen::Matrix<float, 6, 1, 0, 6, 1>, OtherDerived=Eigen::Solve<Eigen::FullPivLU<Eigen::Matrix<float, 6, 6, 0, 6, 6>>, Eigen::Matrix<float, 6, 1, 0, 6, 1>>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/PlainObjectBase.h(531): here
            instantiation of "Eigen::PlainObjectBase<Derived>::PlainObjectBase(const Eigen::DenseBase<OtherDerived> &) [with Derived=Eigen::Matrix<float, 6, 1, 0, 6, 1>, OtherDerived=Eigen::Solve<Eigen::FullPivLU<Eigen::Matrix<float, 6, 6, 0, 6, 6>>, Eigen::Matrix<float, 6, 1, 0, 6, 1>>]" 
<path-to-lib>/eigen-eigen-ab755b9907bc/Eigen/src/Core/Matrix.h(380): here
            instantiation of "Eigen::Matrix<_Scalar, _Rows, _Cols, _Options, _MaxRows, _MaxCols>::Matrix(const Eigen::EigenBase<OtherDerived> &) [with _Scalar=float, _Rows=6, _Cols=1, _Options=0, _MaxRows=6, _MaxCols=1, OtherDerived=Eigen::Solve<Eigen::FullPivLU<Eigen::Matrix<float, 6, 6, 0, 6, 6>>, Eigen::Matrix<float, 6, 1, 0, 6, 1>>]" 
<path-to-main>/main.cu(31): here

There seems to be a list initialization with C++11 code. It can either be suppressed by 

#pragma diag_suppress 2068

which I found during the tests for the first two pragmas, or the initialization can be made C++98 compliant so that the warning does not appear anymore and is truely fixed. Since Eigen claims to not require C++11, I suggest the second approach to tackle this warning.

--------------------------------------------------------------------------------

I hope this helps improving the CUDA support.

I found the reason for the linking problem. In fact NVCC was not able to find the definition since the preprocessor simply removed it. I suggest to remove lines 882 and 889 from the src/LU/FullPivLU.h

882 #ifndef __CUDACC__
883 template<typename Derived>
884 inline const FullPivLU<typename MatrixBase<Derived>::PlainObject>
885 MatrixBase<Derived>::fullPivLu() const
886 {
887   return FullPivLU<PlainObject>(eval());
888 }
889 #endif

and also similarly in all other decompositions. In addition, one should remove the EIGEN_DEVICE_FUNC qualifiers from all decomposition function declarations in lines 331-339 of src/Core/MatrixBase.h.

As far as I see, the decompositions haven't any support for being executed inside a device kernel since there are no EIGEN_DEVICE_FUNC qualifiers elsewhere in the repsective files. Applying above changes would however restore host (CPU) compatibility and at the same time prevent any (unsupported) calls on the device (that feature might be added in the future if possible).

What remains is the last unresolved C++11 list initialization warning. But I couldn't figure out myself why it has triggered since the respective line (src/Core/util/Memory.h(276)) is simply a EIGEN_THROW() macro. Also the other lines in the stacktrace look normal to me, but I might have missed something.

After some research, I finally have found the reason for the C++11 list initialization warning. The lines 957+958 in Eigen/src/Core/util/Macros.h define the EIGEN_THROW and EIGEN_THROUGH_X(X) macros both making use of

return {}

to handle the possible return types of device functions. Unfortunately, any change including the removal of this statement, returning to void (return;), etc. triggers other warnings in my setup.

I would suggest to disable this warning if C++11 in NVCC is not available (version 6.5 and below) using

#pragma diag_suppress 2068

and do not touch the EIGEN_THROW macros.

--------------------------------------------------------------------------------

I have also thought about the addtional tests needed to avoid future occurances of this bug. The only solution without copying all existing code into .cu files would be to force NVCC to handle .cpp files like .cu files.

http://stackoverflow.com/questions/36021233/cmake-cuda-compile-cpp-files-in-cuda-mode-x-cu

The respective compiler flag --x=cu however seems not to work with CMake. Only CMake 3.3+ supports this option via the set_source_files_properties() function opertaing on file-level.

This would guarantee that Eigen works with NVCC on the host (CPU) that should always be the case. But I'm not totally convinced whether this is the best option to tackle the problem. However copying the all test code to .cu files is even worse.

Sorry for only looking at this now.
I can confirm that your test-case does not work with 3.3 and that removing the
  #ifndef __CUDACC__
resolves the issue.
@Gael: was there a reason for those? Otherwise, I think we should remove them.

Regarding the unit-test: Can't we just
  #include "unit-test.cpp"
inside a .cu-file for each unit-test we want to test?

Regarding the warnings, what about this:
  #define EIGEN_THROW_X(X) asm("trap;"); while(1) {}

Your version of the EIGEN_THROW macro compiles in the failing example on my machine without the C++11 warning. Never thought of such a solution, but seems to work.


IMHO, including the unit test in .cu-files makes sense. I never considered this idea because including .cpp-files is in general bad style and should be avoided. However in this particular situation it seems reasonable.


The pragma problem is a bit harder. NVCC defines macros for its version only in CUDA 7.5 and newer, see page 9 in the release notes:

http://developer.download.nvidia.com/compute/cuda/7.5/Prod/docs/sidebar/CUDA_Toolkit_Release_Notes.pdf

Therefore I'm not sure what should be done here. I assume that the defined pragmas are correct for newer versions of NVCC, but in version 6.5 this does not work as mentioned before. On the other hand, we cannot rely on the macro since this requires at least version 7.5. Any ideas?

Regarding the linking issue:

https://bitbucket.org/eigen/eigen/commits/fb1f76ac81b9/
Summary:     Bug 1266: remove CUDA guards on MatrixBase::<decomposition> definitions. (those used to break old nvcc versions that we propably don't care anymore)

For the others, it's difficult for me to test with all nvcc versions. I only have the 7.5 and 8.0 versions available.

Your fix works fine, thank you very much.

Regarding the C++11 warning: Christoph's solution works for me.

The other warnings are indeed difficult. Unfortunately, Nvidia defined the version macro in version 7.5 and above, so we can only distinguish between <7.5, 7.5, 8.0, etc.
So I would suggest to use the currently set pragmas only for versions 7.5 and 8.0.
For older ones, you could either use my suggestion (comment 9, only tested with 6.5) or you can completely leave them out.

(In reply to Patrick Stotko from comment #4)
I've added warning 1222 to the list of warning messages to disable: this suppresses the "warning: invalid error number" messages that you see with older versions of nvcc.

(In reply to Patrick Stotko from comment #9)
I've added warning  2527 and 2529 to the list of messages to suppress.

There is no documentation for error/warning numbers: the only way I found to figure out what to disable is to compile with the -Xcudafe --display_error_number option to force nvcc to print the message numbers alongside the warning.

Unfortunately the message numbers associated with the "calling a __host__ function from a __host__ __device__ function is not allowed"" warnings appear to change with almost every revision of nvcc, so we need to extend the list as we try to compile with more revisions of the compiler.

I can confirm that the warnings related to "invalid error number" and "calling a __host__ function from a __host__ __device__ function is not allowed" are gone with the latest commits, thanks.

NVCC is now only complaining about the C++11 warning. Everything else works fine.

I tried to take care of the EIGEN_THROW warnings in https://bitbucket.org/eigen/eigen/commits/cd0cf5e95bb91b1827daad9618d649e09977d70d. Let me know if this works for you as well.

The last fix works very well. No warnings anymore, thank you Benoit.


Before I would consider this as resolved fixed, I have one last concern. The fix from Gael regarding the link errors might not fully solved the problem. If you look at these lines in Eigen/src/Core/MatrixBase.h

331 EIGEN_DEVICE_FUNC
332 inline const FullPivLU<PlainObject> fullPivLu() const;
333 EIGEN_DEVICE_FUNC
334 inline const PartialPivLU<PlainObject> partialPivLu() const;
335 
336 EIGEN_DEVICE_FUNC
337 inline const PartialPivLU<PlainObject> lu() const;
338 
339 EIGEN_DEVICE_FUNC
340 inline const Inverse<Derived> inverse() const;

you will notice the EIGEN_DEVICE_FUNC qualifiers. But they do not appear in the definitions of the functions, see Eigen/src/LU/FullPivLU.h 

882 template<typename Derived>
883 inline const FullPivLU<typename MatrixBase<Derived>::PlainObject>
884 MatrixBase<Derived>::fullPivLu() const
885 {
886   return FullPivLU<PlainObject>(eval());
887 }

and Eigen/src/LU/FullPivLU.h. So if I would use them inside a kernel, I would expect some error or warning again. I haven't tested it so far, but intuitively there is a problem here.

This only affects these three functions from the LU module. I'm not sure whether inverse() is also affected since I haven't found its definition inside the LU module. The others do have the qualifier and are fine.

Oops, there is a word missing on my last sentence. Of course I mean:

The others do NOT have the qualifier and are fine.

I have added the missing EIGEN_DEVICE_FUNC qualifiers in the function definitions for fullPivLu(), partialPivLu(), lu(), and inverse(). I have also added a few extra missing EIGEN_DEVICE_FUNC.
I have also checked that your example compiles and links correctly with cuda 8.0. I don't have cuda 6.5 installed anymore, so let me know if this fixes also work for you.

I'm sorry, but my comment about the qualifier was a bit misleading. I expected the qualifier to be removed from the declaration instead of adding it to the definition.

Considering the other decompositions, there are also no qualifiers there, so I think that Eigen does not and should not allow them being called inside a kernel (at least for 3.3) due to internal memory allocations or whatever used magic might cause problems on the device.

Furthermore, the qualifier is also missing in the constructor of FullPivLu and the others. This underlines that invocations on the device are currently not intended and users should not try this in their code.

So I suggest to remove the EIGEN_DEVICE_FUNC qualifiers.
As a side note: My example still compiles fine with your changes on my machine. But trying to do this on the device does not compile as I expected.

I removed the offending EIGEN_DEVICE_FUNC qualifiers in https://bitbucket.org/eigen/eigen/commits/36aba87af81ce6697354037b33dbb2633e985027. This code is indeed not ready to run on GPU at the moment.

thank you, I guess this closes this entry.

-- GitLab Migration Automatic Message --

This bug has been migrated to gitlab.com's GitLab instance and has been closed from further activity.

You can subscribe and participate further through the new bug through this link to our GitLab instance: https://gitlab.com/libeigen/eigen/issues/1266.