diff options
Diffstat (limited to 'Eigen/src/Core/arch/SYCL/MathFunctions.h')
| -rw-r--r-- | Eigen/src/Core/arch/SYCL/MathFunctions.h | 207 |
1 files changed, 132 insertions, 75 deletions
diff --git a/Eigen/src/Core/arch/SYCL/MathFunctions.h b/Eigen/src/Core/arch/SYCL/MathFunctions.h index 422839c6c..9e16e6c3f 100644 --- a/Eigen/src/Core/arch/SYCL/MathFunctions.h +++ b/Eigen/src/Core/arch/SYCL/MathFunctions.h @@ -16,7 +16,7 @@ * \brief: * MathFunctions * -*****************************************************************/ + *****************************************************************/ #ifndef EIGEN_MATH_FUNCTIONS_SYCL_H #define EIGEN_MATH_FUNCTIONS_SYCL_H @@ -28,194 +28,251 @@ namespace internal { // Make sure this is only available when targeting a GPU: we don't want to // introduce conflicts between these packet_traits definitions and the ones // we'll use on the host side (SSE, AVX, ...) -//#if defined(__SYCL_DEVICE_ONLY__) && defined(EIGEN_USE_SYCL) -#define SYCL_PLOG(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type plog<packet_type>(const packet_type& a) { return cl::sycl::log(a); } +#if defined(SYCL_DEVICE_ONLY) +#define SYCL_PLOG(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::log(a); \ + } SYCL_PLOG(cl::sycl::cl_float4) SYCL_PLOG(cl::sycl::cl_double2) #undef SYCL_PLOG -#define SYCL_PLOG1P(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type plog1p<packet_type>(const packet_type& a) { return cl::sycl::log1p(a); } +#define SYCL_PLOG1P(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog1p<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::log1p(a); \ + } SYCL_PLOG1P(cl::sycl::cl_float4) SYCL_PLOG1P(cl::sycl::cl_double2) #undef SYCL_PLOG1P -#define SYCL_PLOG10(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type plog10<packet_type>(const packet_type& a) { return cl::sycl::log10(a); } +#define SYCL_PLOG10(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog10<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::log10(a); \ + } SYCL_PLOG10(cl::sycl::cl_float4) SYCL_PLOG10(cl::sycl::cl_double2) #undef SYCL_PLOG10 -#define SYCL_PEXP(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pexp<packet_type>(const packet_type& a) { return cl::sycl::exp(a); } +#define SYCL_PEXP(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pexp<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::exp(a); \ + } SYCL_PEXP(cl::sycl::cl_float4) SYCL_PEXP(cl::sycl::cl_double2) #undef SYCL_PEXP -#define SYCL_PEXPM1(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pexpm1<packet_type>(const packet_type& a) { return cl::sycl::expm1(a); } +#define SYCL_PEXPM1(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pexpm1<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::expm1(a); \ + } SYCL_PEXPM1(cl::sycl::cl_float4) SYCL_PEXPM1(cl::sycl::cl_double2) #undef SYCL_PEXPM1 -#define SYCL_PSQRT(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type psqrt<packet_type>(const packet_type& a) { return cl::sycl::sqrt(a); } +#define SYCL_PSQRT(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psqrt<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::sqrt(a); \ + } SYCL_PSQRT(cl::sycl::cl_float4) SYCL_PSQRT(cl::sycl::cl_double2) #undef SYCL_PSQRT - -#define SYCL_PRSQRT(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type prsqrt<packet_type>(const packet_type& a) { return cl::sycl::rsqrt(a); } +#define SYCL_PRSQRT(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type prsqrt<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::rsqrt(a); \ + } SYCL_PRSQRT(cl::sycl::cl_float4) SYCL_PRSQRT(cl::sycl::cl_double2) #undef SYCL_PRSQRT - /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */ -#define SYCL_PSIN(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type psin<packet_type>(const packet_type& a) { return cl::sycl::sin(a); } +#define SYCL_PSIN(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psin<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::sin(a); \ + } SYCL_PSIN(cl::sycl::cl_float4) SYCL_PSIN(cl::sycl::cl_double2) #undef SYCL_PSIN - /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */ -#define SYCL_PCOS(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pcos<packet_type>(const packet_type& a) { return cl::sycl::cos(a); } +#define SYCL_PCOS(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pcos<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::cos(a); \ + } SYCL_PCOS(cl::sycl::cl_float4) SYCL_PCOS(cl::sycl::cl_double2) #undef SYCL_PCOS /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */ -#define SYCL_PTAN(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type ptan<packet_type>(const packet_type& a) { return cl::sycl::tan(a); } +#define SYCL_PTAN(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ptan<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::tan(a); \ + } SYCL_PTAN(cl::sycl::cl_float4) SYCL_PTAN(cl::sycl::cl_double2) #undef SYCL_PTAN /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */ -#define SYCL_PASIN(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pasin<packet_type>(const packet_type& a) { return cl::sycl::asin(a); } +#define SYCL_PASIN(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pasin<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::asin(a); \ + } SYCL_PASIN(cl::sycl::cl_float4) SYCL_PASIN(cl::sycl::cl_double2) #undef SYCL_PASIN - /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */ -#define SYCL_PACOS(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pacos<packet_type>(const packet_type& a) { return cl::sycl::acos(a); } +#define SYCL_PACOS(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pacos<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::acos(a); \ + } SYCL_PACOS(cl::sycl::cl_float4) SYCL_PACOS(cl::sycl::cl_double2) #undef SYCL_PACOS /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */ -#define SYCL_PATAN(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type patan<packet_type>(const packet_type& a) { return cl::sycl::atan(a); } +#define SYCL_PATAN(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type patan<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::atan(a); \ + } SYCL_PATAN(cl::sycl::cl_float4) SYCL_PATAN(cl::sycl::cl_double2) #undef SYCL_PATAN /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */ -#define SYCL_PSINH(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type psinh<packet_type>(const packet_type& a) { return cl::sycl::sinh(a); } +#define SYCL_PSINH(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psinh<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::sinh(a); \ + } SYCL_PSINH(cl::sycl::cl_float4) SYCL_PSINH(cl::sycl::cl_double2) #undef SYCL_PSINH /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */ -#define SYCL_PCOSH(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pcosh<packet_type>(const packet_type& a) { return cl::sycl::cosh(a); } +#define SYCL_PCOSH(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pcosh<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::cosh(a); \ + } SYCL_PCOSH(cl::sycl::cl_float4) SYCL_PCOSH(cl::sycl::cl_double2) #undef SYCL_PCOSH /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */ -#define SYCL_PTANH(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type ptanh<packet_type>(const packet_type& a) { return cl::sycl::tanh(a); } +#define SYCL_PTANH(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ptanh<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::tanh(a); \ + } SYCL_PTANH(cl::sycl::cl_float4) SYCL_PTANH(cl::sycl::cl_double2) #undef SYCL_PTANH -#define SYCL_PCEIL(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pceil<packet_type>(const packet_type& a) { return cl::sycl::ceil(a); } +#define SYCL_PCEIL(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pceil<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::ceil(a); \ + } SYCL_PCEIL(cl::sycl::cl_float4) SYCL_PCEIL(cl::sycl::cl_double2) #undef SYCL_PCEIL - -#define SYCL_PROUND(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pround<packet_type>(const packet_type& a) { return cl::sycl::round(a); } +#define SYCL_PROUND(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pround<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::round(a); \ + } SYCL_PROUND(cl::sycl::cl_float4) SYCL_PROUND(cl::sycl::cl_double2) #undef SYCL_PROUND -#define SYCL_FLOOR(packet_type) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pfloor<packet_type>(const packet_type& a) { return cl::sycl::floor(a); } +#define SYCL_FLOOR(packet_type) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pfloor<packet_type>( \ + const packet_type& a) { \ + return cl::sycl::floor(a); \ + } SYCL_FLOOR(cl::sycl::cl_float4) SYCL_FLOOR(cl::sycl::cl_double2) #undef SYCL_FLOOR - -#define SYCL_PMIN(packet_type, expr) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pmin<packet_type>(const packet_type& a, const packet_type& b) { return expr; } +#define SYCL_PMIN(packet_type, expr) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pmin<packet_type>( \ + const packet_type& a, const packet_type& b) { \ + return expr; \ + } SYCL_PMIN(cl::sycl::cl_float4, cl::sycl::fmin(a, b)) SYCL_PMIN(cl::sycl::cl_double2, cl::sycl::fmin(a, b)) #undef SYCL_PMIN -#define SYCL_PMAX(packet_type, expr) \ -template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE \ -packet_type pmax<packet_type>(const packet_type& a, const packet_type& b) { return expr; } +#define SYCL_PMAX(packet_type, expr) \ + template <> \ + EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pmax<packet_type>( \ + const packet_type& a, const packet_type& b) { \ + return expr; \ + } SYCL_PMAX(cl::sycl::cl_float4, cl::sycl::fmax(a, b)) SYCL_PMAX(cl::sycl::cl_double2, cl::sycl::fmax(a, b)) #undef SYCL_PMAX -//#endif +#endif -} // end namespace internal +} // end namespace internal -} // end namespace Eigen +} // end namespace Eigen -#endif // EIGEN_MATH_FUNCTIONS_CUDA_H +#endif // EIGEN_MATH_FUNCTIONS_SYCL_H |