diff options
| -rw-r--r-- | CMakeLists.txt | 1 | ||||
| -rw-r--r-- | Eigen/src/Core/GenericPacketMath.h | 75 | ||||
| -rw-r--r-- | Eigen/src/Core/arch/SYCL/InteropHeaders.h | 3 | ||||
| -rw-r--r-- | Eigen/src/Core/arch/SYCL/MathFunctions.h | 4 |
4 files changed, 50 insertions, 33 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt index a6ab9de25..1340b04f3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -498,6 +498,7 @@ option(EIGEN_TEST_SYCL "Add Sycl support." OFF) option(EIGEN_SYCL_TRISYCL "Use the triSYCL Sycl implementation (ComputeCPP by default)." OFF) if(EIGEN_TEST_SYCL) set (CMAKE_MODULE_PATH "${CMAKE_ROOT}/Modules" "cmake/Modules/" "${CMAKE_MODULE_PATH}") + find_package(Threads REQUIRED) if(EIGEN_SYCL_TRISYCL) message(STATUS "Using triSYCL") include(FindTriSYCL) diff --git a/Eigen/src/Core/GenericPacketMath.h b/Eigen/src/Core/GenericPacketMath.h index fad94535f..e2fc7002b 100644 --- a/Eigen/src/Core/GenericPacketMath.h +++ b/Eigen/src/Core/GenericPacketMath.h @@ -364,6 +364,15 @@ struct pminmax_impl<PropagateNumbers> { } }; + +#ifndef SYCL_DEVICE_ONLY +#define EIGEN_BINARY_OP_NAN_PROPAGATION(Type, Func) Func +#else +#define EIGEN_BINARY_OP_NAN_PROPAGATION(Type, Func) \ +[](const Type& a, const Type& b) { \ + return Func(a, b);} +#endif + /** \internal \returns the min of \a a and \a b (coeff-wise). If \a a or \b b is NaN, the return value is implementation defined. */ template<typename Packet> EIGEN_DEVICE_FUNC inline Packet @@ -371,8 +380,10 @@ pmin(const Packet& a, const Packet& b) { return numext::mini(a,b); } /** \internal \returns the min of \a a and \a b (coeff-wise). NaNPropagation determines the NaN propagation semantics. */ -template<int NaNPropagation, typename Packet> EIGEN_DEVICE_FUNC inline Packet -pmin(const Packet& a, const Packet& b) { return pminmax_impl<NaNPropagation>::run(a,b, pmin<Packet>); } +template <int NaNPropagation, typename Packet> +EIGEN_DEVICE_FUNC inline Packet pmin(const Packet& a, const Packet& b) { + return pminmax_impl<NaNPropagation>::run(a, b, EIGEN_BINARY_OP_NAN_PROPAGATION(Packet, (pmin<Packet>))); +} /** \internal \returns the max of \a a and \a b (coeff-wise) If \a a or \b b is NaN, the return value is implementation defined. */ @@ -381,8 +392,10 @@ pmax(const Packet& a, const Packet& b) { return numext::maxi(a, b); } /** \internal \returns the max of \a a and \a b (coeff-wise). NaNPropagation determines the NaN propagation semantics. */ -template<int NaNPropagation, typename Packet> EIGEN_DEVICE_FUNC inline Packet -pmax(const Packet& a, const Packet& b) { return pminmax_impl<NaNPropagation>::run(a,b, pmax<Packet>); } +template <int NaNPropagation, typename Packet> +EIGEN_DEVICE_FUNC inline Packet pmax(const Packet& a, const Packet& b) { + return pminmax_impl<NaNPropagation>::run(a, b, EIGEN_BINARY_OP_NAN_PROPAGATION(Packet,(pmax<Packet>))); +} /** \internal \returns the absolute value of \a a */ template<typename Packet> EIGEN_DEVICE_FUNC inline Packet @@ -705,43 +718,45 @@ predux(const Packet& a) } /** \internal \returns the product of the elements of \a a */ -template<typename Packet> -EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type -predux_mul(const Packet& a) -{ - return predux_helper(a, pmul<typename unpacket_traits<Packet>::type>); +template <typename Packet> +EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_mul( + const Packet& a) { + typedef typename unpacket_traits<Packet>::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmul<Scalar>))); } /** \internal \returns the min of the elements of \a a */ -template<typename Packet> -EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type -predux_min(const Packet& a) -{ - return predux_helper(a, pmin<PropagateFast, typename unpacket_traits<Packet>::type>); +template <typename Packet> +EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_min( + const Packet &a) { + typedef typename unpacket_traits<Packet>::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin<PropagateFast, Scalar>))); } -template<int NaNPropagation, typename Packet> -EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type -predux_min(const Packet& a) -{ - return predux_helper(a, pmin<NaNPropagation, typename unpacket_traits<Packet>::type>); +template <int NaNPropagation, typename Packet> +EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_min( + const Packet& a) { + typedef typename unpacket_traits<Packet>::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin<NaNPropagation, Scalar>))); } -/** \internal \returns the max of the elements of \a a */ -template<typename Packet> -EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type -predux_max(const Packet& a) -{ - return predux_helper(a, pmax<PropagateFast, typename unpacket_traits<Packet>::type>); +/** \internal \returns the min of the elements of \a a */ +template <typename Packet> +EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_max( + const Packet &a) { + typedef typename unpacket_traits<Packet>::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmax<PropagateFast, Scalar>))); } -template<int NaNPropagation, typename Packet> -EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type -predux_max(const Packet& a) -{ - return predux_helper(a, pmax<NaNPropagation, typename unpacket_traits<Packet>::type>); +template <int NaNPropagation, typename Packet> +EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_max( + const Packet& a) { + typedef typename unpacket_traits<Packet>::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmax<NaNPropagation, Scalar>))); } +#undef EIGEN_BINARY_OP_NAN_PROPAGATION + /** \internal \returns true if all coeffs of \a a means "true" * It is supposed to be called on values returned by pcmp_*. */ diff --git a/Eigen/src/Core/arch/SYCL/InteropHeaders.h b/Eigen/src/Core/arch/SYCL/InteropHeaders.h index 710059d50..10856ff5e 100644 --- a/Eigen/src/Core/arch/SYCL/InteropHeaders.h +++ b/Eigen/src/Core/arch/SYCL/InteropHeaders.h @@ -59,6 +59,9 @@ struct sycl_packet_traits : default_packet_traits { HasIGammac = 0, HasBetaInc = 0, HasBlend = has_blend, + // This flag is used to indicate whether packet comparison is supported. + // pcmp_eq, pcmp_lt and pcmp_le should be defined for it to be true. + HasCmp = 1, HasMax = 1, HasMin = 1, HasMul = 1, diff --git a/Eigen/src/Core/arch/SYCL/MathFunctions.h b/Eigen/src/Core/arch/SYCL/MathFunctions.h index 7d8cd132b..2ab0f2a76 100644 --- a/Eigen/src/Core/arch/SYCL/MathFunctions.h +++ b/Eigen/src/Core/arch/SYCL/MathFunctions.h @@ -20,7 +20,6 @@ #ifndef EIGEN_MATH_FUNCTIONS_SYCL_H #define EIGEN_MATH_FUNCTIONS_SYCL_H - namespace Eigen { namespace internal { @@ -238,7 +237,7 @@ SYCL_PROUND(cl::sycl::cl_double2) #undef SYCL_PROUND #define SYCL_PRINT(packet_type) \ - template<> \ + template <> \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type print<packet_type>( \ const packet_type& a) { \ return cl::sycl::rint(a); \ @@ -295,7 +294,6 @@ SYCL_PLDEXP(cl::sycl::cl_double2) #undef SYCL_PLDEXP #endif - } // end namespace internal } // end namespace Eigen |