More Eigen GPU work, with fix for CUDA 11.4

Eigen/src/Core/PlainObjectBase.h

@@ -641,69 +641,91 @@ class PlainObjectBase : public internal::dense_xpr_base<Derived>::type
       * \see class Map
       */
     //@{
+    EIGEN_DEVICE_FUNC
     static inline ConstMapType Map(const Scalar* data)
     { return ConstMapType(data); }
+    EIGEN_DEVICE_FUNC
     static inline MapType Map(Scalar* data)
     { return MapType(data); }
+    EIGEN_DEVICE_FUNC
     static inline ConstMapType Map(const Scalar* data, Index size)
     { return ConstMapType(data, size); }
     EIGEN_DEVICE_FUNC
     static inline MapType Map(Scalar* data, Index size)
     { return MapType(data, size); }
+    EIGEN_DEVICE_FUNC
     static inline ConstMapType Map(const Scalar* data, Index rows, Index cols)
     { return ConstMapType(data, rows, cols); }
     EIGEN_DEVICE_FUNC
     static inline MapType Map(Scalar* data, Index rows, Index cols)
     { return MapType(data, rows, cols); }
 
+    EIGEN_DEVICE_FUNC
     static inline ConstAlignedMapType MapAligned(const Scalar* data)
     { return ConstAlignedMapType(data); }
+    EIGEN_DEVICE_FUNC
     static inline AlignedMapType MapAligned(Scalar* data)
     { return AlignedMapType(data); }
+    EIGEN_DEVICE_FUNC
     static inline ConstAlignedMapType MapAligned(const Scalar* data, Index size)
     { return ConstAlignedMapType(data, size); }
+    EIGEN_DEVICE_FUNC
     static inline AlignedMapType MapAligned(Scalar* data, Index size)
     { return AlignedMapType(data, size); }
+    EIGEN_DEVICE_FUNC
     static inline ConstAlignedMapType MapAligned(const Scalar* data, Index rows, Index cols)
     { return ConstAlignedMapType(data, rows, cols); }
