diff --git a/sycl/include/CL/sycl/builtins.hpp b/sycl/include/CL/sycl/builtins.hpp index b12bad848325a..055db680f8f43 100644 --- a/sycl/include/CL/sycl/builtins.hpp +++ b/sycl/include/CL/sycl/builtins.hpp @@ -32,523 +32,467 @@ namespace sycl { /* ----------------- 4.13.3 Math functions. ---------------------------------*/ // genfloat acos (genfloat x) template -typename std::enable_if::value, T>::type -acos(T x) __NOEXC { +detail::enable_if_t::value, T> acos(T x) __NOEXC { return __sycl_std::__invoke_acos(x); } // genfloat acosh (genfloat x) template -typename std::enable_if::value, T>::type -acosh(T x) __NOEXC { +detail::enable_if_t::value, T> acosh(T x) __NOEXC { return __sycl_std::__invoke_acosh(x); } // genfloat acospi (genfloat x) template -typename std::enable_if::value, T>::type -acospi(T x) __NOEXC { +detail::enable_if_t::value, T> acospi(T x) __NOEXC { return __sycl_std::__invoke_acospi(x); } // genfloat asin (genfloat x) template -typename std::enable_if::value, T>::type -asin(T x) __NOEXC { +detail::enable_if_t::value, T> asin(T x) __NOEXC { return __sycl_std::__invoke_asin(x); } // genfloat asinh (genfloat x) template -typename std::enable_if::value, T>::type -asinh(T x) __NOEXC { +detail::enable_if_t::value, T> asinh(T x) __NOEXC { return __sycl_std::__invoke_asinh(x); } // genfloat asinpi (genfloat x) template -typename std::enable_if::value, T>::type -asinpi(T x) __NOEXC { +detail::enable_if_t::value, T> asinpi(T x) __NOEXC { return __sycl_std::__invoke_asinpi(x); } // genfloat atan (genfloat y_over_x) template -typename std::enable_if::value, T>::type -atan(T y_over_x) __NOEXC { +detail::enable_if_t::value, T> atan(T y_over_x) __NOEXC { return __sycl_std::__invoke_atan(y_over_x); } // genfloat atan2 (genfloat y, genfloat x) template -typename std::enable_if::value, T>::type -atan2(T y, T x) __NOEXC { +detail::enable_if_t::value, T> atan2(T y, T x) __NOEXC { return __sycl_std::__invoke_atan2(y, x); } // genfloat atanh (genfloat x) template -typename std::enable_if::value, T>::type -atanh(T x) __NOEXC { +detail::enable_if_t::value, T> atanh(T x) __NOEXC { return __sycl_std::__invoke_atanh(x); } // genfloat atanpi (genfloat x) template -typename std::enable_if::value, T>::type -atanpi(T x) __NOEXC { +detail::enable_if_t::value, T> atanpi(T x) __NOEXC { return __sycl_std::__invoke_atanpi(x); } // genfloat atan2pi (genfloat y, genfloat x) template -typename std::enable_if::value, T>::type -atan2pi(T y, T x) __NOEXC { +detail::enable_if_t::value, T> atan2pi(T y, + T x) __NOEXC { return __sycl_std::__invoke_atan2pi(y, x); } // genfloat cbrt (genfloat x) template -typename std::enable_if::value, T>::type -cbrt(T x) __NOEXC { +detail::enable_if_t::value, T> cbrt(T x) __NOEXC { return __sycl_std::__invoke_cbrt(x); } // genfloat ceil (genfloat x) template -typename std::enable_if::value, T>::type -ceil(T x) __NOEXC { +detail::enable_if_t::value, T> ceil(T x) __NOEXC { return __sycl_std::__invoke_ceil(x); } // genfloat copysign (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -copysign(T x, T y) __NOEXC { +detail::enable_if_t::value, T> copysign(T x, + T y) __NOEXC { return __sycl_std::__invoke_copysign(x, y); } // genfloat cos (genfloat x) template -typename std::enable_if::value, T>::type -cos(T x) __NOEXC { +detail::enable_if_t::value, T> cos(T x) __NOEXC { return __sycl_std::__invoke_cos(x); } // genfloat cosh (genfloat x) template -typename std::enable_if::value, T>::type -cosh(T x) __NOEXC { +detail::enable_if_t::value, T> cosh(T x) __NOEXC { return __sycl_std::__invoke_cosh(x); } // genfloat cospi (genfloat x) template -typename std::enable_if::value, T>::type -cospi(T x) __NOEXC { +detail::enable_if_t::value, T> cospi(T x) __NOEXC { return __sycl_std::__invoke_cospi(x); } // genfloat erfc (genfloat x) template -typename std::enable_if::value, T>::type -erfc(T x) __NOEXC { +detail::enable_if_t::value, T> erfc(T x) __NOEXC { return __sycl_std::__invoke_erfc(x); } // genfloat erf (genfloat x) template -typename std::enable_if::value, T>::type -erf(T x) __NOEXC { +detail::enable_if_t::value, T> erf(T x) __NOEXC { return __sycl_std::__invoke_erf(x); } // genfloat exp (genfloat x ) template -typename std::enable_if::value, T>::type -exp(T x) __NOEXC { +detail::enable_if_t::value, T> exp(T x) __NOEXC { return __sycl_std::__invoke_exp(x); } // genfloat exp2 (genfloat x) template -typename std::enable_if::value, T>::type -exp2(T x) __NOEXC { +detail::enable_if_t::value, T> exp2(T x) __NOEXC { return __sycl_std::__invoke_exp2(x); } // genfloat exp10 (genfloat x) template -typename std::enable_if::value, T>::type -exp10(T x) __NOEXC { +detail::enable_if_t::value, T> exp10(T x) __NOEXC { return __sycl_std::__invoke_exp10(x); } // genfloat expm1 (genfloat x) template -typename std::enable_if::value, T>::type -expm1(T x) __NOEXC { +detail::enable_if_t::value, T> expm1(T x) __NOEXC { return __sycl_std::__invoke_expm1(x); } // genfloat fabs (genfloat x) template -typename std::enable_if::value, T>::type -fabs(T x) __NOEXC { +detail::enable_if_t::value, T> fabs(T x) __NOEXC { return __sycl_std::__invoke_fabs(x); } // genfloat fdim (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -fdim(T x, T y) __NOEXC { +detail::enable_if_t::value, T> fdim(T x, T y) __NOEXC { return __sycl_std::__invoke_fdim(x, y); } // genfloat floor (genfloat x) template -typename std::enable_if::value, T>::type -floor(T x) __NOEXC { +detail::enable_if_t::value, T> floor(T x) __NOEXC { return __sycl_std::__invoke_floor(x); } // genfloat fma (genfloat a, genfloat b, genfloat c) template -typename std::enable_if::value, T>::type -fma(T a, T b, T c) __NOEXC { +detail::enable_if_t::value, T> fma(T a, T b, + T c) __NOEXC { return __sycl_std::__invoke_fma(a, b, c); } // genfloat fmax (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -fmax(T x, T y) __NOEXC { +detail::enable_if_t::value, T> fmax(T x, T y) __NOEXC { return __sycl_std::__invoke_fmax(x, y); } // genfloat fmax (genfloat x, sgenfloat y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> fmax(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_fmax(x, T(y)); } // genfloat fmin (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -fmin(T x, T y) __NOEXC { +detail::enable_if_t::value, T> fmin(T x, T y) __NOEXC { return __sycl_std::__invoke_fmin(x, y); } // genfloat fmin (genfloat x, sgenfloat y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> fmin(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_fmin(x, T(y)); } // genfloat fmod (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -fmod(T x, T y) __NOEXC { +detail::enable_if_t::value, T> fmod(T x, T y) __NOEXC { return __sycl_std::__invoke_fmod(x, y); } // genfloat fract (genfloat x, genfloatptr iptr) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genfloatptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genfloatptr::value, T> fract(T x, T2 iptr) __NOEXC { return __sycl_std::__invoke_fract(x, iptr); } // genfloat frexp (genfloat x, genintptr exp) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genintptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genintptr::value, T> frexp(T x, T2 exp) __NOEXC { return __sycl_std::__invoke_frexp(x, exp); } // genfloat hypot (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -hypot(T x, T y) __NOEXC { +detail::enable_if_t::value, T> hypot(T x, T y) __NOEXC { return __sycl_std::__invoke_hypot(x, y); } // genint ilogb (genfloat x) -template ::value, T>::type> -typename detail::float_point_to_int::type ilogb(T x) __NOEXC { - return __sycl_std::__invoke_ilogb< - typename detail::float_point_to_int::type>(x); +template ::value, T>> +detail::change_base_type_t ilogb(T x) __NOEXC { + return __sycl_std::__invoke_ilogb>(x); } // float ldexp (float x, int k) // double ldexp (double x, int k) // half ldexp (half x, int k) template -typename std::enable_if::value, T>::type -ldexp(T x, int k) __NOEXC { +detail::enable_if_t::value, T> ldexp(T x, + int k) __NOEXC { return __sycl_std::__invoke_ldexp(x, k); } // vgenfloat ldexp (vgenfloat x, int k) template -typename std::enable_if::value, T>::type -ldexp(T x, int k) __NOEXC { - return __sycl_std::__invoke_ldexp( - x, cl::sycl::vec(k)); +detail::enable_if_t::value, T> ldexp(T x, + int k) __NOEXC { + return __sycl_std::__invoke_ldexp(x, vec(k)); } // vgenfloat ldexp (vgenfloat x, genint k) template -typename std::enable_if< - detail::is_vgenfloat::value && detail::is_intn::value, T>::type +detail::enable_if_t< + detail::is_vgenfloat::value && detail::is_intn::value, T> ldexp(T x, T2 k) __NOEXC { return __sycl_std::__invoke_ldexp(x, k); } // genfloat lgamma (genfloat x) template -typename std::enable_if::value, T>::type -lgamma(T x) __NOEXC { +detail::enable_if_t::value, T> lgamma(T x) __NOEXC { return __sycl_std::__invoke_lgamma(x); } // genfloat lgamma_r (genfloat x, genintptr signp) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genintptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genintptr::value, T> lgamma_r(T x, T2 signp) __NOEXC { return __sycl_std::__invoke_lgamma_r(x, signp); } // genfloat log (genfloat x) template -typename std::enable_if::value, T>::type -log(T x) __NOEXC { +detail::enable_if_t::value, T> log(T x) __NOEXC { return __sycl_std::__invoke_log(x); } // genfloat log2 (genfloat x) template -typename std::enable_if::value, T>::type -log2(T x) __NOEXC { +detail::enable_if_t::value, T> log2(T x) __NOEXC { return __sycl_std::__invoke_log2(x); } // genfloat log10 (genfloat x) template -typename std::enable_if::value, T>::type -log10(T x) __NOEXC { +detail::enable_if_t::value, T> log10(T x) __NOEXC { return __sycl_std::__invoke_log10(x); } // genfloat log1p (genfloat x) template -typename std::enable_if::value, T>::type -log1p(T x) __NOEXC { +detail::enable_if_t::value, T> log1p(T x) __NOEXC { return __sycl_std::__invoke_log1p(x); } // genfloat logb (genfloat x) template -typename std::enable_if::value, T>::type -logb(T x) __NOEXC { +detail::enable_if_t::value, T> logb(T x) __NOEXC { return __sycl_std::__invoke_logb(x); } // genfloat mad (genfloat a, genfloat b, genfloat c) template -typename std::enable_if::value, T>::type -mad(T a, T b, T c) __NOEXC { +detail::enable_if_t::value, T> mad(T a, T b, + T c) __NOEXC { return __sycl_std::__invoke_mad(a, b, c); } // genfloat maxmag (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -maxmag(T x, T y) __NOEXC { +detail::enable_if_t::value, T> maxmag(T x, T y) __NOEXC { return __sycl_std::__invoke_maxmag(x, y); } // genfloat minmag (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -minmag(T x, T y) __NOEXC { +detail::enable_if_t::value, T> minmag(T x, T y) __NOEXC { return __sycl_std::__invoke_minmag(x, y); } // genfloat modf (genfloat x, genfloatptr iptr) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genfloatptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genfloatptr::value, T> modf(T x, T2 iptr) __NOEXC { return __sycl_std::__invoke_modf(x, iptr); } -// genfloath nan (ugenshort nancode) -// genfloatf nan (ugenint nancode) -// genfloatd nan (ugenlonginteger nancode) -template ::value, T>::type> -detail::unsign_integral_to_float_point_t -nan(T nancode) __NOEXC { - return __sycl_std::__invoke_nan< - detail::unsign_integral_to_float_point_t>(nancode); +template ::value, T>> +detail::nan_return_t nan(T nancode) __NOEXC { + return __sycl_std::__invoke_nan>( + detail::convert_data_type>()(nancode)); } // genfloat nextafter (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -nextafter(T x, T y) __NOEXC { +detail::enable_if_t::value, T> nextafter(T x, + T y) __NOEXC { return __sycl_std::__invoke_nextafter(x, y); } // genfloat pow (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -pow(T x, T y) __NOEXC { +detail::enable_if_t::value, T> pow(T x, T y) __NOEXC { return __sycl_std::__invoke_pow(x, y); } // genfloat pown (genfloat x, genint y) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genint::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genint::value, T> pown(T x, T2 y) __NOEXC { return __sycl_std::__invoke_pown(x, y); } // genfloat powr (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -powr(T x, T y) __NOEXC { +detail::enable_if_t::value, T> powr(T x, T y) __NOEXC { return __sycl_std::__invoke_powr(x, y); } // genfloat remainder (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -remainder(T x, T y) __NOEXC { +detail::enable_if_t::value, T> remainder(T x, + T y) __NOEXC { return __sycl_std::__invoke_remainder(x, y); } // genfloat remquo (genfloat x, genfloat y, genintptr quo) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genintptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genintptr::value, T> remquo(T x, T y, T2 quo) __NOEXC { return __sycl_std::__invoke_remquo(x, y, quo); } // genfloat rint (genfloat x) template -typename std::enable_if::value, T>::type -rint(T x) __NOEXC { +detail::enable_if_t::value, T> rint(T x) __NOEXC { return __sycl_std::__invoke_rint(x); } // genfloat rootn (genfloat x, genint y) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genint::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genint::value, T> rootn(T x, T2 y) __NOEXC { return __sycl_std::__invoke_rootn(x, y); } // genfloat round (genfloat x) template -typename std::enable_if::value, T>::type -round(T x) __NOEXC { +detail::enable_if_t::value, T> round(T x) __NOEXC { return __sycl_std::__invoke_round(x); } // genfloat rsqrt (genfloat x) template -typename std::enable_if::value, T>::type -rsqrt(T x) __NOEXC { +detail::enable_if_t::value, T> rsqrt(T x) __NOEXC { return __sycl_std::__invoke_rsqrt(x); } // genfloat sin (genfloat x) template -typename std::enable_if::value, T>::type -sin(T x) __NOEXC { +detail::enable_if_t::value, T> sin(T x) __NOEXC { return __sycl_std::__invoke_sin(x); } // genfloat sincos (genfloat x, genfloatptr cosval) template -typename std::enable_if< - detail::is_genfloat::value && detail::is_genfloatptr::value, T>::type +detail::enable_if_t< + detail::is_genfloat::value && detail::is_genfloatptr::value, T> sincos(T x, T2 cosval) __NOEXC { return __sycl_std::__invoke_sincos(x, cosval); } // genfloat sinh (genfloat x) template -typename std::enable_if::value, T>::type -sinh(T x) __NOEXC { +detail::enable_if_t::value, T> sinh(T x) __NOEXC { return __sycl_std::__invoke_sinh(x); } // genfloat sinpi (genfloat x) template -typename std::enable_if::value, T>::type -sinpi(T x) __NOEXC { +detail::enable_if_t::value, T> sinpi(T x) __NOEXC { return __sycl_std::__invoke_sinpi(x); } // genfloat sqrt (genfloat x) template -typename std::enable_if::value, T>::type -sqrt(T x) __NOEXC { +detail::enable_if_t::value, T> sqrt(T x) __NOEXC { return __sycl_std::__invoke_sqrt(x); } // genfloat tan (genfloat x) template -typename std::enable_if::value, T>::type -tan(T x) __NOEXC { +detail::enable_if_t::value, T> tan(T x) __NOEXC { return __sycl_std::__invoke_tan(x); } // genfloat tanh (genfloat x) template -typename std::enable_if::value, T>::type -tanh(T x) __NOEXC { +detail::enable_if_t::value, T> tanh(T x) __NOEXC { return __sycl_std::__invoke_tanh(x); } // genfloat tanpi (genfloat x) template -typename std::enable_if::value, T>::type -tanpi(T x) __NOEXC { +detail::enable_if_t::value, T> tanpi(T x) __NOEXC { return __sycl_std::__invoke_tanpi(x); } // genfloat tgamma (genfloat x) template -typename std::enable_if::value, T>::type -tgamma(T x) __NOEXC { +detail::enable_if_t::value, T> tgamma(T x) __NOEXC { return __sycl_std::__invoke_tgamma(x); } // genfloat trunc (genfloat x) template -typename std::enable_if::value, T>::type -trunc(T x) __NOEXC { +detail::enable_if_t::value, T> trunc(T x) __NOEXC { return __sycl_std::__invoke_trunc(x); } /* --------------- 4.13.5 Common functions. ---------------------------------*/ // genfloat clamp (genfloat x, genfloat minval, genfloat maxval) template -typename std::enable_if::value, T>::type -clamp(T x, T minval, T maxval) __NOEXC { +detail::enable_if_t::value, T> clamp(T x, T minval, + T maxval) __NOEXC { return __sycl_std::__invoke_fclamp(x, minval, maxval); } @@ -556,7 +500,7 @@ clamp(T x, T minval, T maxval) __NOEXC { // genfloatf clamp (genfloatf x, float minval, float maxval) // genfloatd clamp (genfloatd x, double minval, double maxval) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> clamp(T x, typename T::element_type minval, typename T::element_type maxval) __NOEXC { return __sycl_std::__invoke_fclamp(x, T(minval), T(maxval)); @@ -564,22 +508,20 @@ clamp(T x, typename T::element_type minval, // genfloat degrees (genfloat radians) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> degrees(T radians) __NOEXC { return __sycl_std::__invoke_degrees(radians); } // genfloat abs (genfloat x) template -typename std::enable_if::value, T>::type -abs(T x) __NOEXC { +detail::enable_if_t::value, T> abs(T x) __NOEXC { return __sycl_std::__invoke_fabs(x); } // genfloat max (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -max(T x, T y) __NOEXC { +detail::enable_if_t::value, T> max(T x, T y) __NOEXC { return __sycl_std::__invoke_fmax_common(x, y); } @@ -587,15 +529,14 @@ max(T x, T y) __NOEXC { // genfloatd max (genfloatd x, double y) // genfloath max (genfloath x, half y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> max(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_fmax_common(x, T(y)); } // genfloat min (genfloat x, genfloat y) template -typename std::enable_if::value, T>::type -min(T x, T y) __NOEXC { +detail::enable_if_t::value, T> min(T x, T y) __NOEXC { return __sycl_std::__invoke_fmin_common(x, y); } @@ -603,15 +544,15 @@ min(T x, T y) __NOEXC { // genfloatd min (genfloatd x, double y) // genfloath min (genfloath x, half y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> min(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_fmin_common(x, T(y)); } // genfloat mix (genfloat x, genfloat y, genfloat a) template -typename std::enable_if::value, T>::type -mix(T x, T y, T a) __NOEXC { +detail::enable_if_t::value, T> mix(T x, T y, + T a) __NOEXC { return __sycl_std::__invoke_mix(x, y, a); } @@ -619,22 +560,22 @@ mix(T x, T y, T a) __NOEXC { // genfloatd mix (genfloatd x, genfloatd y, double a) // genfloatd mix (genfloath x, genfloath y, half a) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> mix(T x, T y, typename T::element_type a) __NOEXC { return __sycl_std::__invoke_mix(x, y, T(a)); } // genfloat radians (genfloat degrees) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> radians(T degrees) __NOEXC { return __sycl_std::__invoke_radians(degrees); } // genfloat step (genfloat edge, genfloat x) template -typename std::enable_if::value, T>::type -step(T edge, T x) __NOEXC { +detail::enable_if_t::value, T> step(T edge, + T x) __NOEXC { return __sycl_std::__invoke_step(edge, x); } @@ -642,14 +583,14 @@ step(T edge, T x) __NOEXC { // genfloatd step (double edge, genfloatd x) // genfloatd step (half edge, genfloath x) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> step(typename T::element_type edge, T x) __NOEXC { return __sycl_std::__invoke_step(T(edge), x); } // genfloat smoothstep (genfloat edge0, genfloat edge1, genfloat x) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> smoothstep(T edge0, T edge1, T x) __NOEXC { return __sycl_std::__invoke_smoothstep(edge0, edge1, x); } @@ -658,7 +599,7 @@ smoothstep(T edge0, T edge1, T x) __NOEXC { // genfloatd smoothstep (double edge0, double edge1, genfloatd x) // genfloath smoothstep (half edge0, half edge1, genfloath x) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> smoothstep(typename T::element_type edge0, typename T::element_type edge1, T x) __NOEXC { return __sycl_std::__invoke_smoothstep(T(edge0), T(edge1), x); @@ -666,102 +607,99 @@ smoothstep(typename T::element_type edge0, typename T::element_type edge1, // genfloat sign (genfloat x) template -typename std::enable_if::value, T>::type -sign(T x) __NOEXC { +detail::enable_if_t::value, T> sign(T x) __NOEXC { return __sycl_std::__invoke_sign(x); } /* --------------- 4.13.4 Integer functions. --------------------------------*/ // ugeninteger abs (geninteger x) template -typename std::enable_if::value, T>::type -abs(T x) __NOEXC { +detail::enable_if_t::value, T> abs(T x) __NOEXC { return __sycl_std::__invoke_u_abs(x); } // ugeninteger abs (geninteger x) template -typename std::enable_if::value, - typename detail::make_unsigned::type>::type +detail::enable_if_t::value, + detail::make_unsigned_t> abs(T x) __NOEXC { - return __sycl_std::__invoke_s_abs::type>(x); + return __sycl_std::__invoke_s_abs>(x); } // ugeninteger abs_diff (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -abs_diff(T x, T y) __NOEXC { +detail::enable_if_t::value, T> abs_diff(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_abs_diff(x, y); } // ugeninteger abs_diff (geninteger x, geninteger y) template -typename std::enable_if::value, - typename detail::make_unsigned::type>::type +detail::enable_if_t::value, + detail::make_unsigned_t> abs_diff(T x, T y) __NOEXC { - return __sycl_std::__invoke_s_abs_diff< - typename detail::make_unsigned::type>(x, y); + return __sycl_std::__invoke_s_abs_diff>(x, y); } // geninteger add_sat (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -add_sat(T x, T y) __NOEXC { +detail::enable_if_t::value, T> add_sat(T x, + T y) __NOEXC { return __sycl_std::__invoke_s_add_sat(x, y); } // geninteger add_sat (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -add_sat(T x, T y) __NOEXC { +detail::enable_if_t::value, T> add_sat(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_add_sat(x, y); } // geninteger hadd (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -hadd(T x, T y) __NOEXC { +detail::enable_if_t::value, T> hadd(T x, + T y) __NOEXC { return __sycl_std::__invoke_s_hadd(x, y); } // geninteger hadd (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -hadd(T x, T y) __NOEXC { +detail::enable_if_t::value, T> hadd(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_hadd(x, y); } // geninteger rhadd (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -rhadd(T x, T y) __NOEXC { +detail::enable_if_t::value, T> rhadd(T x, + T y) __NOEXC { return __sycl_std::__invoke_s_rhadd(x, y); } // geninteger rhadd (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -rhadd(T x, T y) __NOEXC { +detail::enable_if_t::value, T> rhadd(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_rhadd(x, y); } // geninteger clamp (geninteger x, geninteger minval, geninteger maxval) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> clamp(T x, T minval, T maxval) __NOEXC { return __sycl_std::__invoke_s_clamp(x, minval, maxval); } // geninteger clamp (geninteger x, geninteger minval, geninteger maxval) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> clamp(T x, T minval, T maxval) __NOEXC { return __sycl_std::__invoke_u_clamp(x, minval, maxval); } // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> clamp(T x, typename T::element_type minval, typename T::element_type maxval) __NOEXC { return __sycl_std::__invoke_s_clamp(x, T(minval), T(maxval)); @@ -769,7 +707,7 @@ clamp(T x, typename T::element_type minval, // geninteger clamp (geninteger x, sgeninteger minval, sgeninteger maxval) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> clamp(T x, typename T::element_type minval, typename T::element_type maxval) __NOEXC { return __sycl_std::__invoke_u_clamp(x, T(minval), T(maxval)); @@ -777,136 +715,130 @@ clamp(T x, typename T::element_type minval, // geninteger clz (geninteger x) template -typename std::enable_if::value, T>::type -clz(T x) __NOEXC { +detail::enable_if_t::value, T> clz(T x) __NOEXC { return __sycl_std::__invoke_clz(x); } namespace intel { // geninteger ctz (geninteger x) template -typename std::enable_if::value, T>::type -ctz(T x) __NOEXC { +detail::enable_if_t::value, T> ctz(T x) __NOEXC { return __sycl_std::__invoke_ctz(x); } -} +} // namespace intel // geninteger mad_hi (geninteger a, geninteger b, geninteger c) template -typename std::enable_if::value, T>::type -mad_hi(T x, T y, T z) __NOEXC { +detail::enable_if_t::value, T> mad_hi(T x, T y, + T z) __NOEXC { return __sycl_std::__invoke_s_mad_hi(x, y, z); } // geninteger mad_hi (geninteger a, geninteger b, geninteger c) template -typename std::enable_if::value, T>::type -mad_hi(T x, T y, T z) __NOEXC { +detail::enable_if_t::value, T> mad_hi(T x, T y, + T z) __NOEXC { return __sycl_std::__invoke_u_mad_hi(x, y, z); } // geninteger mad_sat (geninteger a, geninteger b, geninteger c) template -typename std::enable_if::value, T>::type -mad_sat(T a, T b, T c) __NOEXC { +detail::enable_if_t::value, T> mad_sat(T a, T b, + T c) __NOEXC { return __sycl_std::__invoke_s_mad_sat(a, b, c); } // geninteger mad_sat (geninteger a, geninteger b, geninteger c) template -typename std::enable_if::value, T>::type -mad_sat(T a, T b, T c) __NOEXC { +detail::enable_if_t::value, T> mad_sat(T a, T b, + T c) __NOEXC { return __sycl_std::__invoke_u_mad_sat(a, b, c); } // igeninteger max (igeninteger x, igeninteger y) template -typename std::enable_if::value, T>::type -max(T x, T y) __NOEXC { +detail::enable_if_t::value, T> max(T x, T y) __NOEXC { return __sycl_std::__invoke_s_max(x, y); } // ugeninteger max (ugeninteger x, ugeninteger y) template -typename std::enable_if::value, T>::type -max(T x, T y) __NOEXC { +detail::enable_if_t::value, T> max(T x, T y) __NOEXC { return __sycl_std::__invoke_u_max(x, y); } // igeninteger max (vigeninteger x, sigeninteger y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> max(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_s_max(x, T(y)); } // vugeninteger max (vugeninteger x, sugeninteger y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> max(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_u_max(x, T(y)); } // igeninteger min (igeninteger x, igeninteger y) template -typename std::enable_if::value, T>::type -min(T x, T y) __NOEXC { +detail::enable_if_t::value, T> min(T x, T y) __NOEXC { return __sycl_std::__invoke_s_min(x, y); } // ugeninteger min (ugeninteger x, ugeninteger y) template -typename std::enable_if::value, T>::type -min(T x, T y) __NOEXC { +detail::enable_if_t::value, T> min(T x, T y) __NOEXC { return __sycl_std::__invoke_u_min(x, y); } // vigeninteger min (vigeninteger x, sigeninteger y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> min(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_s_min(x, T(y)); } // vugeninteger min (vugeninteger x, sugeninteger y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> min(T x, typename T::element_type y) __NOEXC { return __sycl_std::__invoke_u_min(x, T(y)); } // geninteger mul_hi (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -mul_hi(T x, T y) __NOEXC { +detail::enable_if_t::value, T> mul_hi(T x, + T y) __NOEXC { return __sycl_std::__invoke_s_mul_hi(x, y); } // geninteger mul_hi (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -mul_hi(T x, T y) __NOEXC { +detail::enable_if_t::value, T> mul_hi(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_mul_hi(x, y); } // geninteger rotate (geninteger v, geninteger i) template -typename std::enable_if::value, T>::type -rotate(T v, T i) __NOEXC { +detail::enable_if_t::value, T> rotate(T v, + T i) __NOEXC { return __sycl_std::__invoke_rotate(v, i); } // geninteger sub_sat (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -sub_sat(T x, T y) __NOEXC { +detail::enable_if_t::value, T> sub_sat(T x, + T y) __NOEXC { return __sycl_std::__invoke_s_sub_sat(x, y); } // geninteger sub_sat (geninteger x, geninteger y) template -typename std::enable_if::value, T>::type -sub_sat(T x, T y) __NOEXC { +detail::enable_if_t::value, T> sub_sat(T x, + T y) __NOEXC { return __sycl_std::__invoke_u_sub_sat(x, y); } @@ -916,74 +848,67 @@ sub_sat(T x, T y) __NOEXC { // ugeninteger16bit upsample (ugeninteger8bit hi, ugeninteger8bit lo) template -typename std::enable_if::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value, + detail::make_larger_t> upsample(T hi, T lo) __NOEXC { - return __sycl_std::__invoke_u_upsample::type>( - hi, lo); + return __sycl_std::__invoke_u_upsample>(hi, lo); } // igeninteger16bit upsample (igeninteger8bit hi, ugeninteger8bit lo) template -typename std::enable_if::value && - detail::is_ugeninteger8bit::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value && + detail::is_ugeninteger8bit::value, + detail::make_larger_t> upsample(T hi, T2 lo) __NOEXC { - return __sycl_std::__invoke_s_upsample::type>( - hi, lo); + return __sycl_std::__invoke_s_upsample>(hi, lo); } // ugeninteger32bit upsample (ugeninteger16bit hi, ugeninteger16bit lo) template -typename std::enable_if::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value, + detail::make_larger_t> upsample(T hi, T lo) __NOEXC { - return __sycl_std::__invoke_u_upsample::type>( - hi, lo); + return __sycl_std::__invoke_u_upsample>(hi, lo); } // igeninteger32bit upsample (igeninteger16bit hi, ugeninteger16bit lo) template -typename std::enable_if::value && - detail::is_ugeninteger16bit::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value && + detail::is_ugeninteger16bit::value, + detail::make_larger_t> upsample(T hi, T2 lo) __NOEXC { - return __sycl_std::__invoke_s_upsample::type>( - hi, lo); + return __sycl_std::__invoke_s_upsample>(hi, lo); } // ugeninteger64bit upsample (ugeninteger32bit hi, ugeninteger32bit lo) template -typename std::enable_if::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value, + detail::make_larger_t> upsample(T hi, T lo) __NOEXC { - return __sycl_std::__invoke_u_upsample::type>( - hi, lo); + return __sycl_std::__invoke_u_upsample>(hi, lo); } // igeninteger64bit upsample (igeninteger32bit hi, ugeninteger32bit lo) template -typename std::enable_if::value && - detail::is_ugeninteger32bit::value, - typename detail::make_upper::type>::type +detail::enable_if_t::value && + detail::is_ugeninteger32bit::value, + detail::make_larger_t> upsample(T hi, T2 lo) __NOEXC { - return __sycl_std::__invoke_s_upsample::type>( - hi, lo); + return __sycl_std::__invoke_s_upsample>(hi, lo); } #undef __invoke_s_upsample // geninteger popcount (geninteger x) template -typename std::enable_if::value, T>::type -popcount(T x) __NOEXC { +detail::enable_if_t::value, T> popcount(T x) __NOEXC { return __sycl_std::__invoke_popcount(x); } // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y, // geninteger32bit z) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> mad24(T x, T y, T z) __NOEXC { return __sycl_std::__invoke_s_mad24(x, y, z); } @@ -991,21 +916,21 @@ mad24(T x, T y, T z) __NOEXC { // geninteger32bit mad24 (geninteger32bit x, geninteger32bit y, // geninteger32bit z) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> mad24(T x, T y, T z) __NOEXC { return __sycl_std::__invoke_u_mad24(x, y, z); } // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> mul24(T x, T y) __NOEXC { return __sycl_std::__invoke_s_mul24(x, y); } // geninteger32bit mul24 (geninteger32bit x, geninteger32bit y) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> mul24(T x, T y) __NOEXC { return __sycl_std::__invoke_u_mul24(x, y); } @@ -1018,8 +943,8 @@ mul24(T x, T y) __NOEXC { // half3 cross (half3 p0, half3 p1) // half4 cross (half4 p0, half4 p1) template -typename std::enable_if::value, T>::type -cross(T p0, T p1) __NOEXC { +detail::enable_if_t::value, T> cross(T p0, + T p1) __NOEXC { return __sycl_std::__invoke_cross(p0, p1); } @@ -1027,136 +952,131 @@ cross(T p0, T p1) __NOEXC { // double dot (double p0, double p1) // half dot (half p0, half p1) template -typename std::enable_if::value, T>::type -dot(T p0, T p1) __NOEXC { +detail::enable_if_t::value, T> dot(T p0, T p1) __NOEXC { return __sycl_std::__invoke_FMul(p0, p1); } // float dot (vgengeofloat p0, vgengeofloat p1) template -typename std::enable_if::value, - cl::sycl::cl_float>::type +detail::enable_if_t::value, float> dot(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_Dot(p0, p1); + return __sycl_std::__invoke_Dot(p0, p1); } // double dot (vgengeodouble p0, vgengeodouble p1) template -typename std::enable_if::value, - cl::sycl::cl_double>::type +detail::enable_if_t::value, double> dot(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_Dot(p0, p1); + return __sycl_std::__invoke_Dot(p0, p1); } // half dot (vgengeohalf p0, vgengeohalf p1) template -typename std::enable_if::value, - cl::sycl::cl_half>::type -dot(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_Dot(p0, p1); +detail::enable_if_t::value, half> dot(T p0, + T p1) __NOEXC { + return __sycl_std::__invoke_Dot(p0, p1); } // float distance (gengeofloat p0, gengeofloat p1) -template ::value, T>::type> -cl::sycl::cl_float distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_distance(p0, p1); +template ::value, T>> +float distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); } // double distance (gengeodouble p0, gengeodouble p1) -template ::value, T>::type> -cl::sycl::cl_double distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_distance(p0, p1); +template ::value, T>> +double distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); } // half distance (gengeohalf p0, gengeohalf p1) -template ::value, T>::type> -cl::sycl::cl_half distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_distance(p0, p1); +template ::value, T>> +half distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_distance(p0, p1); } // float length (gengeofloat p) -template ::value, T>::type> -cl::sycl::cl_float length(T p) __NOEXC { - return __sycl_std::__invoke_length(p); +template ::value, T>> +float length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); } // double length (gengeodouble p) -template ::value, T>::type> -cl::sycl::cl_double length(T p) __NOEXC { - return __sycl_std::__invoke_length(p); +template ::value, T>> +double length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); } // half length (gengeohalf p) -template ::value, T>::type> -cl::sycl::cl_half length(T p) __NOEXC { - return __sycl_std::__invoke_length(p); +template ::value, T>> +half length(T p) __NOEXC { + return __sycl_std::__invoke_length(p); } // gengeofloat normalize (gengeofloat p) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> normalize(T p) __NOEXC { return __sycl_std::__invoke_normalize(p); } // gengeodouble normalize (gengeodouble p) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> normalize(T p) __NOEXC { return __sycl_std::__invoke_normalize(p); } // gengeohalf normalize (gengeohalf p) template -typename std::enable_if::value, T>::type -normalize(T p) __NOEXC { +detail::enable_if_t::value, T> normalize(T p) __NOEXC { return __sycl_std::__invoke_normalize(p); } // float fast_distance (gengeofloat p0, gengeofloat p1) -template ::value, T>::type> -cl::sycl::cl_float fast_distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_fast_distance(p0, p1); +template ::value, T>> +float fast_distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_fast_distance(p0, p1); } // double fast_distance (gengeodouble p0, gengeodouble p1) -template ::value, T>::type> -cl::sycl::cl_double fast_distance(T p0, T p1) __NOEXC { - return __sycl_std::__invoke_fast_distance(p0, p1); +template ::value, T>> +double fast_distance(T p0, T p1) __NOEXC { + return __sycl_std::__invoke_fast_distance(p0, p1); } // float fast_length (gengeofloat p) -template ::value, T>::type> -cl::sycl::cl_float fast_length(T p) __NOEXC { - return __sycl_std::__invoke_fast_length(p); +template ::value, T>> +float fast_length(T p) __NOEXC { + return __sycl_std::__invoke_fast_length(p); } // double fast_length (gengeodouble p) -template ::value, T>::type> -cl::sycl::cl_double fast_length(T p) __NOEXC { - return __sycl_std::__invoke_fast_length(p); +template ::value, T>> +double fast_length(T p) __NOEXC { + return __sycl_std::__invoke_fast_length(p); } // gengeofloat fast_normalize (gengeofloat p) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } // gengeodouble fast_normalize (gengeodouble p) template -typename std::enable_if::value, T>::type +detail::enable_if_t::value, T> fast_normalize(T p) __NOEXC { return __sycl_std::__invoke_fast_normalize(p); } @@ -1167,8 +1087,8 @@ fast_normalize(T p) __NOEXC { // igeninteger32bit isequal (genfloatf x, genfloatf y) // int isequal (double x,double y); // longn isequal (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isequal(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FOrdEqual>(x, y)); @@ -1179,8 +1099,8 @@ detail::common_rel_ret_t isequal(T x, T y) __NOEXC { // igeninteger32bit isnotequal (genfloatf x, genfloatf y) // int isnotequal (double x, double y) // longn isnotequal (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isnotequal(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FUnordNotEqual>(x, y)); @@ -1191,8 +1111,8 @@ detail::common_rel_ret_t isnotequal(T x, T y) __NOEXC { // igeninteger32bit isgreater (genfloatf x, genfloatf y) // int isgreater (double x, double y) // longn isgreater (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isgreater(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FOrdGreaterThan>(x, y)); @@ -1203,8 +1123,8 @@ detail::common_rel_ret_t isgreater(T x, T y) __NOEXC { // igeninteger32bit isgreaterequal (genfloatf x, genfloatf y) // int isgreaterequal (double x, double y) // longn isgreaterequal (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isgreaterequal(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FOrdGreaterThanEqual>(x, y)); @@ -1215,8 +1135,8 @@ detail::common_rel_ret_t isgreaterequal(T x, T y) __NOEXC { // igeninteger32bit isless (genfloatf x, genfloatf y) // int isless (long x, long y) // longn isless (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isless(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FOrdLessThan>(x, y)); @@ -1227,8 +1147,8 @@ detail::common_rel_ret_t isless(T x, T y) __NOEXC { // igeninteger32bit islessequal (genfloatf x, genfloatf y) // int islessequal (double x, double y) // longn islessequal (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t islessequal(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_FOrdLessThanEqual>(x, y)); @@ -1239,8 +1159,8 @@ detail::common_rel_ret_t islessequal(T x, T y) __NOEXC { // igeninteger32bit islessgreater (genfloatf x, genfloatf y) // int islessgreater (double x, double y) // longn islessgreater (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_LessOrGreater>(x, y)); @@ -1251,8 +1171,8 @@ detail::common_rel_ret_t islessgreater(T x, T y) __NOEXC { // igeninteger32bit isfinite (genfloatf x) // int isfinite (double x) // longn isfinite (doublen x) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isfinite(T x) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_IsFinite>(x)); @@ -1263,8 +1183,8 @@ detail::common_rel_ret_t isfinite(T x) __NOEXC { // igeninteger32bit isinf (genfloatf x) // int isinf (double x) // longn isinf (doublen x) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isinf(T x) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_IsInf>(x)); @@ -1275,8 +1195,8 @@ detail::common_rel_ret_t isinf(T x) __NOEXC { // igeninteger32bit isnan (genfloatf x) // int isnan (double x) // longn isnan (doublen x) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isnan(T x) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_IsNan>(x)); @@ -1287,8 +1207,8 @@ detail::common_rel_ret_t isnan(T x) __NOEXC { // igeninteger32bit isnormal (genfloatf x) // int isnormal (double x) // longn isnormal (doublen x) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isnormal(T x) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_IsNormal>(x)); @@ -1299,8 +1219,8 @@ detail::common_rel_ret_t isnormal(T x) __NOEXC { // igeninteger32bit isordered (genfloatf x, genfloatf y) // int isordered (double x, double y) // longn isordered (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isordered(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_Ordered>(x, y)); @@ -1311,8 +1231,8 @@ detail::common_rel_ret_t isordered(T x, T y) __NOEXC { // igeninteger32bit isunordered (genfloatf x, genfloatf y) // int isunordered (double x, double y) // longn isunordered (doublen x, doublen y) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t isunordered(T x, T y) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_Unordered>(x, y)); @@ -1323,8 +1243,8 @@ detail::common_rel_ret_t isunordered(T x, T y) __NOEXC { // igeninteger32bit signbit (genfloatf x) // int signbit (double) // longn signbit (doublen x) -template ::value, T>::type> +template ::value, T>> detail::common_rel_ret_t signbit(T x) __NOEXC { return detail::RelConverter::apply( __sycl_std::__invoke_SignBitSet>(x)); @@ -1332,17 +1252,13 @@ detail::common_rel_ret_t signbit(T x) __NOEXC { // int any (sigeninteger x) template -typename std::enable_if::value, - cl::sycl::cl_int>::type -any(T x) __NOEXC { - return detail::Boolean<1>(cl::sycl::cl_int(detail::msbIsSet(x))); +detail::enable_if_t::value, int> any(T x) __NOEXC { + return detail::Boolean<1>(int(detail::msbIsSet(x))); } // int any (vigeninteger x) template -typename std::enable_if::value, - cl::sycl::cl_int>::type -any(T x) __NOEXC { +detail::enable_if_t::value, int> any(T x) __NOEXC { return detail::rel_sign_bit_test_ret_t( __sycl_std::__invoke_Any>( detail::rel_sign_bit_test_arg_t(x))); @@ -1350,17 +1266,13 @@ any(T x) __NOEXC { // int all (sigeninteger x) template -typename std::enable_if::value, - cl::sycl::cl_int>::type -all(T x) __NOEXC { - return detail::Boolean<1>(cl::sycl::cl_int(detail::msbIsSet(x))); +detail::enable_if_t::value, int> all(T x) __NOEXC { + return detail::Boolean<1>(int(detail::msbIsSet(x))); } // int all (vigeninteger x) template -typename std::enable_if::value, - cl::sycl::cl_int>::type -all(T x) __NOEXC { +detail::enable_if_t::value, int> all(T x) __NOEXC { return detail::rel_sign_bit_test_ret_t( __sycl_std::__invoke_All>( detail::rel_sign_bit_test_arg_t(x))); @@ -1368,77 +1280,71 @@ all(T x) __NOEXC { // gentype bitselect (gentype a, gentype b, gentype c) template -typename std::enable_if::value, T>::type -bitselect(T a, T b, T c) __NOEXC { +detail::enable_if_t::value, T> bitselect(T a, T b, + T c) __NOEXC { return __sycl_std::__invoke_bitselect(a, b, c); } // geninteger select (geninteger a, geninteger b, igeninteger c) template -typename std::enable_if::value && - detail::is_igeninteger::value, - T>::type +detail::enable_if_t< + detail::is_geninteger::value && detail::is_igeninteger::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // geninteger select (geninteger a, geninteger b, ugeninteger c) template -typename std::enable_if::value && - detail::is_ugeninteger::value, - T>::type +detail::enable_if_t< + detail::is_geninteger::value && detail::is_ugeninteger::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloatf select (genfloatf a, genfloatf b, genint c) template -typename std::enable_if< - detail::is_genfloatf::value && detail::is_genint::value, T>::type +detail::enable_if_t< + detail::is_genfloatf::value && detail::is_genint::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloatf select (genfloatf a, genfloatf b, ugenint c) template -typename std::enable_if< - detail::is_genfloatf::value && detail::is_ugenint::value, T>::type +detail::enable_if_t< + detail::is_genfloatf::value && detail::is_ugenint::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloatd select (genfloatd a, genfloatd b, igeninteger64 c) template -typename std::enable_if::value && - detail::is_igeninteger64bit::value, - T>::type +detail::enable_if_t< + detail::is_genfloatd::value && detail::is_igeninteger64bit::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloatd select (genfloatd a, genfloatd b, ugeninteger64 c) template -typename std::enable_if::value && - detail::is_ugeninteger64bit::value, - T>::type +detail::enable_if_t< + detail::is_genfloatd::value && detail::is_ugeninteger64bit::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloath select (genfloath a, genfloath b, igeninteger16 c) template -typename std::enable_if::value && - detail::is_igeninteger16bit::value, - T>::type +detail::enable_if_t< + detail::is_genfloath::value && detail::is_igeninteger16bit::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } // genfloath select (genfloath a, genfloath b, ugeninteger16 c) template -typename std::enable_if::value && - detail::is_ugeninteger16bit::value, - T>::type +detail::enable_if_t< + detail::is_genfloath::value && detail::is_ugeninteger16bit::value, T> select(T a, T b, T2 c) __NOEXC { return __sycl_std::__invoke_Select(detail::select_arg_c_t(c), b, a); } @@ -1447,99 +1353,86 @@ namespace native { /* ----------------- 4.13.3 Math functions. ---------------------------------*/ // genfloatf cos (genfloatf x) template -typename std::enable_if::value, T>::type -cos(T x) __NOEXC { +detail::enable_if_t::value, T> cos(T x) __NOEXC { return __sycl_std::__invoke_native_cos(x); } // genfloatf divide (genfloatf x, genfloatf y) template -typename std::enable_if::value, T>::type -divide(T x, T y) __NOEXC { +detail::enable_if_t::value, T> divide(T x, + T y) __NOEXC { return __sycl_std::__invoke_native_divide(x, y); } // genfloatf exp (genfloatf x) template -typename std::enable_if::value, T>::type -exp(T x) __NOEXC { +detail::enable_if_t::value, T> exp(T x) __NOEXC { return __sycl_std::__invoke_native_exp(x); } // genfloatf exp2 (genfloatf x) template -typename std::enable_if::value, T>::type -exp2(T x) __NOEXC { +detail::enable_if_t::value, T> exp2(T x) __NOEXC { return __sycl_std::__invoke_native_exp2(x); } // genfloatf exp10 (genfloatf x) template -typename std::enable_if::value, T>::type -exp10(T x) __NOEXC { +detail::enable_if_t::value, T> exp10(T x) __NOEXC { return __sycl_std::__invoke_native_exp10(x); } // genfloatf log (genfloatf x) template -typename std::enable_if::value, T>::type -log(T x) __NOEXC { +detail::enable_if_t::value, T> log(T x) __NOEXC { return __sycl_std::__invoke_native_log(x); } // genfloatf log2 (genfloatf x) template -typename std::enable_if::value, T>::type -log2(T x) __NOEXC { +detail::enable_if_t::value, T> log2(T x) __NOEXC { return __sycl_std::__invoke_native_log2(x); } // genfloatf log10 (genfloatf x) template -typename std::enable_if::value, T>::type -log10(T x) __NOEXC { +detail::enable_if_t::value, T> log10(T x) __NOEXC { return __sycl_std::__invoke_native_log10(x); } // genfloatf powr (genfloatf x, genfloatf y) template -typename std::enable_if::value, T>::type -powr(T x, T y) __NOEXC { +detail::enable_if_t::value, T> powr(T x, T y) __NOEXC { return __sycl_std::__invoke_native_powr(x, y); } // genfloatf recip (genfloatf x) template -typename std::enable_if::value, T>::type -recip(T x) __NOEXC { +detail::enable_if_t::value, T> recip(T x) __NOEXC { return __sycl_std::__invoke_native_recip(x); } // genfloatf rsqrt (genfloatf x) template -typename std::enable_if::value, T>::type -rsqrt(T x) __NOEXC { +detail::enable_if_t::value, T> rsqrt(T x) __NOEXC { return __sycl_std::__invoke_native_rsqrt(x); } // genfloatf sin (genfloatf x) template -typename std::enable_if::value, T>::type -sin(T x) __NOEXC { +detail::enable_if_t::value, T> sin(T x) __NOEXC { return __sycl_std::__invoke_native_sin(x); } // genfloatf sqrt (genfloatf x) template -typename std::enable_if::value, T>::type -sqrt(T x) __NOEXC { +detail::enable_if_t::value, T> sqrt(T x) __NOEXC { return __sycl_std::__invoke_native_sqrt(x); } // genfloatf tan (genfloatf x) template -typename std::enable_if::value, T>::type -tan(T x) __NOEXC { +detail::enable_if_t::value, T> tan(T x) __NOEXC { return __sycl_std::__invoke_native_tan(x); } @@ -1548,99 +1441,86 @@ namespace half_precision { /* ----------------- 4.13.3 Math functions. ---------------------------------*/ // genfloatf cos (genfloatf x) template -typename std::enable_if::value, T>::type -cos(T x) __NOEXC { +detail::enable_if_t::value, T> cos(T x) __NOEXC { return __sycl_std::__invoke_half_cos(x); } // genfloatf divide (genfloatf x, genfloatf y) template -typename std::enable_if::value, T>::type -divide(T x, T y) __NOEXC { +detail::enable_if_t::value, T> divide(T x, + T y) __NOEXC { return __sycl_std::__invoke_half_divide(x, y); } // genfloatf exp (genfloatf x) template -typename std::enable_if::value, T>::type -exp(T x) __NOEXC { +detail::enable_if_t::value, T> exp(T x) __NOEXC { return __sycl_std::__invoke_half_exp(x); } // genfloatf exp2 (genfloatf x) template -typename std::enable_if::value, T>::type -exp2(T x) __NOEXC { +detail::enable_if_t::value, T> exp2(T x) __NOEXC { return __sycl_std::__invoke_half_exp2(x); } // genfloatf exp10 (genfloatf x) template -typename std::enable_if::value, T>::type -exp10(T x) __NOEXC { +detail::enable_if_t::value, T> exp10(T x) __NOEXC { return __sycl_std::__invoke_half_exp10(x); } // genfloatf log (genfloatf x) template -typename std::enable_if::value, T>::type -log(T x) __NOEXC { +detail::enable_if_t::value, T> log(T x) __NOEXC { return __sycl_std::__invoke_half_log(x); } // genfloatf log2 (genfloatf x) template -typename std::enable_if::value, T>::type -log2(T x) __NOEXC { +detail::enable_if_t::value, T> log2(T x) __NOEXC { return __sycl_std::__invoke_half_log2(x); } // genfloatf log10 (genfloatf x) template -typename std::enable_if::value, T>::type -log10(T x) __NOEXC { +detail::enable_if_t::value, T> log10(T x) __NOEXC { return __sycl_std::__invoke_half_log10(x); } // genfloatf powr (genfloatf x, genfloatf y) template -typename std::enable_if::value, T>::type -powr(T x, T y) __NOEXC { +detail::enable_if_t::value, T> powr(T x, T y) __NOEXC { return __sycl_std::__invoke_half_powr(x, y); } // genfloatf recip (genfloatf x) template -typename std::enable_if::value, T>::type -recip(T x) __NOEXC { +detail::enable_if_t::value, T> recip(T x) __NOEXC { return __sycl_std::__invoke_half_recip(x); } // genfloatf rsqrt (genfloatf x) template -typename std::enable_if::value, T>::type -rsqrt(T x) __NOEXC { +detail::enable_if_t::value, T> rsqrt(T x) __NOEXC { return __sycl_std::__invoke_half_rsqrt(x); } // genfloatf sin (genfloatf x) template -typename std::enable_if::value, T>::type -sin(T x) __NOEXC { +detail::enable_if_t::value, T> sin(T x) __NOEXC { return __sycl_std::__invoke_half_sin(x); } // genfloatf sqrt (genfloatf x) template -typename std::enable_if::value, T>::type -sqrt(T x) __NOEXC { +detail::enable_if_t::value, T> sqrt(T x) __NOEXC { return __sycl_std::__invoke_half_sqrt(x); } // genfloatf tan (genfloatf x) template -typename std::enable_if::value, T>::type -tan(T x) __NOEXC { +detail::enable_if_t::value, T> tan(T x) __NOEXC { return __sycl_std::__invoke_half_tan(x); } diff --git a/sycl/include/CL/sycl/detail/generic_type_lists.hpp b/sycl/include/CL/sycl/detail/generic_type_lists.hpp new file mode 100644 index 0000000000000..c777a0553d768 --- /dev/null +++ b/sycl/include/CL/sycl/detail/generic_type_lists.hpp @@ -0,0 +1,375 @@ +//==-------- generic_type_lists.hpp - SYCL Generic type lists --------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include + +// Generic type name description, which serves as a description for all valid +// types of parameters to kernel functions + +// Forward declaration +namespace cl { +namespace sycl { +template class vec; +} // namespace sycl +} // namespace cl + +namespace cl { +namespace sycl { +namespace detail { +namespace gtl { +// floating point types +using scalar_half_list = type_list; + +using vector_half_list = type_list, vec, vec, + vec, vec, vec>; + +using half_list = type_list; + +using scalar_float_list = type_list; + +using vector_float_list = + type_list, vec, vec, vec, + vec, vec>; + +using float_list = type_list; + +using scalar_double_list = type_list; + +using vector_double_list = + type_list, vec, vec, vec, + vec, vec>; + +using double_list = type_list; + +using scalar_floating_list = + type_list; + +using vector_floating_list = + type_list; + +using floating_list = type_list; + +// geometric floating point types +using scalar_geo_half_list = type_list; + +using scalar_geo_float_list = type_list; + +using scalar_geo_double_list = type_list; + +using vector_geo_half_list = + type_list, vec, vec, vec>; + +using vector_geo_float_list = + type_list, vec, vec, vec>; + +using vector_geo_double_list = + type_list, vec, vec, vec>; + +using geo_half_list = type_list; + +using geo_float_list = type_list; + +using geo_double_list = + type_list; + +using scalar_geo_list = type_list; + +using vector_geo_list = type_list; + +using geo_list = type_list; + +// cross floating point types +using cross_half_list = type_list, vec>; + +using cross_float_list = type_list, vec>; + +using cross_double_list = type_list, vec>; + +using cross_floating_list = + type_list; + +using scalar_default_char_list = type_list; + +using vector_default_char_list = + type_list, vec, vec, vec, + vec, vec>; + +using default_char_list = + type_list; + +using scalar_signed_char_list = type_list; + +using vector_signed_char_list = + type_list, vec, vec, + vec, vec, vec>; + +using signed_char_list = + type_list; + +using scalar_unsigned_char_list = type_list; + +using vector_unsigned_char_list = + type_list, vec, + vec, vec, + vec, vec>; + +using unsigned_char_list = + type_list; + +using scalar_char_list = + type_list; + +using vector_char_list = + type_list; + +using char_list = type_list; + +// short int types +using scalar_signed_short_list = type_list; + +using vector_signed_short_list = + type_list, vec, vec, + vec, vec, + vec>; + +using signed_short_list = + type_list; + +using scalar_unsigned_short_list = type_list; + +using vector_unsigned_short_list = + type_list, vec, + vec, vec, + vec, vec>; + +using unsigned_short_list = + type_list; + +using scalar_short_list = + type_list; + +using vector_short_list = + type_list; + +using short_list = type_list; + +// int types +using scalar_signed_int_list = type_list; + +using vector_signed_int_list = + type_list, vec, vec, + vec, vec, vec>; + +using signed_int_list = + type_list; + +using scalar_unsigned_int_list = type_list; + +using vector_unsigned_int_list = + type_list, vec, vec, + vec, vec, + vec>; + +using unsigned_int_list = + type_list; + +using scalar_int_list = + type_list; + +using vector_int_list = + type_list; + +using int_list = type_list; + +// long types +using scalar_signed_long_list = type_list; + +using vector_signed_long_list = + type_list, vec, vec, + vec, vec, vec>; + +using signed_long_list = + type_list; + +using scalar_unsigned_long_list = type_list; + +using vector_unsigned_long_list = + type_list, vec, + vec, vec, + vec, vec>; + +using unsigned_long_list = + type_list; + +using scalar_long_list = + type_list; + +using vector_long_list = + type_list; + +using long_list = type_list; + +// long long types +using scalar_signed_longlong_list = type_list; + +using vector_signed_longlong_list = + type_list, vec, + vec, vec, + vec, vec>; + +using signed_longlong_list = + type_list; + +using scalar_unsigned_longlong_list = type_list; + +using vector_unsigned_longlong_list = + type_list, vec, + vec, vec, + vec, vec>; + +using unsigned_longlong_list = + type_list; + +using scalar_longlong_list = + type_list; + +using vector_longlong_list = + type_list; + +using longlong_list = type_list; + +// long integer types +using scalar_signed_long_integer_list = + type_list; + +using vector_signed_long_integer_list = + type_list; + +using signed_long_integer_list = + type_list; + +using scalar_unsigned_long_integer_list = + type_list; + +using vector_unsigned_long_integer_list = + type_list; + +using unsigned_long_integer_list = type_list; + +using scalar_long_integer_list = type_list; + +using vector_long_integer_list = type_list; + +using long_integer_list = + type_list; + +// integer types +using scalar_signed_integer_list = type_list< + conditional_t::value, + type_list, + scalar_signed_char_list>, + scalar_signed_short_list, scalar_signed_int_list, scalar_signed_long_list, + scalar_signed_longlong_list>; + +using vector_signed_integer_list = type_list< + conditional_t::value, + type_list, + vector_signed_char_list>, + vector_signed_short_list, vector_signed_int_list, vector_signed_long_list, + vector_signed_longlong_list>; + +using signed_integer_list = + type_list; + +using scalar_unsigned_integer_list = + type_list::value, + type_list, + scalar_unsigned_char_list>, + scalar_unsigned_short_list, scalar_unsigned_int_list, + scalar_unsigned_long_list, scalar_unsigned_longlong_list>; + +using vector_unsigned_integer_list = + type_list::value, + type_list, + vector_unsigned_char_list>, + vector_unsigned_short_list, vector_unsigned_int_list, + vector_unsigned_long_list, vector_unsigned_longlong_list>; + +using unsigned_integer_list = + type_list; + +using scalar_integer_list = + type_list; + +using vector_integer_list = + type_list; + +using integer_list = type_list; + +// basic types +using scalar_signed_basic_list = + type_list; + +using vector_signed_basic_list = + type_list; + +using signed_basic_list = + type_list; + +using scalar_unsigned_basic_list = type_list; + +using vector_unsigned_basic_list = type_list; + +using unsigned_basic_list = + type_list; + +using scalar_basic_list = + type_list; + +using vector_basic_list = + type_list; + +using basic_list = type_list; + +// nan builtin types +using nan_list = type_list; +} // namespace gtl +namespace gvl { +// address spaces +using all_address_space_list = + address_space_list; + +using nonconst_address_space_list = + address_space_list; + +using nonlocal_address_space_list = + address_space_list; +} // namespace gvl +} // namespace detail +} // namespace sycl +} // namespace cl diff --git a/sycl/include/CL/sycl/detail/generic_type_traits.hpp b/sycl/include/CL/sycl/detail/generic_type_traits.hpp index 1816793ce6e64..459fea38a469e 100644 --- a/sycl/include/CL/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/CL/sycl/detail/generic_type_traits.hpp @@ -8,8 +8,12 @@ #pragma once -#include -#include +#include +#include +#include +#include +#include +#include #include #include @@ -18,810 +22,246 @@ namespace cl { namespace sycl { namespace detail { -template struct type_list; +template using is_floatn = is_contained; -template struct type_list { - using head = H; - using tail = type_list; -}; +template using is_genfloatf = is_contained; + +template +using is_doublen = is_contained; + +template using is_genfloatd = is_contained; -template <> struct type_list<> {}; +template using is_halfn = is_contained; -template -struct is_contained - : std::conditional::type, - typename TL::head>::value, - std::true_type, - is_contained>::type {}; +template using is_genfloath = is_contained; -template struct is_contained> : std::false_type {}; +template using is_genfloat = is_contained; -// floatn: float2, float3, float4, float8, float16 template -using is_floatn = typename is_contained< - T, type_list>::type; +using is_sgenfloat = is_contained; -// genfloatf: float, floatn template -using is_genfloatf = - std::integral_constant>::value || - is_floatn::value>; +using is_vgenfloat = is_contained; -// doublen: double2, double3, double4, double8, double16 template -using is_doublen = - typename is_contained>::type; +using is_gengeofloat = is_contained; -// genfloatd: double, doublen template -using is_genfloatd = - std::integral_constant>::value || - is_doublen::value>; +using is_gengeodouble = is_contained; + +template using is_gengeohalf = is_contained; -// halfn: half2, half3, half4, half8, half16 template -using is_halfn = typename is_contained< - T, type_list>::type; +using is_vgengeofloat = is_contained; -// genfloath: half, halfn template -using is_genfloath = - std::integral_constant>::value || - is_halfn::value>; +using is_vgengeodouble = is_contained; -// genfloat: genfloatf, genfloatd, genfloath template -using is_genfloat = std::integral_constant::value || - is_genfloatd::value || - is_genfloath::value>; +using is_vgengeohalf = is_contained; + +template using is_sgengeo = is_contained; + +template using is_vgengeo = is_contained; -// sgenfloat: float, double, half template -using is_sgenfloat = - typename is_contained>::type; +using is_gencrossfloat = is_contained; -// vgenfloat: floatn, doublen, halfn template -using is_vgenfloat = - std::integral_constant::value || is_doublen::value || - is_halfn::value>; +using is_gencrossdouble = is_contained; -// gengeofloat: float, float2, float3, float4 template -using is_gengeofloat = typename is_contained< - T, type_list>::type; +using is_gencrosshalf = is_contained; -// gengeodouble: double, double2, double3, double4 template -using is_gengeodouble = typename is_contained< - T, type_list>::type; +using is_gencross = is_contained; -// gengeohalf: half, half2, half3, half4 template -using is_gengeohalf = typename is_contained< - T, type_list>::type; +using is_charn = is_contained; -// gengeofloat: float, float2, float3, float4 template -using is_vgengeofloat = - typename is_contained>::type; +using is_scharn = is_contained; -// gengeodouble: double, double2, double3, double4 template -using is_vgengeodouble = - typename is_contained>::type; +using is_ucharn = is_contained; -// gengeohalf: half2, half3, half4 template -using is_vgengeohalf = - typename is_contained>::type; +using is_igenchar = is_contained; -// sgengeo: float, double, half template -using is_sgengeo = std::integral_constant< - bool, is_contained>::value>; +using is_ugenchar = is_contained; + +template using is_genchar = is_contained; -// vgengeo: vgengeofloat, vgengeodouble, vgengeohalf template -using is_vgengeo = - std::integral_constant::value || - is_vgengeodouble::value || - is_vgengeohalf::value>; +using is_shortn = is_contained; -// gencrossfloat: float3, float4 template -using is_gencrossfloat = - typename is_contained>::type; +using is_genshort = is_contained; -// gencrossdouble: double3, double4 template -using is_gencrossdouble = - typename is_contained>::type; +using is_ushortn = is_contained; -// gencrosshalf: half3, half4 template -using is_gencrosshalf = - typename is_contained>::type; +using is_ugenshort = is_contained; -// gencross: gencrossfloat, gencrossdouble, gencrosshalf template -using is_gencross = - std::integral_constant::value || - is_gencrossdouble::value || - is_gencrosshalf::value>; +using is_uintn = is_contained; -// charn: char2, char3, char4, char8, char16 template -using is_charn = typename is_contained< - T, type_list>::type; +using is_ugenint = is_contained; -// scharn: schar2, schar3, schar4, schar8, schar16 template -using is_scharn = typename is_contained< - T, type_list>::type; +using is_intn = is_contained; + +template using is_genint = is_contained; -// ucharn: uchar2, uchar3, uchar4, uchar8, uchar16 template -using is_ucharn = typename is_contained< - T, type_list>::type; +using is_ulongn = is_contained; -// igenchar: signed char, scharn template -using is_igenchar = - std::integral_constant>::value || - is_scharn::value>; +using is_ugenlong = is_contained; -// ugenchar: unsigned char, ucharn template -using is_ugenchar = - std::integral_constant>::value || - is_ucharn::value>; +using is_longn = is_contained; + +template using is_genlong = is_contained; -// genchar: char, charn, igenchar, ugenchar template -using is_genchar = std::integral_constant< - bool, is_contained>::value || is_charn::value || - is_igenchar::value || is_ugenchar::value>; +using is_ulonglongn = is_contained; -// shortn: short2, short3, short4, short8, short16 template -using is_shortn = typename is_contained< - T, type_list>::type; +using is_ugenlonglong = is_contained; -// genshort: short, shortn template -using is_genshort = - std::integral_constant>::value || - is_shortn::value>; +using is_longlongn = is_contained; -// ushortn: ushort2, ushort3, ushort4, ushort8, ushort16 template -using is_ushortn = - typename is_contained>::type; +using is_genlonglong = is_contained; -// genushort: ushort, ushortn template -using is_ugenshort = - std::integral_constant>::value || - is_ushortn::value>; +using is_igenlonginteger = is_contained; -// uintn: uint2, uint3, uint4, uint8, uint16 template -using is_uintn = typename is_contained< - T, type_list>::type; +using is_ugenlonginteger = is_contained; + +template using is_geninteger = is_contained; -// ugenint: unsigned int, uintn template -using is_ugenint = - std::integral_constant>::value || - is_uintn::value>; +using is_igeninteger = is_contained; -// intn: int2, int3, int4, int8, int16 template -using is_intn = typename is_contained< - T, type_list>::type; +using is_ugeninteger = is_contained; -// genint: int, intn template -using is_genint = - std::integral_constant>::value || - is_intn::value>; +using is_sgeninteger = is_contained; -// ulongn: ulong2, ulong3, ulong4, ulong8,ulong16 template -using is_ulongn = typename is_contained< - T, type_list>::type; +using is_vgeninteger = is_contained; -// ugenlong: unsigned long int, ulongn template -using is_ugenlong = - std::integral_constant>::value || - is_ulongn::value>; +using is_sigeninteger = is_contained; -// longn: long2, long3, long4, long8, long16 template -using is_longn = typename is_contained< - T, type_list>::type; +using is_sugeninteger = is_contained; -// genlong: long int, longn template -using is_genlong = - std::integral_constant>::value || - is_longn::value>; +using is_vigeninteger = is_contained; -// ulonglongn: ulonglong2, ulonglong3, ulonglong4,ulonglong8, ulonglong16 template -using is_ulonglongn = - typename is_contained>::type; +using is_vugeninteger = is_contained; + +template using is_gentype = is_contained; -// ugenlonglong: unsigned long long int, ulonglongn template -using is_ugenlonglong = - std::integral_constant>::value || - is_ulonglongn::value>; +using is_vgentype = is_contained; -// longlongn: longlong2, longlong3, longlong4,longlong8, longlong16 template -using is_longlongn = typename is_contained< - T, type_list>::type; +using is_sgentype = is_contained; -// genlonglong: long long int, longlongn template -using is_genlonglong = - std::integral_constant>::value || - is_longlongn::value>; +using is_igeninteger8bit = is_gen_based_on_type_sizeof; -// igenlonginteger: genlong, genlonglong template -using is_igenlonginteger = - std::integral_constant::value || is_genlonglong::value>; +using is_igeninteger16bit = is_gen_based_on_type_sizeof; -// ugenlonginteger ugenlong, ugenlonglong template -using is_ugenlonginteger = - std::integral_constant::value || is_ugenlonglong::value>; +using is_igeninteger32bit = is_gen_based_on_type_sizeof; -// geninteger: genchar, genshort, ugenshort, genint, ugenint, igenlonginteger, -// ugenlonginteger template -using is_geninteger = std::integral_constant< - bool, is_genchar::value || is_genshort::value || - is_ugenshort::value || is_genint::value || - is_ugenint::value || is_igenlonginteger::value || - is_ugenlonginteger::value>; +using is_igeninteger64bit = is_gen_based_on_type_sizeof; -// igeninteger: igenchar, genshort, genint, igenlonginteger template -using is_igeninteger = std::integral_constant< - bool, is_igenchar::value || is_genshort::value || - is_genint::value || is_igenlonginteger::value>; +using is_ugeninteger8bit = is_gen_based_on_type_sizeof; -// ugeninteger: ugenchar, ugenshort, ugenint, ugenlonginteger template -using is_ugeninteger = std::integral_constant< - bool, is_ugenchar::value || is_ugenshort::value || - is_ugenint::value || is_ugenlonginteger::value>; +using is_ugeninteger16bit = is_gen_based_on_type_sizeof; -// sgeninteger: char, signed char, unsigned char, short, unsigned short, int, -// unsigned int, long int, unsigned long int, long long int, unsigned long long -// int template -using is_sgeninteger = typename is_contained< - T, type_list>::type; +using is_ugeninteger32bit = is_gen_based_on_type_sizeof; -// vgeninteger: charn, scharn, ucharn, shortn, ushortn, intn, uintn, longn, -// ulongn, longlongn, ulonglongn template -using is_vgeninteger = std::integral_constant< - bool, is_charn::value || is_scharn::value || is_ucharn::value || - is_shortn::value || is_ushortn::value || - is_intn::value || is_uintn::value || is_longn::value || - is_ulongn::value || is_longlongn::value || - is_ulonglongn::value>; +using is_ugeninteger64bit = is_gen_based_on_type_sizeof; -// sigeninteger: char, signed char, short, int, long int, , long long int template -using is_sigeninteger = typename is_contained< - T, type_list>::type; +using is_geninteger8bit = is_gen_based_on_type_sizeof; -// sugeninteger: unsigned char, unsigned short, unsigned int, unsigned long -// int, unsigned long long int template -using is_sugeninteger = typename is_contained< - T, type_list>::type; +using is_geninteger16bit = is_gen_based_on_type_sizeof; -// vigeninteger: charn, scharn, shortn, intn, longn, longlongn template -using is_vigeninteger = - std::integral_constant::value || is_scharn::value || - is_shortn::value || is_intn::value || - is_longn::value || - is_longlongn::value>; +using is_geninteger32bit = is_gen_based_on_type_sizeof; -// vugeninteger: ucharn, ushortn, uintn, ulongn, ulonglongn template -using is_vugeninteger = std::integral_constant< - bool, is_ucharn::value || is_ushortn::value || is_uintn::value || - is_ulongn::value || is_ulonglongn::value>; +using is_geninteger64bit = is_gen_based_on_type_sizeof; -// gentype: genfloat, geninteger template -using is_gentype = std::integral_constant::value || - is_geninteger::value>; +using is_genintptr = bool_constant< + is_pointer::value && is_genint>::value && + is_address_space_compliant::value>; -// vgentype: vgeninteger || vgenfloat template -using is_vgentype = std::integral_constant::value || is_vgenfloat::value>; +using is_genfloatptr = bool_constant< + is_pointer::value && is_genfloat>::value && + is_address_space_compliant::value>; -// sgentype: sgeninteger || sgenfloat template -using is_sgentype = std::integral_constant::value || is_sgenfloat::value>; +using is_genptr = bool_constant< + is_pointer::value && is_gentype>::value && + is_address_space_compliant::value>; -// forward declarations -template class TryToGetElementType; - -// genintegerNbit All types within geninteger whose base type are N bits in -// size, where N = 8, 16, 32, 64 -template -using is_igenintegerNbit = typename std::integral_constant< - bool, is_igeninteger::value && - (sizeof(typename TryToGetElementType::type) == N)>; - -// igeninteger8bit All types within igeninteger whose base type are 8 bits in -// size -template using is_igeninteger8bit = is_igenintegerNbit; - -// igeninteger16bit All types within igeninteger whose base type are 16 bits in -// size -template using is_igeninteger16bit = is_igenintegerNbit; - -// igeninteger32bit All types within igeninteger whose base type are 32 bits in -// size -template using is_igeninteger32bit = is_igenintegerNbit; - -// igeninteger64bit All types within igeninteger whose base type are 64 bits in -// size -template using is_igeninteger64bit = is_igenintegerNbit; - -// ugenintegerNbit All types within ugeninteger whose base type are N bits in -// size, where N = 8, 16, 32, 64. -template -using is_ugenintegerNbit = typename std::integral_constant< - bool, is_ugeninteger::value && - (sizeof(typename TryToGetElementType::type) == N)>; - -// ugeninteger8bit All types within ugeninteger whose base type are 8 bits in -// size -template using is_ugeninteger8bit = is_ugenintegerNbit; - -// ugeninteger16bit All types within ugeninteger whose base type are 16 bits in -// size -template using is_ugeninteger16bit = is_ugenintegerNbit; - -// ugeninteger32bit All types within ugeninteger whose base type are 32 bits in -// size -template using is_ugeninteger32bit = is_ugenintegerNbit; - -// ugeninteger64bit All types within ugeninteger whose base type are 64 bits in -// size -template using is_ugeninteger64bit = is_ugenintegerNbit; - -// genintegerNbit All types within geninteger whose base type are N bits in -// size, where N = 8, 16, 32, 64. -template -using is_genintegerNbit = typename std::integral_constant< - bool, is_geninteger::value && - (sizeof(typename TryToGetElementType::type) == N)>; - -// geninteger8bit All types within geninteger whose base type are 8 bits in size -template using is_geninteger8bit = is_genintegerNbit; - -// geninteger16bit All types within geninteger whose base type are 16 bits in -// size -template using is_geninteger16bit = is_genintegerNbit; - -// geninteger32bit All types within geninteger whose base type are 32 bits in -// size -template using is_geninteger32bit = is_genintegerNbit; - -// geninteger64bit All types within geninteger whose base type are 64 bits in -// size -template using is_geninteger64bit = is_genintegerNbit; - -template -using is_MultiPtrOfGLR = - std::integral_constant>::value || - std::is_same>::value || - std::is_same>::value || - std::is_same::value>; - -// genintptr All permutations of multi_ptr where dataT is -// all types within genint and addressSpace is -// access::address_space::global_space, access::address_space::local_space and -// access::address_space::private_space -template -using is_genintptr = - std::integral_constant::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value>; - -// is_genintegralptr is multi_ptr on any of intergral type including 'int' -// and 'long long'. -template -using is_genintegralptr = - std::integral_constant::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_genintptr

::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value>; - -// genfloatptr All permutations of multi_ptr where dataT is -// all types within genfloat and addressSpace is -// access::address_space::global_space, access::address_space::local_space and -// access::address_space::private_space -template -using is_genfloatptr = - std::integral_constant::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value || - is_MultiPtrOfGLR::value>; - -// is_genptr = is_genintegralptr || is_genfloatptr -template -using is_genptr = - std::integral_constant::value || - is_genfloatptr::value>; - -template struct unsign_integral_to_float_point; -template <> struct unsign_integral_to_float_point { - using type = cl_float; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_float2; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_float3; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_float4; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_float8; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_float16; -}; +template using is_nan_type = is_contained; -template <> struct unsign_integral_to_float_point { - using type = cl_half; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_half2; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_half3; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_half4; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_half8; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_half16; -}; +template using nan_return_t = typename nan_types::ret_type; -template <> struct unsign_integral_to_float_point { - using type = cl_double; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_double2; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_double3; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_double4; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_double8; -}; -template <> struct unsign_integral_to_float_point { - using type = cl_double16; -}; +template +using nan_argument_base_t = typename nan_types::arg_type; template -using is_nan_type = - std::integral_constant::value || - detail::is_ugenint::value || - detail::is_ugenlonginteger::value>; +using make_floating_point_t = make_type_t; -// Used for some relational functions -template struct float_point_to_sign_integral; -template <> struct float_point_to_sign_integral { - using type = cl_int; -}; -template <> struct float_point_to_sign_integral { - using type = cl_int2; -}; -template <> struct float_point_to_sign_integral { - using type = cl_int3; -}; -template <> struct float_point_to_sign_integral { - using type = cl_int4; -}; -template <> struct float_point_to_sign_integral { - using type = cl_int8; -}; -template <> struct float_point_to_sign_integral { - using type = cl_int16; -}; +template +using make_singed_integer_t = make_type_t; -template <> struct float_point_to_sign_integral { - using type = cl_int; -}; -template <> struct float_point_to_sign_integral { - using type = cl_short2; -}; -template <> struct float_point_to_sign_integral { - using type = cl_short3; -}; -template <> struct float_point_to_sign_integral { - using type = cl_short4; -}; -template <> struct float_point_to_sign_integral { - using type = cl_short8; -}; -template <> struct float_point_to_sign_integral { - using type = cl_short16; -}; +template +using make_unsinged_integer_t = + make_type_t; -template <> struct float_point_to_sign_integral { - using type = cl_int; -}; -template <> struct float_point_to_sign_integral { - using type = cl_long2; -}; -template <> struct float_point_to_sign_integral { - using type = cl_long3; -}; -template <> struct float_point_to_sign_integral { - using type = cl_long4; -}; -template <> struct float_point_to_sign_integral { - using type = cl_long8; +template +struct convert_data_type_impl; + +template +struct convert_data_type_impl::value, T>> { + B operator()(T t) { return static_cast(t); } }; -template <> struct float_point_to_sign_integral { - using type = cl_long16; + +template +struct convert_data_type_impl::value, T>> { + vec operator()(T t) { return t.template convert(); } }; -// Used for ilogb built-in -template struct float_point_to_int; -template <> struct float_point_to_int { using type = cl_int; }; -template <> struct float_point_to_int { using type = cl_int2; }; -template <> struct float_point_to_int { using type = cl_int3; }; -template <> struct float_point_to_int { using type = cl_int4; }; -template <> struct float_point_to_int { using type = cl_int8; }; -template <> struct float_point_to_int { using type = cl_int16; }; - -template <> struct float_point_to_int { using type = cl_int; }; -template <> struct float_point_to_int { using type = cl_int2; }; -template <> struct float_point_to_int { using type = cl_int3; }; -template <> struct float_point_to_int { using type = cl_int4; }; -template <> struct float_point_to_int { using type = cl_int8; }; -template <> struct float_point_to_int { using type = cl_int16; }; - -template <> struct float_point_to_int { using type = cl_int; }; -template <> struct float_point_to_int { using type = cl_int2; }; -template <> struct float_point_to_int { using type = cl_int3; }; -template <> struct float_point_to_int { using type = cl_int4; }; -template <> struct float_point_to_int { using type = cl_int8; }; -template <> struct float_point_to_int { using type = cl_int16; }; - -// Used for abs and abs_diff built-in -template struct make_unsigned { using type = T; }; - -template <> struct make_unsigned { - using type = unsigned char; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; - -template <> struct make_unsigned { - using type = unsigned short; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; - -template <> struct make_unsigned { - using type = unsigned int; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; - -template <> struct make_unsigned { - using type = unsigned long; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; - -template <> struct make_unsigned { - using type = unsigned long long; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; -template <> struct make_unsigned> { - using type = cl::sycl::vec; }; - -// Used for upsample built-in -// Bases on Table 4.93: Scalar data type aliases supported by SYCL -template struct make_upper; - -template <> struct make_upper { using type = cl_short; }; -template <> struct make_upper { using type = cl_short2; }; -template <> struct make_upper { using type = cl_short3; }; -template <> struct make_upper { using type = cl_short4; }; -template <> struct make_upper { using type = cl_short8; }; -template <> struct make_upper { using type = cl_short16; }; - -template <> struct make_upper { using type = cl_ushort; }; -template <> struct make_upper { using type = cl_ushort2; }; -template <> struct make_upper { using type = cl_ushort3; }; -template <> struct make_upper { using type = cl_ushort4; }; -template <> struct make_upper { using type = cl_ushort8; }; -template <> struct make_upper { using type = cl_ushort16; }; - -template <> struct make_upper { using type = cl_int; }; -template <> struct make_upper { using type = cl_int2; }; -template <> struct make_upper { using type = cl_int3; }; -template <> struct make_upper { using type = cl_int4; }; -template <> struct make_upper { using type = cl_int8; }; -template <> struct make_upper { using type = cl_int16; }; - -template <> struct make_upper { using type = cl_uint; }; -template <> struct make_upper { using type = cl_uint2; }; -template <> struct make_upper { using type = cl_uint3; }; -template <> struct make_upper { using type = cl_uint4; }; -template <> struct make_upper { using type = cl_uint8; }; -template <> struct make_upper { using type = cl_uint16; }; - -template <> struct make_upper { using type = cl_long; }; -template <> struct make_upper { using type = cl_long2; }; -template <> struct make_upper { using type = cl_long3; }; -template <> struct make_upper { using type = cl_long4; }; -template <> struct make_upper { using type = cl_long8; }; -template <> struct make_upper { using type = cl_long16; }; - -template <> struct make_upper { using type = cl_ulong; }; -template <> struct make_upper { using type = cl_ulong2; }; -template <> struct make_upper { using type = cl_ulong3; }; -template <> struct make_upper { using type = cl_ulong4; }; -template <> struct make_upper { using type = cl_ulong8; }; -template <> struct make_upper { using type = cl_ulong16; }; - -template <> struct make_upper { using type = longlong; }; -template <> struct make_upper { using type = longlong2; }; -template <> struct make_upper { using type = longlong3; }; -template <> struct make_upper { using type = longlong4; }; -template <> struct make_upper { using type = longlong8; }; -template <> struct make_upper { using type = longlong16; }; - -template <> struct make_upper { using type = ulonglong; }; -template <> struct make_upper { using type = ulonglong2; }; -template <> struct make_upper { using type = ulonglong3; }; -template <> struct make_upper { using type = ulonglong4; }; -template <> struct make_upper { using type = ulonglong8; }; -template <> struct make_upper { using type = ulonglong16; }; +template +using convert_data_type = convert_data_type_impl; // Try to get pointer_t, otherwise T template class TryToGetPointerT { @@ -830,8 +270,8 @@ template class TryToGetPointerT { public: using type = decltype(check(T())); - static constexpr bool value = std::is_pointer::value || - !std::is_same::value; + static constexpr bool value = + std::is_pointer::value || !std::is_same::value; }; // Try to get element_type, otherwise T @@ -854,8 +294,8 @@ template class TryToGetVectorT { static constexpr bool value = !std::is_same::value; }; -// Try to get pointer_t (if pointer_t indicates on the type with vector_t -// creates a pointer type on vector_t), otherwise T +// Try to get pointer_t (if pointer_t indicates on the type with_remainder +// vector_t creates a pointer type on vector_t), otherwise T template class TryToGetPointerVecT { static T check(...); template @@ -864,7 +304,7 @@ template class TryToGetPointerVecT { A::address_space>::type * check(const A &); template - static typename TryToGetVectorT::type * check(const A *); + static typename TryToGetVectorT::type *check(const A *); public: using type = decltype(check(T())); @@ -891,62 +331,59 @@ T TryToGetPointer(T &t) { return t; } -// Use conditional_t to improve code readablity. -template -using conditional_t = typename std::conditional::type; - // select_apply_cl_scalar_t selects from T8/T16/T32/T64 basing on // sizeof(IN). expected to handle scalar types. template using select_apply_cl_scalar_t = - conditional_t>>; + conditional_t>>; // Shortcuts for selecting scalar int/unsigned int/fp type. template using select_cl_scalar_intergal_signed_t = - select_apply_cl_scalar_t; + select_apply_cl_scalar_t; template using select_cl_scalar_intergal_unsigned_t = - select_apply_cl_scalar_t; + select_apply_cl_scalar_t; template using select_cl_scalar_float_t = - select_apply_cl_scalar_t; + select_apply_cl_scalar_t; template using select_cl_scalar_intergal_t = - conditional_t::value, - select_cl_scalar_intergal_signed_t, - select_cl_scalar_intergal_unsigned_t>; + conditional_t::value, + select_cl_scalar_intergal_signed_t, + select_cl_scalar_intergal_unsigned_t>; // select_cl_scalar_t picks corresponding cl_* type for input // scalar T or returns T if T is not scalar. template using select_cl_scalar_t = - conditional_t::value, - select_cl_scalar_intergal_t, - conditional_t::value, - select_cl_scalar_float_t, T>>; + conditional_t::value, select_cl_scalar_intergal_t, + conditional_t::value, + select_cl_scalar_float_t, T>>; // select_cl_vector_or_scalar does cl_* type selection for element type of // a vector type T and does scalar type substitution. If T is not // vector or scalar unmodified T is returned. -template -struct select_cl_vector_or_scalar; +template struct select_cl_vector_or_scalar; template struct select_cl_vector_or_scalar< - T, typename std::enable_if::value>::type> { - using type = vec, - T::get_count()>; + T, typename std::enable_if::value>::type> { + using type = + vec, T::get_count()>; }; template struct select_cl_vector_or_scalar< - T, typename std::enable_if::value>::type> { + T, typename std::enable_if::value>::type> { using type = select_cl_scalar_t; }; @@ -959,61 +396,57 @@ struct select_cl_mptr_or_vector_or_scalar; template struct select_cl_mptr_or_vector_or_scalar< - T, typename std::enable_if::value && - !std::is_pointer::value>::type> { + T, typename std::enable_if::value && + !std::is_pointer::value>::type> { using type = multi_ptr< - typename select_cl_vector_or_scalar::type, - T::address_space>; + typename select_cl_vector_or_scalar::type, + T::address_space>; }; template struct select_cl_mptr_or_vector_or_scalar< - T, typename std::enable_if::value || - std::is_pointer::value>::type> { + T, typename std::enable_if::value || + std::is_pointer::value>::type> { using type = typename select_cl_vector_or_scalar::type; }; // All types converting shortcut. template using SelectMatchingOpenCLType_t = - typename select_cl_mptr_or_vector_or_scalar::type; + typename select_cl_mptr_or_vector_or_scalar::type; // Converts T to OpenCL friendly // template -using ConvertToOpenCLType_t = - conditional_t>::value, +using ConvertToOpenCLType_t = conditional_t< + TryToGetVectorT>::value, typename TryToGetVectorT>::type, - conditional_t>::value, - typename TryToGetPointerVecT>::type, - SelectMatchingOpenCLType_t>>; + conditional_t< + TryToGetPointerT>::value, + typename TryToGetPointerVecT>::type, + SelectMatchingOpenCLType_t>>; // convertDataToType() function converts data from FROM type to TO type using // 'as' method for vector type and copy otherwise. template -typename std::enable_if::value && - is_vgentype::value && - sizeof(TO) == sizeof(FROM), TO>::type +typename std::enable_if::value && is_vgentype::value && + sizeof(TO) == sizeof(FROM), + TO>::type convertDataToType(FROM t) { return t.template as(); } template -typename std::enable_if::value && - is_vgentype::value) && - sizeof(TO) == sizeof(FROM), TO>::type +typename std::enable_if::value && is_vgentype::value) && + sizeof(TO) == sizeof(FROM), + TO>::type convertDataToType(FROM t) { return TryToGetPointer(t); } -template - using unsign_integral_to_float_point_t = - typename unsign_integral_to_float_point< - SelectMatchingOpenCLType_t>::type; - // Used for all,any and select relational built-in functions template inline constexpr T msbMask(T) { - using UT = typename std::make_unsigned::type; + using UT = make_unsigned_t; return T(UT(1) << (sizeof(T) * 8 - 1)); } @@ -1022,7 +455,8 @@ template inline constexpr bool msbIsSet(const T x) { } template -using common_rel_ret_t = typename detail::float_point_to_sign_integral::type; +using common_rel_ret_t = + conditional_t::value, make_singed_integer_t, int>; // forward declaration template struct Boolean; @@ -1070,7 +504,7 @@ template struct RelationalTestForSignBitType { using return_type = detail::Boolean<1>; using argument_type = detail::Boolean::value>; #else - using return_type = cl::sycl::cl_int; + using return_type = int; using argument_type = T; #endif }; diff --git a/sycl/include/CL/sycl/detail/stl_type_traits.hpp b/sycl/include/CL/sycl/detail/stl_type_traits.hpp new file mode 100644 index 0000000000000..6cb10b01ae3b5 --- /dev/null +++ b/sycl/include/CL/sycl/detail/stl_type_traits.hpp @@ -0,0 +1,74 @@ +//==------- stl_type_traits.hpp - SYCL STL type traits analogs -------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +namespace cl { +namespace sycl { +namespace detail { + +template using bool_constant = std::integral_constant; + +template +using allocator_value_type_t = typename std::allocator_traits::value_type; + +template +using allocator_pointer_t = typename std::allocator_traits::pointer; + +template +using enable_if_t = typename std::enable_if::type; + +template +using conditional_t = typename std::conditional::type; + +template +using remove_const_t = typename std::remove_const::type; + +template using remove_cv_t = typename std::remove_cv::type; + +template using add_pointer_t = typename std::add_pointer::type; + +template +using iterator_category_t = typename std::iterator_traits::iterator_category; + +template +using iterator_value_type_t = typename std::iterator_traits::value_type; + +template +using iterator_pointer_t = typename std::iterator_traits::pointer; + +template +using iterator_to_const_type_t = + std::is_const>::type>; + +template using requirements_list = void; + +// TODO Align with C++ named requirements: LegacyOutputIterator +// https://en.cppreference.com/w/cpp/named_req/OutputIterator +template +using output_iterator_requirements = + requirements_list, + decltype(*std::declval() = + std::declval>())>; + +template struct is_output_iterator { + static constexpr bool value = false; +}; + +template +struct is_output_iterator> { + static constexpr bool value = true; +}; + +} // namespace detail +} // namespace sycl +} // namespace cl diff --git a/sycl/include/CL/sycl/detail/type_list.hpp b/sycl/include/CL/sycl/detail/type_list.hpp new file mode 100644 index 0000000000000..fef9edbfbe388 --- /dev/null +++ b/sycl/include/CL/sycl/detail/type_list.hpp @@ -0,0 +1,140 @@ +//==----------- type_list.hpp - SYCL list of types utils -------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +#include + +namespace cl { +namespace sycl { +namespace detail { + +template using head_t = typename T::head; + +template using tail_t = typename T::tail; + +// type_list +template class type_list; + +using empty_type_list = type_list<>; + +template +struct is_empty_type_list + : conditional_t::value, std::true_type, + std::false_type> {}; + +template <> struct type_list<> {}; + +template struct type_list { + using head = H; + using tail = type_list; +}; + +template +struct type_list, T2...> { +private: + using remainder = tail_t>; + static constexpr bool has_remainder = !is_empty_type_list::value; + using without_remainder = type_list; + using with_remainder = type_list; + +public: + using head = head_t>; + using tail = conditional_t; +}; + +// is_contained +template +struct is_contained + : conditional_t, head_t>::value, + std::true_type, is_contained>> {}; + +template +struct is_contained : std::false_type {}; + +// value_list +template struct value_list; + +template struct value_list { + static constexpr T head = H; + using tail = value_list; +}; + +template struct value_list {}; + +// is_contained_value +template +struct is_contained_value + : conditional_t>> {}; + +template +struct is_contained_value> : std::false_type {}; + +// address_space_list +template +using address_space_list = value_list; + +template +using is_one_of_spaces = is_contained_value; + +// size type predicates +template +struct is_type_size_equal : bool_constant<(sizeof(T1) == sizeof(T2))> {}; + +template +struct is_type_size_greater : bool_constant<(sizeof(T1) > sizeof(T2))> {}; + +template +struct is_type_size_double_of + : bool_constant<(sizeof(T1) == (sizeof(T2) * 2))> {}; + +template +struct is_type_size_less : bool_constant<(sizeof(T1) < sizeof(T2))> {}; + +template +struct is_type_size_half_of : bool_constant<(sizeof(T1) == (sizeof(T2) / 2))> { +}; + +// find required type +template class C, typename T> +struct find_type { + using head = head_t; + using tail = typename find_type, C, T>::type; + using type = conditional_t::value, head, tail>; +}; + +template