diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp index f760d825d6a9d..1655c78c96352 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp @@ -16,7 +16,6 @@ #include // to define C++14,17 extensions #include #include -#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { @@ -208,8 +207,8 @@ template struct computation_type { template constexpr bool is_type() { return false; } template constexpr bool is_type() { - using UU = typename std::remove_const::type; - using TT = typename std::remove_const::type; + using UU = typename detail::remove_const_t; + using TT = typename detail::remove_const_t; return std::is_same::value || is_type(); } @@ -228,10 +227,10 @@ struct bitcast_helper { // Change the element type of a simd vector. template ::value>> -ESIMD_INLINE typename std::conditional< +ESIMD_INLINE typename detail::conditional_t< std::is_same::value, vector_type_t, vector_type_t::nToElems()>>::type + bitcast_helper::nToElems()>> bitcast(vector_type_t Val) { // Noop. if constexpr (std::is_same::value) diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp index 4bd5755c8f33d..6d0be7f5521ff 100755 --- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp @@ -100,10 +100,10 @@ struct is_esimd_scalar template struct is_dword_type : std::integral_constant< - bool, std::is_same::type>::value || - std::is_same::type>::value> { -}; + bool, + std::is_same>::value || + std::is_same>::value> {}; template struct is_dword_type> { @@ -119,9 +119,10 @@ template struct is_word_type : std::integral_constant< bool, - std::is_same::type>::value || + std::is_same>::value || std::is_same::type>::value> {}; + typename sycl::detail::remove_const_t>::value> {}; template struct is_word_type> { @@ -136,9 +137,9 @@ template struct is_byte_type : std::integral_constant< bool, - std::is_same::type>::value || + std::is_same>::value || std::is_same::type>::value> {}; + typename sycl::detail::remove_const_t>::value> {}; template struct is_byte_type> { @@ -152,31 +153,36 @@ template struct is_byte_type> { template struct is_fp_type : std::integral_constant< - bool, - std::is_same::type>::value> {}; + bool, std::is_same>::value> { +}; template struct is_df_type : std::integral_constant< - bool, - std::is_same::type>::value> {}; + bool, std::is_same>::value> { +}; template struct is_fp_or_dword_type : std::integral_constant< bool, - std::is_same::type>::value || - std::is_same::type>::value || + std::is_same>::value || + std::is_same>::value || std::is_same::type>::value> {}; + typename sycl::detail::remove_const_t>::value> {}; template struct is_qword_type : std::integral_constant< bool, - std::is_same::type>::value || + std::is_same>::value || std::is_same::type>::value> {}; + typename sycl::detail::remove_const_t>::value> {}; template struct is_qword_type> { diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp index 25f8e339fefd6..6a9251cbc33b6 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp @@ -328,8 +328,9 @@ template class simd { // // @return 1 if any element is set, 0 otherwise - template ::value, T2>> + template < + typename T1 = element_type, typename T2 = Ty, + typename = sycl::detail::enable_if_t::value, T2>> uint16_t any() { return __esimd_any(data()); } @@ -338,8 +339,9 @@ template class simd { // // @return 1 if all elements are set, 0 otherwise - template ::value, T2>> + template < + typename T1 = element_type, typename T2 = Ty, + typename = sycl::detail::enable_if_t::value, T2>> uint16_t all() { return __esimd_all(data()); } diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp index 3ac5f8d792769..bda406c12beb9 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp @@ -58,13 +58,12 @@ __esimd_abs_common_internal(simd src0, int flag = GENX_NOSAT) { } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - __esimd_abs_common_internal(T1 src0, int flag = GENX_NOSAT) { - typedef typename std::remove_const::type TT0; - typedef typename std::remove_const::type TT1; +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +__esimd_abs_common_internal(T1 src0, int flag = GENX_NOSAT) { + typedef typename sycl::detail::remove_const_t TT0; + typedef typename sycl::detail::remove_const_t TT1; simd Src0 = src0; simd Result = __esimd_abs_common_internal(Src0, flag); @@ -73,21 +72,21 @@ ESIMD_NODEBUG ESIMD_INLINE } // namespace detail template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - !std::is_same::type, - typename std::remove_const::type>::value, - simd>::type +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + !std::is_same, + typename sycl::detail::remove_const_t>::value, + simd> esimd_abs(simd src0, int flag = GENX_NOSAT) { return detail::__esimd_abs_common_internal(src0, flag); } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - !std::is_same::type, - typename std::remove_const::type>::value && +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + !std::is_same, + typename sycl::detail::remove_const_t>::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_abs(T1 src0, int flag = GENX_NOSAT) { return detail::__esimd_abs_common_internal(src0, flag); } @@ -99,20 +98,20 @@ ESIMD_NODEBUG ESIMD_INLINE simd esimd_abs(simd src0, } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - typename std::remove_const::type>::type - esimd_abs(T1 src0, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_abs(T1 src0, int flag = GENX_NOSAT) { return detail::__esimd_abs_common_internal(src0, flag); } // esimd_shl template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_shl(simd src0, U src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -146,11 +145,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_shl(T1 src0, T2 src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -162,10 +161,10 @@ esimd_shl(T1 src0, T2 src1, int flag = GENX_NOSAT) { // esimd_shr template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_shr(simd src0, U src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -180,11 +179,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_shr(T1 src0, T2 src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -195,20 +194,18 @@ esimd_shr(T1 src0, T2 src1, int flag = GENX_NOSAT) { // esimd_rol template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value, - simd>::type - esimd_rol(simd src0, simd src1) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && std::is_integral::value, simd> +esimd_rol(simd src0, simd src1) { return __esimd_rol(src0, src1); } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_rol(simd src0, U src1) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -217,11 +214,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_rol(T1 src0, T2 src1) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -232,20 +229,18 @@ esimd_rol(T1 src0, T2 src1) { // esimd_ror template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value, - simd>::type - esimd_ror(simd src0, simd src1) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && std::is_integral::value, simd> +esimd_ror(simd src0, simd src1) { return __esimd_ror(src0, src1); } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_ror(simd src0, U src1) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -254,11 +249,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_ror(T1 src0, T2 src1) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -270,10 +265,10 @@ esimd_ror(T1 src0, T2 src1) { // esimd_lsr template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_lsr(simd src0, U src1, int flag = GENX_NOSAT) { typedef typename computation_type::type IntermedTy; typedef typename std::make_unsigned::type ComputationTy; @@ -287,11 +282,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -301,11 +296,11 @@ esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) { } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_vector::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - decltype(esimd_lsr(T2(), T1()))>::type + decltype(esimd_lsr(T2(), T1()))> esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) { return esimd_lsr(src1, src0, flag); } @@ -313,10 +308,10 @@ esimd_lsr(T1 src0, T2 src1, int flag = GENX_NOSAT) { // esimd_asr template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - std::is_integral::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + std::is_integral::value, + simd> esimd_asr(simd src0, U src1, int flag = GENX_NOSAT) { typedef typename computation_type::type IntermedTy; typedef typename std::make_signed::type ComputationTy; @@ -330,11 +325,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -344,11 +339,11 @@ esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) { } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_vector::value && std::is_integral::value && std::is_integral::value && std::is_integral::value, - decltype(esimd_asr(T2(), T1()))>::type + decltype(esimd_asr(T2(), T1()))> esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) { return esimd_asr(src1, src0, flag); } @@ -358,10 +353,10 @@ esimd_asr(T1 src0, T2 src1, int flag = GENX_NOSAT) { // use mulh instruction for high half template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_dword_type::value && - detail::is_dword_type::value, - simd>::type + typename sycl::detail::enable_if_t::value && + detail::is_dword_type::value && + detail::is_dword_type::value, + simd> esimd_imul(simd &rmd, simd src0, U src1) { typedef typename computation_type::type ComputationTy; typename detail::simd_type::type Src0 = src0; @@ -378,12 +373,11 @@ ESIMD_NODEBUG ESIMD_INLINE // We need to special case SZ==1 to avoid "error: when select size is 1, the // stride must also be 1" on the selects. template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_dword_type::value && - detail::is_dword_type::value && SZ == 1, - simd>::type - esimd_imul(simd &rmd, simd src0, U src1) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_dword_type::value && detail::is_dword_type::value && + detail::is_dword_type::value && SZ == 1, + simd> +esimd_imul(simd &rmd, simd src0, U src1) { typedef typename computation_type::type ComputationTy; ComputationTy Product = convert(src0); @@ -393,12 +387,11 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_dword_type::value && - detail::is_dword_type::value && SZ != 1, - simd>::type - esimd_imul(simd &rmd, simd src0, U src1) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_dword_type::value && detail::is_dword_type::value && + detail::is_dword_type::value && SZ != 1, + simd> +esimd_imul(simd &rmd, simd src0, U src1) { typedef typename computation_type::type ComputationTy; ComputationTy Product = convert(src0); @@ -412,18 +405,18 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_imul(simd &rmd, U src0, simd src1) { return esimd_imul(rmd, src1, src0); } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value && - detail::is_esimd_scalar::value, - T0>::type + typename sycl::detail::enable_if_t::value && + detail::is_esimd_scalar::value && + detail::is_esimd_scalar::value, + T0> esimd_imul(simd &rmd, T src0, U src1) { simd src_0 = src0; simd src_1 = src1; @@ -433,42 +426,42 @@ ESIMD_NODEBUG ESIMD_INLINE // esimd_quot template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - std::is_integral::value && std::is_integral::value, simd>::type +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && std::is_integral::value, simd> esimd_quot(simd src0, U src1) { return src0 / src1; } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_quot(T0 src0, T1 src1) { return src0 / src1; } // esimd_mod template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - std::is_integral::value && std::is_integral::value, simd>::type +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && std::is_integral::value, simd> esimd_mod(simd src0, U src1) { return src0 % src1; } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && std::is_integral::value && std::is_integral::value, - typename std::remove_const::type>::type + typename sycl::detail::remove_const_t> esimd_mod(T0 src0, T1 src1) { return src0 % src1; } // esimd_div, compute quotient and remainder of division. template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - std::is_integral::value && std::is_integral::value, simd>::type +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && std::is_integral::value, simd> esimd_div(simd &remainder, simd src0, U src1) { remainder = src0 % src1; return src0 / src1; @@ -476,23 +469,22 @@ esimd_div(simd &remainder, simd src0, U src1) { template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_integral::value && - detail::is_esimd_scalar::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_integral::value && + detail::is_esimd_scalar::value, + simd> esimd_div(simd &remainder, U src0, simd src1) { remainder = src0 % src1; return src0 / src1; } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - esimd_div(simd::type, 1> &remainder, T0 src0, - T1 src1) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value && + detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_div(simd, 1> &remainder, T0 src0, + T1 src1) { remainder[0] = src0 % src1; return src0 / src1; } @@ -522,8 +514,8 @@ esimd_max(simd src0, simd src1, int flag = GENX_NOSAT) { template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_max(simd src0, T src1, int flag = GENX_NOSAT) { simd Src1 = src1; simd Result = esimd_max(src0, Src1, flag); @@ -532,8 +524,8 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_max(T src0, simd src1, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_max(Src0, src1, flag); @@ -542,7 +534,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, T>::type + typename sycl::detail::enable_if_t::value, T> esimd_max(T src0, T src1, int flag = GENX_NOSAT) { simd Src0 = src0; simd Src1 = src1; @@ -569,8 +561,8 @@ esimd_min(simd src0, simd src1, int flag = GENX_NOSAT) { template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_min(simd src0, T src1, int flag = GENX_NOSAT) { simd Src1 = src1; simd Result = esimd_min(src0, Src1, flag); @@ -579,8 +571,8 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, - simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_min(T src0, simd src1, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_min(Src0, src1, flag); @@ -588,7 +580,7 @@ ESIMD_NODEBUG ESIMD_INLINE } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, T>::type + typename sycl::detail::enable_if_t::value, T> esimd_min(T src0, T src1, int flag = GENX_NOSAT) { simd Src0 = src0; simd Src1 = src1; @@ -689,13 +681,13 @@ esimd_line(float P, float Q, simd src1, int flag = GENX_NOSAT) { // If the gen is not specified we warn the programmer that they are potentially // using a less efficient implementation if not on GEN10 or above. template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value && - detail::is_fp_or_dword_type::value && - std::is_floating_point::value, - simd>::type - esimd_dp2(simd src0, U src1, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_fp_or_dword_type::value && + std::is_floating_point::value && + detail::is_fp_or_dword_type::value && + std::is_floating_point::value, + simd> +esimd_dp2(simd src0, U src1, int flag = GENX_NOSAT) { static_assert(SZ % 4 == 0, "result size is not a multiple of 4"); simd Src1 = src1; @@ -711,13 +703,13 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value && - detail::is_fp_or_dword_type::value && - std::is_floating_point::value, - simd>::type - esimd_dp3(simd src0, U src1, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_fp_or_dword_type::value && + std::is_floating_point::value && + detail::is_fp_or_dword_type::value && + std::is_floating_point::value, + simd> +esimd_dp3(simd src0, U src1, int flag = GENX_NOSAT) { static_assert(SZ % 4 == 0, "result size is not a multiple of 4"); simd Src1 = src1; @@ -734,13 +726,13 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value && - detail::is_fp_or_dword_type::value && - std::is_floating_point::value, - simd>::type - esimd_dp4(simd src0, U src1, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_fp_or_dword_type::value && + std::is_floating_point::value && + detail::is_fp_or_dword_type::value && + std::is_floating_point::value, + simd> +esimd_dp4(simd src0, U src1, int flag = GENX_NOSAT) { static_assert(SZ % 4 == 0, "result size is not a multiple of 4"); simd Src1 = src1; @@ -758,13 +750,12 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value && - detail::is_fp_or_dword_type::value && - std::is_floating_point::value, - simd>::type - esimd_dph(simd src0, U src1, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_fp_or_dword_type::value && std::is_floating_point::value && + detail::is_fp_or_dword_type::value && + std::is_floating_point::value, + simd> +esimd_dph(simd src0, U src1, int flag = GENX_NOSAT) { static_assert(SZ % 4 == 0, "result size is not a multiple of 4"); simd Src1 = src1; @@ -782,9 +773,9 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_floating_point::value, + simd> esimd_line(simd src0, simd src1, int flag = GENX_NOSAT) { static_assert(SZ % 4 == 0, "result size is not a multiple of 4"); @@ -803,9 +794,9 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value, - simd>::type + typename sycl::detail::enable_if_t::value && + std::is_floating_point::value, + simd> esimd_line(float P, float Q, simd src1, int flag = GENX_NOSAT) { simd Src0 = P; Src0(3) = Q; @@ -836,11 +827,10 @@ ESIMD_NODEBUG ESIMD_INLINE simd esimd_lzd(simd src0, } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - esimd_lzd(T0 src0, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_lzd(T0 src0, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_lzd(Src0); return Result[0]; @@ -877,13 +867,12 @@ esimd_lrp(simd src0, U src1, V src2, int flag = GENX_NOSAT) { // If the gen is not specified we warn the programmer that they are potentially // using less efficient implementation. template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - std::is_floating_point::value && - detail::is_fp_or_dword_type::value && - std::is_floating_point::value, - simd>::type - esimd_lrp(simd src0, U src1, V src2, int flag = GENX_NOSAT) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_fp_or_dword_type::value && std::is_floating_point::value && + detail::is_fp_or_dword_type::value && + std::is_floating_point::value, + simd> +esimd_lrp(simd src0, U src1, V src2, int flag = GENX_NOSAT) { simd Src1 = src1; simd Src2 = src2; @@ -932,11 +921,10 @@ ESIMD_NODEBUG ESIMD_INLINE simd esimd_bf_reverse(simd src0) { } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - esimd_bf_reverse(T1 src0) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_bf_reverse(T1 src0) { simd Src0 = src0; simd Result = esimd_bf_reverse(Src0); return Result[0]; @@ -945,9 +933,10 @@ ESIMD_NODEBUG ESIMD_INLINE // esimd_bf_insert template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_bf_insert(U src0, V src1, W src2, simd src3) { - typedef typename detail::dword_type::type DT1; + typedef typename detail::dword_type DT1; static_assert(std::is_integral::value && sizeof(DT1) == sizeof(int), "operand conversion failed"); simd Src0 = src0; @@ -959,11 +948,10 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - esimd_bf_insert(T1 src0, T2 src1, T3 src2, T4 src3) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_bf_insert(T1 src0, T2 src1, T3 src2, T4 src3) { simd Src3 = src3; simd Result = esimd_bf_insert(src0, src1, src2, Src3); return Result[0]; @@ -972,9 +960,10 @@ ESIMD_NODEBUG ESIMD_INLINE // esimd_bf_extract template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_bf_extract(U src0, V src1, simd src2) { - typedef typename detail::dword_type::type DT1; + typedef typename detail::dword_type DT1; static_assert(std::is_integral::value && sizeof(DT1) == sizeof(int), "operand conversion failed"); simd Src0 = src0; @@ -985,11 +974,10 @@ ESIMD_NODEBUG ESIMD_INLINE } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - typename std::remove_const::type>::type - esimd_bf_extract(T1 src0, T2 src1, T3 src2) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_esimd_scalar::value && detail::is_esimd_scalar::value, + typename sycl::detail::remove_const_t> +esimd_bf_extract(T1 src0, T2 src1, T3 src2) { simd Src2 = src2; simd Result = esimd_bf_extract(src0, src1, Src2); return Result[0]; @@ -1072,8 +1060,8 @@ ESIMD_INTRINSIC_DEF(double, sqrt_ieee) } \ template \ ESIMD_NODEBUG ESIMD_INLINE \ - typename std::enable_if::value, \ - simd>::type \ + typename sycl::detail::enable_if_t::value, \ + simd> \ esimd_##name(U src0, simd src1, int flag = GENX_NOSAT) { \ simd Src0 = src0; \ return esimd_##name(Src0, src1, flag); \ @@ -1107,7 +1095,8 @@ esimd_sincos(simd &dstcos, U src0, int flag = GENX_NOSAT) { template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_atan(simd src0, int flag = GENX_NOSAT) { simd Src0 = esimd_abs(src0); @@ -1138,7 +1127,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, T>::type + typename sycl::detail::enable_if_t::value, T> esimd_atan(T src0, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_atan(Src0, flag); @@ -1149,7 +1138,8 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_acos(simd src0, int flag = GENX_NOSAT) { simd Src0 = esimd_abs(src0); @@ -1182,7 +1172,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, T>::type + typename sycl::detail::enable_if_t::value, T> esimd_acos(T src0, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_acos(Src0, flag); @@ -1193,7 +1183,8 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_asin(simd src0, int flag = GENX_NOSAT) { simd Neg = src0 < T(0.0); @@ -1210,7 +1201,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, T>::type + typename sycl::detail::enable_if_t::value, T> esimd_asin(T src0, int flag = GENX_NOSAT) { simd Src0 = src0; simd Result = esimd_asin(Src0, flag); @@ -1247,22 +1238,22 @@ ESIMD_INTRINSIC_DEF(rndz) template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if<(N == 8 || N == 16 || N == 32), uint>::type + typename sycl::detail::enable_if_t<(N == 8 || N == 16 || N == 32), uint> esimd_pack_mask(simd src0) { return __esimd_pack_mask(src0.data()); } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if<(N == 8 || N == 16 || N == 32), - simd>::type + typename sycl::detail::enable_if_t<(N == 8 || N == 16 || N == 32), + simd> esimd_unpack_mask(uint src0) { return __esimd_unpack_mask(src0); } template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if<(N != 8 && N != 16 && N < 32), uint>::type + typename sycl::detail::enable_if_t<(N != 8 && N != 16 && N < 32), uint> esimd_pack_mask(simd src0) { simd src_0 = 0; src_0.template select() = src0.template format(); @@ -1272,14 +1263,15 @@ ESIMD_NODEBUG ESIMD_INLINE /// Count component-wise the total bits set in source operand. template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value, simd>::type + typename sycl::detail::enable_if_t::value, + simd> esimd_cbit(simd src0) { return __esimd_cbit(src0.data()); } template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< - std::is_integral::value && detail::is_esimd_scalar::value, uint>::type +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + std::is_integral::value && detail::is_esimd_scalar::value, uint> esimd_cbit(T src) { simd Src = src; simd Result = esimd_cbit(Src); @@ -1312,11 +1304,9 @@ esimd_fbh(simd src) { } template -ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if::value && - detail::is_esimd_scalar::value, - T>::type - esimd_fbh(T src) { +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< + detail::is_dword_type::value && detail::is_esimd_scalar::value, T> +esimd_fbh(T src) { simd Src = src; simd Result = esimd_fbh(Src); return Result[0]; @@ -1325,10 +1315,10 @@ ESIMD_NODEBUG ESIMD_INLINE template simd esimd_rdtsc(); template -ESIMD_NODEBUG ESIMD_INLINE typename std::enable_if< +ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t< detail::is_dword_type::value && detail::is_dword_type::value && detail::is_dword_type::value && detail::is_dword_type::value, - simd>::type + simd> esimd_dp4a(simd src0, simd src1, simd src2, int flag = GENX_NOSAT) { simd Src0 = src0; diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp index ad2dd7e9fbe97..f1662cf6c2713 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp @@ -24,8 +24,8 @@ namespace INTEL { namespace gpu { template > + typename = sycl::detail::enable_if_t< + (ElemsPerAddr == 1 || ElemsPerAddr == 2 || ElemsPerAddr == 4)>> constexpr unsigned int ElemsPerAddrEncoding() { // encoding requires log2 of ElemsPerAddr if constexpr (ElemsPerAddr == 1) @@ -83,12 +83,11 @@ constexpr unsigned int ElemsPerAddrEncoding() { /// flat-address gather template -ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<((n == 8 || n == 16 || n == 32) && - (ElemsPerAddr == 1 || ElemsPerAddr == 2 || - ElemsPerAddr == 4)), - simd>::type - gather(T *p, simd offsets, simd pred = 1) { +ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< + ((n == 8 || n == 16 || n == 32) && + (ElemsPerAddr == 1 || ElemsPerAddr == 2 || ElemsPerAddr == 4)), + simd> +gather(T *p, simd offsets, simd pred = 1) { simd offsets_i = convert(offsets); simd addrs(reinterpret_cast(p)); @@ -125,13 +124,12 @@ ESIMD_INLINE ESIMD_NODEBUG /// flat-address scatter template -ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<((n == 8 || n == 16 || n == 32) && - (ElemsPerAddr == 1 || ElemsPerAddr == 2 || - ElemsPerAddr == 4)), - void>::type - scatter(T *p, simd vals, simd offsets, - simd pred = 1) { +ESIMD_INLINE ESIMD_NODEBUG typename sycl::detail::enable_if_t< + ((n == 8 || n == 16 || n == 32) && + (ElemsPerAddr == 1 || ElemsPerAddr == 2 || ElemsPerAddr == 4)), + void> +scatter(T *p, simd vals, simd offsets, + simd pred = 1) { simd offsets_i = convert(offsets); simd addrs(reinterpret_cast(p)); addrs = addrs + offsets_i; @@ -263,9 +261,10 @@ ESIMD_INLINE ESIMD_NODEBUG void block_store(AccessorTy acc, uint32_t offset, template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(sizeof(T) <= 4) && (N == 1 || N == 8 || N == 16) && - !std::is_pointer::value, - simd>::type + typename sycl::detail::enable_if_t<(sizeof(T) <= 4) && + (N == 1 || N == 8 || N == 16) && + !std::is_pointer::value, + simd> gather(AccessorTy acc, simd offsets, uint32_t glob_offset = 0) { @@ -282,8 +281,9 @@ ESIMD_INLINE ESIMD_NODEBUG if constexpr (sizeof(T) < 4) { static_assert(std::is_integral::value, "only integral 1- & 2-byte types are supported"); - using PromoT = typename std::conditional::value, int32_t, - uint32_t>::type; + using PromoT = + typename sycl::detail::conditional_t::value, int32_t, + uint32_t>; #if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__) const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc); const simd promo_vals = @@ -328,9 +328,10 @@ ESIMD_INLINE ESIMD_NODEBUG template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(sizeof(T) <= 4) && (N == 1 || N == 8 || N == 16) && - !std::is_pointer::value, - void>::type + typename sycl::detail::enable_if_t<(sizeof(T) <= 4) && + (N == 1 || N == 8 || N == 16) && + !std::is_pointer::value, + void> scatter(AccessorTy acc, simd vals, simd offsets, uint32_t glob_offset = 0, simd pred = 1) { @@ -347,8 +348,9 @@ ESIMD_INLINE ESIMD_NODEBUG if constexpr (sizeof(T) < 4) { static_assert(std::is_integral::value, "only integral 1- & 2-byte types are supported"); - using PromoT = typename std::conditional::value, int32_t, - uint32_t>::type; + using PromoT = + typename sycl::detail::conditional_t::value, int32_t, + uint32_t>; const simd promo_vals = sycl::INTEL::gpu::convert(vals); #if defined(__SYCL_DEVICE_ONLY__) && defined(__SYCL_EXPLICIT_SIMD__) const auto surf_ind = AccessorPrivateProxy::getNativeImageObj(acc); @@ -395,8 +397,8 @@ ESIMD_INLINE ESIMD_NODEBUG void scalar_store(AccessorTy acc, uint32_t offset, template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(n == 16 || n == 32) && (sizeof(T) == 4), - simd>::type + typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4), + simd> gather4(T *p, simd offsets, simd pred = 1) { simd offsets_i = convert(offsets); @@ -409,8 +411,8 @@ ESIMD_INLINE ESIMD_NODEBUG template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(n == 16 || n == 32) && (sizeof(T) == 4), - void>::type + typename sycl::detail::enable_if_t<(n == 16 || n == 32) && (sizeof(T) == 4), + void> scatter4(T *p, simd vals, simd offsets, simd pred = 1) { simd offsets_i = convert(offsets); @@ -530,7 +532,7 @@ constexpr bool check_atomic() { template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> flat_atomic(T *p, simd offset, simd pred) { simd vAddr(reinterpret_cast(p)); simd offset_i1 = convert(offset); @@ -542,7 +544,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> flat_atomic(T *p, simd offset, simd src0, simd pred) { simd vAddr(reinterpret_cast(p)); @@ -556,7 +558,7 @@ ESIMD_NODEBUG ESIMD_INLINE template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> flat_atomic(T *p, simd offset, simd src0, simd src1, simd pred) { simd vAddr(reinterpret_cast(p)); @@ -594,7 +596,7 @@ SYCL_EXTERNAL void slm_init(uint32_t size); /// only allow simd-16 and simd-32 template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(n == 16 || n == 32), simd>::type + typename sycl::detail::enable_if_t<(n == 16 || n == 32), simd> slm_load(simd offsets, simd pred = 1) { return __esimd_slm_read(offsets.data(), pred.data()); } @@ -602,7 +604,7 @@ ESIMD_INLINE ESIMD_NODEBUG /// SLM scatter template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(n == 16 || n == 32), void>::type + typename sycl::detail::enable_if_t<(n == 16 || n == 32), void> slm_store(simd vals, simd offsets, simd pred = 1) { __esimd_slm_write(offsets.data(), vals.data(), pred.data()); @@ -612,16 +614,17 @@ ESIMD_INLINE ESIMD_NODEBUG /// only allow simd-8, simd-16 and simd-32 template ESIMD_INLINE ESIMD_NODEBUG - typename std::enable_if<(n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), - simd>::type + typename sycl::detail::enable_if_t<(n == 8 || n == 16 || n == 32) && + (sizeof(T) == 4), + simd> slm_load4(simd offsets, simd pred = 1) { return __esimd_slm_read4(offsets.data(), pred.data()); } /// SLM scatter4 template -typename std::enable_if<(n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), - void>::type +typename sycl::detail::enable_if_t< + (n == 8 || n == 16 || n == 32) && (sizeof(T) == 4), void> slm_store4(simd vals, simd offsets, simd pred = 1) { __esimd_slm_write4(offsets.data(), vals.data(), pred.data()); @@ -662,7 +665,7 @@ ESIMD_INLINE ESIMD_NODEBUG void slm_block_store(uint32_t offset, /// SLM atomic, zero source operand: inc and dec template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> slm_atomic(simd offsets, simd pred) { return __esimd_slm_atomic0(offsets.data(), pred.data()); } @@ -670,7 +673,7 @@ ESIMD_NODEBUG ESIMD_INLINE /// SLM atomic, one source operand, add/sub/min/max etc template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> slm_atomic(simd offsets, simd src0, simd pred) { return __esimd_slm_atomic1(offsets.data(), src0.data(), @@ -680,7 +683,7 @@ ESIMD_NODEBUG ESIMD_INLINE /// SLM atomic, two source operands template ESIMD_NODEBUG ESIMD_INLINE - typename std::enable_if(), simd>::type + typename sycl::detail::enable_if_t(), simd> slm_atomic(simd offsets, simd src0, simd src1, simd pred) { return __esimd_slm_atomic2(offsets.data(), src0.data(), src1.data(), diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_view.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_view.hpp index e7043e3063089..16d79190fd228 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd_view.hpp +++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_view.hpp @@ -143,7 +143,7 @@ template class simd_view { // @return the representing region object. // template > + typename = sycl::detail::enable_if_t> auto select(uint16_t Offset = 0) { using TopRegionTy = region1d_t; using NewRegionTy = std::pair; @@ -171,7 +171,8 @@ template class simd_view { // @return the representing region object. // template > + typename T = simd_view, + typename = sycl::detail::enable_if_t> auto select(uint16_t OffsetY = 0, uint16_t OffsetX = 0) { using TopRegionTy = region2d_t; @@ -263,19 +264,22 @@ template class simd_view { } // Reference a row from a 2D region. This returns a 1D region. - template > + template > auto row(int i) { return select<1, 0, getSizeX(), 1>(i, 0).template format(); } // Reference a column from a 2D region. This returns a 2D region. - template > + template > auto column(int i) { return select(0, i); } // Read a single element from a 1D region, by value only. - template > + template > element_type operator[](int i) const { return read()[i]; } @@ -336,8 +340,9 @@ template class simd_view { // // @return 1 if any element is set, 0 otherwise - template ::value, T2>> + template < + typename T1 = element_type, typename T2 = BaseTy, + typename = sycl::detail::enable_if_t::value, T2>> uint16_t any() { return read().any(); } @@ -346,8 +351,9 @@ template class simd_view { // // @return 1 if all elements are set, 0 otherwise - template ::value, T2>> + template < + typename T1 = element_type, typename T2 = BaseTy, + typename = sycl::detail::enable_if_t::value, T2>> uint16_t all() { return read().all(); } diff --git a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp index d2f70ad20f60e..1fc758f4b4f8b 100644 --- a/sycl/include/CL/sycl/INTEL/fpga_utils.hpp +++ b/sycl/include/CL/sycl/INTEL/fpga_utils.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -25,8 +26,8 @@ template