GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-09-19 04:43:13 +08:00
Code Issues Packages Projects Releases Wiki Activity

[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.

Browse Source
This commit is contained in:
mehdi-goli 2020-11-02 11:13:37 +00:00 committed by Rasmus Munk Larsen
parent 6961468915
commit e24a1f57e3
4 changed files with 50 additions and 33 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
CMakeLists.txt
View File

@ -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) option(EIGEN_SYCL_TRISYCL "Use the triSYCL Sycl implementation (ComputeCPP by default)." OFF)
if(EIGEN_TEST_SYCL) if(EIGEN_TEST_SYCL)
set (CMAKE_MODULE_PATH "${CMAKE_ROOT}/Modules" "cmake/Modules/" "${CMAKE_MODULE_PATH}") set (CMAKE_MODULE_PATH "${CMAKE_ROOT}/Modules" "cmake/Modules/" "${CMAKE_MODULE_PATH}")
find_package(Threads REQUIRED)
if(EIGEN_SYCL_TRISYCL) if(EIGEN_SYCL_TRISYCL)
message(STATUS "Using triSYCL") message(STATUS "Using triSYCL")
include(FindTriSYCL) include(FindTriSYCL)

75
Eigen/src/Core/GenericPacketMath.h
View File

@ -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). /** \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. */ If \a a or \b b is NaN, the return value is implementation defined. */
template<typename Packet> EIGEN_DEVICE_FUNC inline Packet 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). /** \internal \returns the min of \a a and \a b (coeff-wise).
NaNPropagation determines the NaN propagation semantics. */ NaNPropagation determines the NaN propagation semantics. */
template<int NaNPropagation, typename Packet> EIGEN_DEVICE_FUNC inline Packet template <int NaNPropagation, typename Packet>
pmin(const Packet& a, const Packet& b) { return pminmax_impl<NaNPropagation>::run(a,b, pmin<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) /** \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. */ 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). /** \internal \returns the max of \a a and \a b (coeff-wise).
NaNPropagation determines the NaN propagation semantics. */ NaNPropagation determines the NaN propagation semantics. */
template<int NaNPropagation, typename Packet> EIGEN_DEVICE_FUNC inline Packet template <int NaNPropagation, typename Packet>
pmax(const Packet& a, const Packet& b) { return pminmax_impl<NaNPropagation>::run(a,b, pmax<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 */ /** \internal \returns the absolute value of \a a */
template<typename Packet> EIGEN_DEVICE_FUNC inline Packet 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 */ /** \internal \returns the product of the elements of \a a */
template<typename Packet> template <typename Packet>
EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_mul(
predux_mul(const Packet& a) const Packet& a) {
{ typedef typename unpacket_traits<Packet>::type Scalar;
return predux_helper(a, pmul<typename unpacket_traits<Packet>::type>); return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmul<Scalar>)));
} }
/** \internal \returns the min of the elements of \a a */ /** \internal \returns the min of the elements of \a a */
template<typename Packet> template <typename Packet>
EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_min(
predux_min(const Packet& a) const Packet &a) {
{ typedef typename unpacket_traits<Packet>::type Scalar;
return predux_helper(a, pmin<PropagateFast, typename unpacket_traits<Packet>::type>); return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin<PropagateFast, Scalar>)));
} }
template<int NaNPropagation, typename Packet> template <int NaNPropagation, typename Packet>
EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_min(
predux_min(const Packet& a) const Packet& a) {
{ typedef typename unpacket_traits<Packet>::type Scalar;
return predux_helper(a, pmin<NaNPropagation, typename unpacket_traits<Packet>::type>); return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmin<NaNPropagation, Scalar>)));
} }
/** \internal \returns the max of the elements of \a a */ /** \internal \returns the min of the elements of \a a */
template<typename Packet> template <typename Packet>
EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_max(
predux_max(const Packet& a) const Packet &a) {
{ typedef typename unpacket_traits<Packet>::type Scalar;
return predux_helper(a, pmax<PropagateFast, typename unpacket_traits<Packet>::type>); return predux_helper(a, EIGEN_BINARY_OP_NAN_PROPAGATION(Scalar, (pmax<PropagateFast, Scalar>)));
} }
template<int NaNPropagation, typename Packet> template <int NaNPropagation, typename Packet>
EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type EIGEN_DEVICE_FUNC inline typename unpacket_traits<Packet>::type predux_max(
predux_max(const Packet& a) const Packet& a) {
{ typedef typename unpacket_traits<Packet>::type Scalar;
return predux_helper(a, pmax<NaNPropagation, typename unpacket_traits<Packet>::type>); 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" /** \internal \returns true if all coeffs of \a a means "true"
* It is supposed to be called on values returned by pcmp_*. * It is supposed to be called on values returned by pcmp_*.
*/ */

3
Eigen/src/Core/arch/SYCL/InteropHeaders.h
View File

@ -59,6 +59,9 @@ struct sycl_packet_traits : default_packet_traits {
HasIGammac = 0, HasIGammac = 0,
HasBetaInc = 0, HasBetaInc = 0,
HasBlend = has_blend, 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, HasMax = 1,
HasMin = 1, HasMin = 1,
HasMul = 1, HasMul = 1,

4
Eigen/src/Core/arch/SYCL/MathFunctions.h
View File

@ -20,7 +20,6 @@
#ifndef EIGEN_MATH_FUNCTIONS_SYCL_H #ifndef EIGEN_MATH_FUNCTIONS_SYCL_H
#define EIGEN_MATH_FUNCTIONS_SYCL_H #define EIGEN_MATH_FUNCTIONS_SYCL_H
namespace Eigen { namespace Eigen {
namespace internal { namespace internal {
@ -238,7 +237,7 @@ SYCL_PROUND(cl::sycl::cl_double2)
#undef SYCL_PROUND #undef SYCL_PROUND
#define SYCL_PRINT(packet_type) \ #define SYCL_PRINT(packet_type) \
template<> \ template <> \
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type print<packet_type>( \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type print<packet_type>( \
const packet_type& a) { \ const packet_type& a) { \
return cl::sycl::rint(a); \ return cl::sycl::rint(a); \
@ -295,7 +294,6 @@ SYCL_PLDEXP(cl::sycl::cl_double2)
#undef SYCL_PLDEXP #undef SYCL_PLDEXP
#endif #endif
} // end namespace internal } // end namespace internal
} // end namespace Eigen } // end namespace Eigen