+    EIGEN_DEVICE_FUNC
     static inline AlignedMapType MapAligned(Scalar* data, Index rows, Index cols)
     { return AlignedMapType(data, rows, cols); }
 
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstMapType<Stride<Outer, Inner> >::type Map(const Scalar* data, const Stride<Outer, Inner>& stride)
     { return typename StridedConstMapType<Stride<Outer, Inner> >::type(data, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedMapType<Stride<Outer, Inner> >::type Map(Scalar* data, const Stride<Outer, Inner>& stride)
     { return typename StridedMapType<Stride<Outer, Inner> >::type(data, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstMapType<Stride<Outer, Inner> >::type Map(const Scalar* data, Index size, const Stride<Outer, Inner>& stride)
     { return typename StridedConstMapType<Stride<Outer, Inner> >::type(data, size, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedMapType<Stride<Outer, Inner> >::type Map(Scalar* data, Index size, const Stride<Outer, Inner>& stride)
     { return typename StridedMapType<Stride<Outer, Inner> >::type(data, size, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstMapType<Stride<Outer, Inner> >::type Map(const Scalar* data, Index rows, Index cols, const Stride<Outer, Inner>& stride)
     { return typename StridedConstMapType<Stride<Outer, Inner> >::type(data, rows, cols, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedMapType<Stride<Outer, Inner> >::type Map(Scalar* data, Index rows, Index cols, const Stride<Outer, Inner>& stride)
     { return typename StridedMapType<Stride<Outer, Inner> >::type(data, rows, cols, stride); }
 
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type MapAligned(const Scalar* data, const Stride<Outer, Inner>& stride)
     { return typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type(data, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedAlignedMapType<Stride<Outer, Inner> >::type MapAligned(Scalar* data, const Stride<Outer, Inner>& stride)
     { return typename StridedAlignedMapType<Stride<Outer, Inner> >::type(data, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type MapAligned(const Scalar* data, Index size, const Stride<Outer, Inner>& stride)
     { return typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type(data, size, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedAlignedMapType<Stride<Outer, Inner> >::type MapAligned(Scalar* data, Index size, const Stride<Outer, Inner>& stride)
     { return typename StridedAlignedMapType<Stride<Outer, Inner> >::type(data, size, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type MapAligned(const Scalar* data, Index rows, Index cols, const Stride<Outer, Inner>& stride)
     { return typename StridedConstAlignedMapType<Stride<Outer, Inner> >::type(data, rows, cols, stride); }
     template<int Outer, int Inner>
+    EIGEN_DEVICE_FUNC
     static inline typename StridedAlignedMapType<Stride<Outer, Inner> >::type MapAligned(Scalar* data, Index rows, Index cols, const Stride<Outer, Inner>& stride)
     { return typename StridedAlignedMapType<Stride<Outer, Inner> >::type(data, rows, cols, stride); }
     //@}

Eigen/src/Core/Solve.h

@@ -65,6 +65,7 @@ class Solve : public SolveImpl<Decomposition,RhsType,typename internal::traits<R
   typedef typename internal::traits<Solve>::PlainObject PlainObject;
   typedef typename internal::traits<Solve>::StorageIndex StorageIndex;
 
+  EIGEN_DEVICE_FUNC
   Solve(const Decomposition &dec, const RhsType &rhs)
     : m_dec(dec), m_rhs(rhs)
   {}

Eigen/src/Core/SolverBase.h

@@ -17,6 +17,7 @@ namespace internal {
 template<typename Derived>
 struct solve_assertion {
     template<bool Transpose_, typename Rhs>
+    EIGEN_DEVICE_FUNC
     static void run(const Derived& solver, const Rhs& b) { solver.template _check_solve_assertion<Transpose_>(b); }
 };
 
@@ -26,6 +27,7 @@ struct solve_assertion<Transpose<Derived> >
     typedef Transpose<Derived> type;
 
     template<bool Transpose_, typename Rhs>
+    EIGEN_DEVICE_FUNC
     static void run(const type& transpose, const Rhs& b)
     {
         internal::solve_assertion<typename internal::remove_all<Derived>::type>::template run<true>(transpose.nestedExpression(), b);
@@ -38,6 +40,7 @@ struct solve_assertion<CwiseUnaryOp<Eigen::internal::scalar_conjugate_op<Scalar>
     typedef CwiseUnaryOp<Eigen::internal::scalar_conjugate_op<Scalar>, const Transpose<Derived> > type;
 
     template<bool Transpose_, typename Rhs>
+    EIGEN_DEVICE_FUNC
     static void run(const type& adjoint, const Rhs& b)
     {
         internal::solve_assertion<typename internal::remove_all<Transpose<Derived> >::type>::template run<true>(adjoint.nestedExpression(), b);
@@ -105,6 +108,7 @@ class SolverBase : public EigenBase<Derived>
       */
     template<typename Rhs>
     inline const Solve<Derived, Rhs>
+    EIGEN_DEVICE_FUNC
     solve(const MatrixBase<Rhs>& b) const
     {
       internal::solve_assertion<typename internal::remove_all<Derived>::type>::template run<false>(derived(), b);
@@ -147,6 +151,7 @@ class SolverBase : public EigenBase<Derived>
   protected:
 
     template<bool Transpose_, typename Rhs>
+    EIGEN_DEVICE_FUNC
     void _check_solve_assertion(const Rhs& b) const {
         EIGEN_ONLY_USED_FOR_DEBUG(b);
         eigen_assert(derived().m_isInitialized && "Solver is not initialized.");

Eigen/src/Core/products/GeneralBlockPanelKernel.h

@@ -432,10 +432,10 @@ template <typename Packet>
 struct QuadPacket
 {
   Packet B_0, B1, B2, B3;
-  const Packet& get(const FixedInt<0>&) const { return B_0; }
-  const Packet& get(const FixedInt<1>&) const { return B1; }
-  const Packet& get(const FixedInt<2>&) const { return B2; }
-  const Packet& get(const FixedInt<3>&) const { return B3; }
+  EIGEN_DEVICE_FUNC const Packet& get(const FixedInt<0>&) const { return B_0; }
+  EIGEN_DEVICE_FUNC const Packet& get(const FixedInt<1>&) const { return B1; }
+  EIGEN_DEVICE_FUNC const Packet& get(const FixedInt<2>&) const { return B2; }
+  EIGEN_DEVICE_FUNC const Packet& get(const FixedInt<3>&) const { return B3; }
 };
 
 template <int N, typename T1, typename T2, typename T3>
@@ -1240,7 +1240,8 @@ struct last_row_process_16_packets
   typedef typename SwappedTraits::ResPacket SResPacket;
   typedef typename SwappedTraits::AccPacket SAccPacket;
 
-  EIGEN_STRONG_INLINE void operator()(const DataMapper& res, SwappedTraits &straits, const LhsScalar* blA,
+  EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
+  void operator()(const DataMapper& res, SwappedTraits &straits, const LhsScalar* blA,
                   const RhsScalar* blB, Index depth, const Index endk, Index i, Index j2,
                   ResScalar alpha, SAccPacket &C0)
     {

Eigen/src/Core/util/IntegralConstant.h

@@ -52,41 +52,56 @@ template<int N> class FixedInt
 {
 public:
   static const int value = N;
-  operator int() const { return value; }
-  FixedInt() {}
-  FixedInt( VariableAndFixedInt<N> other) {
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC operator int() const { return value; }
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC FixedInt() {}
+#if EIGEN_HAS_CXX14
+  EIGEN_CONSTEXPR
+#endif
+  EIGEN_DEVICE_FUNC FixedInt(VariableAndFixedInt<N> other) {
     #ifndef EIGEN_INTERNAL_DEBUGGING
     EIGEN_UNUSED_VARIABLE(other);
     #endif
     eigen_internal_assert(int(other)==N);
   }
 
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<-N> operator-() const { return FixedInt<-N>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N+M> operator+( FixedInt<M>) const { return FixedInt<N+M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N-M> operator-( FixedInt<M>) const { return FixedInt<N-M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N*M> operator*( FixedInt<M>) const { return FixedInt<N*M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N/M> operator/( FixedInt<M>) const { return FixedInt<N/M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N%M> operator%( FixedInt<M>) const { return FixedInt<N%M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N|M> operator|( FixedInt<M>) const { return FixedInt<N|M>(); }
   template<int M>
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt<N&M> operator&( FixedInt<M>) const { return FixedInt<N&M>(); }
 
 #if EIGEN_HAS_CXX14
   // Needed in C++14 to allow fix<N>():
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt operator() () const { return *this; }
 
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   VariableAndFixedInt<N> operator() (int val) const { return VariableAndFixedInt<N>(val); }
 #else
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt ( FixedInt<N> (*)() ) {}
 #endif
 
 #if EIGEN_HAS_CXX11
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
   FixedInt(std::integral_constant<int,N>) {}
 #endif
 };
@@ -124,8 +139,8 @@ template<int N> class VariableAndFixedInt
 {
 public:
   static const int value = N;
-  operator int() const { return m_value; }
-  VariableAndFixedInt(int val) { m_value = val; }
+  EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC operator int() const { return m_value; }
+  EIGEN_DEVICE_FUNC VariableAndFixedInt(int val) { m_value = val; }
 protected:
   int m_value;
 };
@@ -153,9 +168,14 @@ struct get_fixed_value<variable_if_dynamic<T,N>,Default> {
   static const int value = N;
 };
 
-template<typename T> EIGEN_DEVICE_FUNC Index get_runtime_value(const T &x) { return x; }
+template<typename T>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
+Index get_runtime_value(const T &x) { return x; }
+
 #if !EIGEN_HAS_CXX14
-template<int N> EIGEN_DEVICE_FUNC Index get_runtime_value(FixedInt<N> (*)()) { return N; }
+template<int N>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
+Index get_runtime_value(FixedInt<N> (*)()) { return N; }
 #endif
 
 // Cleanup integer/FixedInt/VariableAndFixedInt/etc types:
@@ -186,14 +206,17 @@ template<int N, int DynamicKey> struct cleanup_index_type<std::integral_constant
 
 #if EIGEN_HAS_CXX14
 template<int N>
+EIGEN_DEVICE_CONST
 static const internal::FixedInt<N> fix{};
 #else
 template<int N>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
 inline internal::FixedInt<N> fix() { return internal::FixedInt<N>(); }
 
 // The generic typename T is mandatory. Otherwise, a code like fix<N> could refer to either the function above or this next overload.
 // This way a code like fix<N> can only refer to the previous function.
 template<int N,typename T>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
 inline internal::VariableAndFixedInt<N> fix(T val) { return internal::VariableAndFixedInt<N>(internal::convert_index<int>(val)); }
 #endif
 
@@ -232,6 +255,7 @@ inline internal::VariableAndFixedInt<N> fix(T val) { return internal::VariableAn
   * \sa fix<N>(int), seq, seqN
   */
 template<int N>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
 static const auto fix();
 
 /** \fn fix<N>(int)
@@ -263,6 +287,7 @@ static const auto fix();
   * \sa fix, seqN, class ArithmeticSequence
   */
 template<int N>
+EIGEN_CONSTEXPR EIGEN_DEVICE_FUNC
 static const auto fix(int val);
 
 #endif // EIGEN_PARSED_BY_DOXYGEN

Eigen/src/Core/util/Macros.h

@@ -970,6 +970,12 @@
   #define EIGEN_DEVICE_FUNC
 #endif
 
+// Define compile-time constants for use by both the host and device functions
+#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIP_DEVICE_COMPILE)
+  #define EIGEN_DEVICE_CONST __device__
+#else
+  #define EIGEN_DEVICE_CONST
+#endif
 
 // this macro allows to get rid of linking errors about multiply defined functions.
 //  - static is not very good because it prevents definitions from different object files to be merged.