From e24a1f57e35f3f3894a5612bb8b4e34bf68ebb26 Mon Sep 17 00:00:00 2001 From: mehdi-goli Date: Mon, 2 Nov 2020 11:13:37 +0000 Subject: [SYCL Function pointer Issue]: SYCL does not support function pointer inside the kernel, due to the portability issue of a function pointer and memory address space among host and accelerators. To fix the issue, function pointers have been replaced by function objects. --- CMakeLists.txt | 1 + Eigen/src/Core/GenericPacketMath.h | 75 ++++++++++++++++++------------- Eigen/src/Core/arch/SYCL/InteropHeaders.h | 3 ++ 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 { } }; + +#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 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 EIGEN_DEVICE_FUNC inline Packet -pmin(const Packet& a, const Packet& b) { return pminmax_impl::run(a,b, pmin); } +template +EIGEN_DEVICE_FUNC inline Packet pmin(const Packet& a, const Packet& b) { + return pminmax_impl::run(a, b, EIGEN_BINARY_OP_NAN_PROPAGATION(Packet, (pmin))); +} /** \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 EIGEN_DEVICE_FUNC inline Packet -pmax(const Packet& a, const Packet& b) { return pminmax_impl::run(a,b, pmax); } +template +EIGEN_DEVICE_FUNC inline Packet pmax(const Packet& a, const Packet& b) { + return pminmax_impl::run(a, b, EIGEN_BINARY_OP_NAN_PROPAGATION(Packet,(pmax))); +} /** \internal \returns the absolute value of \a a */ template EIGEN_DEVICE_FUNC inline Packet @@ -705,43 +718,45 @@ predux(const Packet& a) } /** \internal \returns the product of the elements of \a a */ -template -EIGEN_DEVICE_FUNC inline typename unpacket_traits::type -predux_mul(const Packet& a) -{ - return predux_helper(a, pmul::type>); +template +EIGEN_DEVICE_FUNC inline typename unpacket_traits::type predux_mul( + const Packet& a) { + typedef typename unpacket_traits::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmul))); } /** \internal \returns the min of the elements of \a a */ -template -EIGEN_DEVICE_FUNC inline typename unpacket_traits::type -predux_min(const Packet& a) -{ - return predux_helper(a, pmin::type>); +template +EIGEN_DEVICE_FUNC inline typename unpacket_traits::type predux_min( + const Packet &a) { + typedef typename unpacket_traits::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin))); } -template -EIGEN_DEVICE_FUNC inline typename unpacket_traits::type -predux_min(const Packet& a) -{ - return predux_helper(a, pmin::type>); +template +EIGEN_DEVICE_FUNC inline typename unpacket_traits::type predux_min( + const Packet& a) { + typedef typename unpacket_traits::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin))); } -/** \internal \returns the max of the elements of \a a */ -template -EIGEN_DEVICE_FUNC inline typename unpacket_traits::type -predux_max(const Packet& a) -{ - return predux_helper(a, pmax::type>); +/** \internal \returns the min of the elements of \a a */ +template +EIGEN_DEVICE_FUNC inline typename unpacket_traits::type predux_max( + const Packet &a) { + typedef typename unpacket_traits::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmax))); } -template -EIGEN_DEVICE_FUNC inline typename unpacket_traits::type -predux_max(const Packet& a) -{ - return predux_helper(a, pmax::type>); +template +EIGEN_DEVICE_FUNC inline typename unpacket_traits::type predux_max( + const Packet& a) { + typedef typename unpacket_traits::type Scalar; + return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmax))); } +#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( \ 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 -- cgit v1.2.3