arch/SYCL/MathFunctions.h

0001 // This file is part of Eigen, a lightweight C++ template library
0002 // for linear algebra.
0003 //
0004 // Mehdi Goli    Codeplay Software Ltd.
0005 // Ralph Potter  Codeplay Software Ltd.
0006 // Luke Iwanski  Codeplay Software Ltd.
0007 // Contact: <eigen@codeplay.com>
0008 //
0009 // This Source Code Form is subject to the terms of the Mozilla
0010 // Public License v. 2.0. If a copy of the MPL was not distributed
0011 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
0012
0013 /*****************************************************************
0014  * MathFunctions.h
0015  *
0016  * \brief:
0017  *  MathFunctions
0018  *
0019  *****************************************************************/
0020
0021 #ifndef EIGEN_MATH_FUNCTIONS_SYCL_H
0022 #define EIGEN_MATH_FUNCTIONS_SYCL_H
0023 namespace Eigen {
0024
0025 namespace internal {
0026
0027 // Make sure this is only available when targeting a GPU: we don't want to
0028 // introduce conflicts between these packet_traits definitions and the ones
0029 // we'll use on the host side (SSE, AVX, ...)
0030 #if defined(SYCL_DEVICE_ONLY)
0031 #define SYCL_PLOG(packet_type)                                         \
0032   template <>                                                          \
0033   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog<packet_type>( \
0034       const packet_type& a) {                                          \
0035     return cl::sycl::log(a);                                           \
0036   }
0037
0038 SYCL_PLOG(cl::sycl::cl_float4)
0039 SYCL_PLOG(cl::sycl::cl_double2)
0040 #undef SYCL_PLOG
0041
0042 #define SYCL_PLOG1P(packet_type)                                         \
0043   template <>                                                            \
0044   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog1p<packet_type>( \
0045       const packet_type& a) {                                            \
0046     return cl::sycl::log1p(a);                                           \
0047   }
0048
0049 SYCL_PLOG1P(cl::sycl::cl_float4)
0050 SYCL_PLOG1P(cl::sycl::cl_double2)
0051 #undef SYCL_PLOG1P
0052
0053 #define SYCL_PLOG10(packet_type)                                         \
0054   template <>                                                            \
0055   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type plog10<packet_type>( \
0056       const packet_type& a) {                                            \
0057     return cl::sycl::log10(a);                                           \
0058   }
0059
0060 SYCL_PLOG10(cl::sycl::cl_float4)
0061 SYCL_PLOG10(cl::sycl::cl_double2)
0062 #undef SYCL_PLOG10
0063
0064 #define SYCL_PEXP(packet_type)                                         \
0065   template <>                                                          \
0066   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pexp<packet_type>( \
0067       const packet_type& a) {                                          \
0068     return cl::sycl::exp(a);                                           \
0069   }
0070
0071 SYCL_PEXP(cl::sycl::cl_float4)
0072 SYCL_PEXP(cl::sycl::cl_float)
0073 SYCL_PEXP(cl::sycl::cl_double2)
0074 #undef SYCL_PEXP
0075
0076 #define SYCL_PEXPM1(packet_type)                                         \
0077   template <>                                                            \
0078   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pexpm1<packet_type>( \
0079       const packet_type& a) {                                            \
0080     return cl::sycl::expm1(a);                                           \
0081   }
0082
0083 SYCL_PEXPM1(cl::sycl::cl_float4)
0084 SYCL_PEXPM1(cl::sycl::cl_double2)
0085 #undef SYCL_PEXPM1
0086
0087 #define SYCL_PSQRT(packet_type)                                         \
0088   template <>                                                           \
0089   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psqrt<packet_type>( \
0090       const packet_type& a) {                                           \
0091     return cl::sycl::sqrt(a);                                           \
0092   }
0093
0094 SYCL_PSQRT(cl::sycl::cl_float4)
0095 SYCL_PSQRT(cl::sycl::cl_double2)
0096 #undef SYCL_PSQRT
0097
0098 #define SYCL_PRSQRT(packet_type)                                         \
0099   template <>                                                            \
0100   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type prsqrt<packet_type>( \
0101       const packet_type& a) {                                            \
0102     return cl::sycl::rsqrt(a);                                           \
0103   }
0104
0105 SYCL_PRSQRT(cl::sycl::cl_float4)
0106 SYCL_PRSQRT(cl::sycl::cl_double2)
0107 #undef SYCL_PRSQRT
0108
0109 /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */
0110 #define SYCL_PSIN(packet_type)                                         \
0111   template <>                                                          \
0112   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psin<packet_type>( \
0113       const packet_type& a) {                                          \
0114     return cl::sycl::sin(a);                                           \
0115   }
0116
0117 SYCL_PSIN(cl::sycl::cl_float4)
0118 SYCL_PSIN(cl::sycl::cl_double2)
0119 #undef SYCL_PSIN
0120
0121 /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */
0122 #define SYCL_PCOS(packet_type)                                         \
0123   template <>                                                          \
0124   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pcos<packet_type>( \
0125       const packet_type& a) {                                          \
0126     return cl::sycl::cos(a);                                           \
0127   }
0128
0129 SYCL_PCOS(cl::sycl::cl_float4)
0130 SYCL_PCOS(cl::sycl::cl_double2)
0131 #undef SYCL_PCOS
0132
0133 /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */
0134 #define SYCL_PTAN(packet_type)                                         \
0135   template <>                                                          \
0136   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ptan<packet_type>( \
0137       const packet_type& a) {                                          \
0138     return cl::sycl::tan(a);                                           \
0139   }
0140
0141 SYCL_PTAN(cl::sycl::cl_float4)
0142 SYCL_PTAN(cl::sycl::cl_double2)
0143 #undef SYCL_PTAN
0144
0145 /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */
0146 #define SYCL_PASIN(packet_type)                                         \
0147   template <>                                                           \
0148   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pasin<packet_type>( \
0149       const packet_type& a) {                                           \
0150     return cl::sycl::asin(a);                                           \
0151   }
0152
0153 SYCL_PASIN(cl::sycl::cl_float4)
0154 SYCL_PASIN(cl::sycl::cl_double2)
0155 #undef SYCL_PASIN
0156
0157 /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */
0158 #define SYCL_PACOS(packet_type)                                         \
0159   template <>                                                           \
0160   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pacos<packet_type>( \
0161       const packet_type& a) {                                           \
0162     return cl::sycl::acos(a);                                           \
0163   }
0164
0165 SYCL_PACOS(cl::sycl::cl_float4)
0166 SYCL_PACOS(cl::sycl::cl_double2)
0167 #undef SYCL_PACOS
0168
0169 /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */
0170 #define SYCL_PATAN(packet_type)                                         \
0171   template <>                                                           \
0172   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type patan<packet_type>( \
0173       const packet_type& a) {                                           \
0174     return cl::sycl::atan(a);                                           \
0175   }
0176
0177 SYCL_PATAN(cl::sycl::cl_float4)
0178 SYCL_PATAN(cl::sycl::cl_double2)
0179 #undef SYCL_PATAN
0180
0181 /** \internal \returns the hyperbolic sine of \a a (coeff-wise) */
0182 #define SYCL_PSINH(packet_type)                                         \
0183   template <>                                                           \
0184   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type psinh<packet_type>( \
0185       const packet_type& a) {                                           \
0186     return cl::sycl::sinh(a);                                           \
0187   }
0188
0189 SYCL_PSINH(cl::sycl::cl_float4)
0190 SYCL_PSINH(cl::sycl::cl_double2)
0191 #undef SYCL_PSINH
0192
0193 /** \internal \returns the hyperbolic cosine of \a a (coeff-wise) */
0194 #define SYCL_PCOSH(packet_type)                                         \
0195   template <>                                                           \
0196   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pcosh<packet_type>( \
0197       const packet_type& a) {                                           \
0198     return cl::sycl::cosh(a);                                           \
0199   }
0200
0201 SYCL_PCOSH(cl::sycl::cl_float4)
0202 SYCL_PCOSH(cl::sycl::cl_double2)
0203 #undef SYCL_PCOSH
0204
0205 /** \internal \returns the hyperbolic tan of \a a (coeff-wise) */
0206 #define SYCL_PTANH(packet_type)                                         \
0207   template <>                                                           \
0208   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ptanh<packet_type>( \
0209       const packet_type& a) {                                           \
0210     return cl::sycl::tanh(a);                                           \
0211   }
0212
0213 SYCL_PTANH(cl::sycl::cl_float4)
0214 SYCL_PTANH(cl::sycl::cl_double2)
0215 #undef SYCL_PTANH
0216
0217 #define SYCL_PCEIL(packet_type)                                         \
0218   template <>                                                           \
0219   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pceil<packet_type>( \
0220       const packet_type& a) {                                           \
0221     return cl::sycl::ceil(a);                                           \
0222   }
0223
0224 SYCL_PCEIL(cl::sycl::cl_float4)
0225 SYCL_PCEIL(cl::sycl::cl_double2)
0226 #undef SYCL_PCEIL
0227
0228 #define SYCL_PROUND(packet_type)                                         \
0229   template <>                                                            \
0230   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pround<packet_type>( \
0231       const packet_type& a) {                                            \
0232     return cl::sycl::round(a);                                           \
0233   }
0234
0235 SYCL_PROUND(cl::sycl::cl_float4)
0236 SYCL_PROUND(cl::sycl::cl_double2)
0237 #undef SYCL_PROUND
0238
0239 #define SYCL_PRINT(packet_type)                                         \
0240   template <>                                                           \
0241   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type print<packet_type>( \
0242       const packet_type& a) {                                           \
0243     return cl::sycl::rint(a);                                           \
0244   }
0245
0246 SYCL_PRINT(cl::sycl::cl_float4)
0247 SYCL_PRINT(cl::sycl::cl_double2)
0248 #undef SYCL_PRINT
0249
0250 #define SYCL_FLOOR(packet_type)                                          \
0251   template <>                                                            \
0252   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pfloor<packet_type>( \
0253       const packet_type& a) {                                            \
0254     return cl::sycl::floor(a);                                           \
0255   }
0256
0257 SYCL_FLOOR(cl::sycl::cl_float4)
0258 SYCL_FLOOR(cl::sycl::cl_double2)
0259 #undef SYCL_FLOOR
0260
0261 #define SYCL_PMIN(packet_type, expr)                                   \
0262   template <>                                                          \
0263   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pmin<packet_type>( \
0264       const packet_type& a, const packet_type& b) {                    \
0265     return expr;                                                       \
0266   }
0267
0268 SYCL_PMIN(cl::sycl::cl_float4, cl::sycl::fmin(a, b))
0269 SYCL_PMIN(cl::sycl::cl_double2, cl::sycl::fmin(a, b))
0270 #undef SYCL_PMIN
0271
0272 #define SYCL_PMAX(packet_type, expr)                                   \
0273   template <>                                                          \
0274   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pmax<packet_type>( \
0275       const packet_type& a, const packet_type& b) {                    \
0276     return expr;                                                       \
0277   }
0278
0279 SYCL_PMAX(cl::sycl::cl_float4, cl::sycl::fmax(a, b))
0280 SYCL_PMAX(cl::sycl::cl_double2, cl::sycl::fmax(a, b))
0281 #undef SYCL_PMAX
0282
0283 #define SYCL_PLDEXP(packet_type)                                             \
0284   template <>                                                                \
0285   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pldexp(                  \
0286       const packet_type& a, const packet_type& exponent) {                   \
0287     return cl::sycl::ldexp(                                                  \
0288         a, exponent.template convert<cl::sycl::cl_int,                       \
0289                                      cl::sycl::rounding_mode::automatic>()); \
0290   }
0291
0292 SYCL_PLDEXP(cl::sycl::cl_float4)
0293 SYCL_PLDEXP(cl::sycl::cl_double2)
0294 #undef SYCL_PLDEXP
0295
0296 #endif
0297 }  // end namespace internal
0298
0299 }  // end namespace Eigen
0300
0301 #endif  // EIGEN_MATH_FUNCTIONS_SYCL_H