From 3dbab10d62c8666a3cf8627f875a6ec3d701c300 Mon Sep 17 00:00:00 2001
From: Artem Gindinson <artem.gindinson@intel.com>
Date: Fri, 13 Nov 2020 15:49:53 +0300
Subject: [PATCH 1/3] [SYCL][NFC] Employ the is_same_v helper in SYCL library

Similarly to [intel/llvm@b469f03], refactor the usages of the
`is_same` type trait to reduce code verbosity. This particular
trait was chosen because of its popularity, and will probably
be the last in this refactoring series.

Signed-off-by: Artem Gindinson <artem.gindinson@intel.com>
---
 .../sycl/INTEL/esimd/detail/esimd_intrin.hpp  |   2 +-
 .../sycl/INTEL/esimd/detail/esimd_types.hpp   |   7 +-
 .../CL/sycl/INTEL/esimd/detail/esimd_util.hpp |  54 ++++----
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp    |   2 +-
 .../CL/sycl/INTEL/esimd/esimd_math.hpp        |  16 +--
 .../CL/sycl/ONEAPI/accessor_property_list.hpp |   2 +-
 sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp    |  12 +-
 .../CL/sycl/ONEAPI/group_algorithm.hpp        | 128 +++++++++---------
 sycl/include/CL/sycl/ONEAPI/reduction.hpp     | 108 +++++++--------
 sycl/include/CL/sycl/ONEAPI/sub_group.hpp     |   4 +-
 sycl/include/CL/sycl/accessor.hpp             |  35 +++--
 sycl/include/CL/sycl/atomic.hpp               |  15 +-
 sycl/include/CL/sycl/backend/level_zero.hpp   |  12 +-
 sycl/include/CL/sycl/backend/opencl.hpp       |  14 +-
 sycl/include/CL/sycl/buffer.hpp               |   2 +-
 sycl/include/CL/sycl/detail/accessor_impl.hpp |  42 +++---
 sycl/include/CL/sycl/detail/cg_types.hpp      |  12 +-
 .../CL/sycl/detail/generic_type_traits.hpp    |  10 +-
 .../CL/sycl/detail/stl_type_traits.hpp        |   3 +
 .../CL/sycl/detail/sycl_mem_obj_allocator.hpp |   4 +-
 .../include/CL/sycl/detail/sycl_mem_obj_t.hpp |   4 +-
 sycl/include/CL/sycl/detail/type_list.hpp     |   7 +-
 sycl/include/CL/sycl/detail/type_traits.hpp   |   6 +-
 sycl/include/CL/sycl/handler.hpp              |  21 +--
 sycl/include/CL/sycl/multi_ptr.hpp            |  33 +++--
 sycl/include/CL/sycl/stream.hpp               |  10 +-
 sycl/include/CL/sycl/types.hpp                | 113 ++++++++--------
 .../source/detail/scheduler/graph_builder.cpp |   2 +-
 sycl/source/detail/scheduler/scheduler.hpp    |   2 +-
 29 files changed, 338 insertions(+), 344 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
index 34bbc905ceb44..111fd37c3e024 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
@@ -179,7 +179,7 @@ readRegion(const vector_type_t<BT, BN> &Base, std::pair<T, U> Region) {
     return readRegion<BT1, BN1>(Base1, Region.first);
   else {
     static_assert(T::Is_2D);
-    static_assert(std::is_same<ElemTy, BT1>::value);
+    static_assert(sycl::detail::is_same_v<ElemTy, BT1>);
     // To read a 2D region, we need the parent region
     // Read full rows with non-trivial vertical and horizontal stride = 1.
     constexpr int M = T::Size_y * PaTy::Size_x;
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 1655c78c96352..aae6535aff487 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp
@@ -209,7 +209,7 @@ template <typename U> constexpr bool is_type() { return false; }
 template <typename U, typename T, typename... Ts> constexpr bool is_type() {
   using UU = typename detail::remove_const_t<U>;
   using TT = typename detail::remove_const_t<T>;
-  return std::is_same<UU, TT>::value || is_type<UU, Ts...>();
+  return sycl::detail::is_same_v<UU, TT> || is_type<UU, Ts...>();
 }
 
 // calculates the number of elements in "To" type
@@ -228,12 +228,13 @@ struct bitcast_helper {
 template <typename ToEltTy, typename FromEltTy, int FromN,
           typename = csd::enable_if_t<is_vectorizable<ToEltTy>::value>>
 ESIMD_INLINE typename detail::conditional_t<
-    std::is_same<FromEltTy, ToEltTy>::value, vector_type_t<FromEltTy, FromN>,
+    sycl::detail::is_same_v<FromEltTy, ToEltTy>,
+    vector_type_t<FromEltTy, FromN>,
     vector_type_t<ToEltTy,
                   bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems()>>
 bitcast(vector_type_t<FromEltTy, FromN> Val) {
   // Noop.
-  if constexpr (std::is_same<FromEltTy, ToEltTy>::value)
+  if constexpr (sycl::detail::is_same_v<FromEltTy, ToEltTy>)
     return Val;
 
   // Bitcast
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 6d0be7f5521ff..5bd1c5db2d62b 100755
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp
@@ -101,9 +101,10 @@ template <typename T>
 struct is_dword_type
     : std::integral_constant<
           bool,
-          std::is_same<int, typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<unsigned int,
-                           typename sycl::detail::remove_const_t<T>>::value> {};
+          sycl::detail::is_same_v<int,
+                                  typename sycl::detail::remove_const_t<T>> ||
+              sycl::detail::is_same_v<
+                  unsigned int, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_dword_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -119,10 +120,10 @@ template <typename T>
 struct is_word_type
     : std::integral_constant<
           bool,
-          std::is_same<short,
-                       typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<unsigned short,
-                           typename sycl::detail::remove_const_t<T>>::value> {};
+          sycl::detail::is_same_v<short,
+                                  typename sycl::detail::remove_const_t<T>> ||
+              sycl::detail::is_same_v<
+                  unsigned short, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_word_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -137,9 +138,10 @@ template <typename T>
 struct is_byte_type
     : std::integral_constant<
           bool,
-          std::is_same<char, typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<unsigned char,
-                           typename sycl::detail::remove_const_t<T>>::value> {};
+          sycl::detail::is_same_v<char,
+                                  typename sycl::detail::remove_const_t<T>> ||
+              sycl::detail::is_same_v<
+                  unsigned char, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_byte_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -153,36 +155,34 @@ template <typename T, int N> struct is_byte_type<sycl::INTEL::gpu::simd<T, N>> {
 template <typename T>
 struct is_fp_type
     : std::integral_constant<
-          bool, std::is_same<float,
-                             typename sycl::detail::remove_const_t<T>>::value> {
-};
+          bool, sycl::detail::is_same_v<
+                    float, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_df_type
     : std::integral_constant<
-          bool, std::is_same<double,
-                             typename sycl::detail::remove_const_t<T>>::value> {
-};
+          bool, sycl::detail::is_same_v<
+                    double, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_fp_or_dword_type
     : std::integral_constant<
           bool,
-          std::is_same<float,
-                       typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<int,
-                           typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<unsigned int,
-                           typename sycl::detail::remove_const_t<T>>::value> {};
+          sycl::detail::is_same_v<float,
+                                  typename sycl::detail::remove_const_t<T>> ||
+              sycl::detail::is_same_v<
+                  int, typename sycl::detail::remove_const_t<T>> ||
+              sycl::detail::is_same_v<
+                  unsigned int, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_qword_type
     : std::integral_constant<
-          bool,
-          std::is_same<long long,
-                       typename sycl::detail::remove_const_t<T>>::value ||
-              std::is_same<unsigned long long,
-                           typename sycl::detail::remove_const_t<T>>::value> {};
+          bool, sycl::detail::is_same_v<
+                    long long, typename sycl::detail::remove_const_t<T>> ||
+                    sycl::detail::is_same_v<
+                        unsigned long long,
+                        typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_qword_type<sycl::INTEL::gpu::vector_type<T, N>> {
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index 6a9251cbc33b6..b3437c7c9dcf6 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -403,7 +403,7 @@ template <typename Ty, int N> class simd {
         Base1 = __esimd_wrregion<ElemTy, BN1, M,
                                  /*VS*/ 0, M, Stride>(Base1, Val, Offset);
       } else {
-        static_assert(std::is_same<ElemTy, BT>::value);
+        static_assert(sycl::detail::is_same_v<ElemTy, BT>);
         // Read columns with non-trivial horizontal stride.
         constexpr int M = TR::length;
         constexpr int VS = PaTy::Size_x * TR::Stride_y;
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
index bda406c12beb9..dc1bf201ac3d9 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
@@ -73,8 +73,8 @@ __esimd_abs_common_internal(T1 src0, int flag = GENX_NOSAT) {
 
 template <typename T0, typename T1, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
-    !std::is_same<typename sycl::detail::remove_const_t<T0>,
-                  typename sycl::detail::remove_const_t<T1>>::value,
+    !sycl::detail::is_same_v<typename sycl::detail::remove_const_t<T0>,
+                             typename sycl::detail::remove_const_t<T1>>,
     simd<T0, SZ>>
 esimd_abs(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
   return detail::__esimd_abs_common_internal<T0, T1, SZ>(src0, flag);
@@ -82,8 +82,8 @@ esimd_abs(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
 
 template <typename T0, typename T1>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
-    !std::is_same<typename sycl::detail::remove_const_t<T0>,
-                  typename sycl::detail::remove_const_t<T1>>::value &&
+    !sycl::detail::is_same_v<typename sycl::detail::remove_const_t<T0>,
+                             typename sycl::detail::remove_const_t<T1>> &&
         detail::is_esimd_scalar<T0>::value &&
         detail::is_esimd_scalar<T1>::value,
     typename sycl::detail::remove_const_t<T0>>
@@ -1905,12 +1905,12 @@ ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
 
 template <typename T0, typename T1, int SZ, typename BinaryOperation>
 ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
-  if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
-                             std::plus<>>::value) {
+  if constexpr (sycl::detail::is_same_v<remove_cvref_t<BinaryOperation>,
+                                        std::plus<>>) {
     T0 retv = detail::esimd_sum<T0>(v);
     return retv;
-  } else if constexpr (std::is_same<remove_cvref_t<BinaryOperation>,
-                                    std::multiplies<>>::value) {
+  } else if constexpr (sycl::detail::is_same_v<remove_cvref_t<BinaryOperation>,
+                                               std::multiplies<>>) {
     T0 retv = detail::esimd_prod<T0>(v);
     return retv;
   }
diff --git a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
index b83e4ed43d316..c0a72f6b9a440 100644
--- a/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/accessor_property_list.hpp
@@ -86,7 +86,7 @@ class accessor_property_list : protected sycl::detail::PropertyListBase {
             auto... Args>
   struct ContainsPropertyInstance
       : detail::conditional_t<
-            !std::is_same_v<typename ContainerT::Head, void> &&
+            !detail::is_same_v<typename ContainerT::Head, void> &&
                 AreSameTemplate<PropT<Args...>,
                                 typename ContainerT::Head>::value,
             std::true_type,
diff --git a/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp b/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
index c98c4f55857b8..e3a5f504c46d1 100644
--- a/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/atomic_ref.hpp
@@ -126,14 +126,14 @@ class atomic_ref_base {
   static_assert(
       detail::IsValidAtomicType<T>::value,
       "Invalid atomic type.  Valid types are arithmetic and pointer types");
-  static_assert(!std::is_same<T, bool>::value,
+  static_assert(!detail::is_same_v<T, bool>,
                 "ONEAPI::atomic_ref does not support bool type");
-  static_assert(!(std::is_same<T, char>::value ||
-                  std::is_same<T, signed char>::value ||
-                  std::is_same<T, unsigned char>::value),
+  static_assert(!(detail::is_same_v<T, char> ||
+                  detail::is_same_v<T, signed char> ||
+                  detail::is_same_v<T, unsigned char>),
                 "ONEAPI::atomic_ref does not support char type");
-  static_assert(!(std::is_same<T, short>::value ||
-                  std::is_same<T, unsigned short>::value),
+  static_assert(!(detail::is_same_v<T, short> ||
+                  detail::is_same_v<T, unsigned short>),
                 "ONEAPI::atomic_ref does not support short type");
   static_assert(detail::IsValidAtomicAddressSpace<AddressSpace>::value,
                 "Invalid atomic address_space.  Valid address spaces are: "
diff --git a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
index 3aaab50b236aa..1b0c7ad38e7a7 100644
--- a/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/group_algorithm.hpp
@@ -421,9 +421,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 reduce(Group, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(x, x)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(x, x)), float>::value),
+      detail::is_same_v<decltype(binary_op(x, x)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(x, x)), float>),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<T, __spv::GroupOperation::Reduce,
@@ -443,10 +443,10 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 reduce(Group g, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(x[0], x[0])),
-                   typename T::element_type>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
+      detail::is_same_v<decltype(binary_op(x[0], x[0])),
+                        typename T::element_type> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(x[0], x[0])), float>),
       "Result type of binary_op must match reduction accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -482,9 +482,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 reduce(Group g, V x, T init, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init, x)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init, x)), float>::value),
+      detail::is_same_v<decltype(binary_op(init, x)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init, x)), float>),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return binary_op(init, reduce(g, x, binary_op));
@@ -505,10 +505,10 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 reduce(Group g, V x, T init, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init[0], x[0])),
-                   typename T::element_type>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
+      detail::is_same_v<decltype(binary_op(init[0], x[0])),
+                        typename T::element_type> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init[0], x[0])), float>),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T result = init;
@@ -551,9 +551,9 @@ reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
   using T = typename detail::remove_pointer<Ptr>::type;
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(*first, *first)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(*first, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(*first, *first)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(*first, *first)), float>),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   typename Ptr::element_type partial =
@@ -582,9 +582,9 @@ detail::enable_if_t<
 reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init, *first)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(init, *first)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init, *first)), float>),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T partial = sycl::detail::identity<T, BinaryOperation>::value;
@@ -608,9 +608,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 exclusive_scan(Group, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
-  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value ||
-                    (std::is_same<T, half>::value &&
-                     std::is_same<decltype(binary_op(x, x)), float>::value),
+  static_assert(detail::is_same_v<decltype(binary_op(x, x)), T> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(x, x)), float>),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<T, __spv::GroupOperation::ExclusiveScan,
@@ -629,12 +629,11 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 exclusive_scan(Group g, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
-  static_assert(
-      std::is_same<decltype(binary_op(x[0], x[0])),
-                   typename T::element_type>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
-      "Result type of binary_op must match scan accumulation type.");
+  static_assert(detail::is_same_v<decltype(binary_op(x[0], x[0])),
+                                  typename T::element_type> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(x[0], x[0])), float>),
+                "Result type of binary_op must match scan accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
     result[s] = exclusive_scan(g, x[s], binary_op);
@@ -652,10 +651,10 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 exclusive_scan(Group g, V x, T init, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init[0], x[0])),
-                   typename T::element_type>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
+      detail::is_same_v<decltype(binary_op(init[0], x[0])),
+                        typename T::element_type> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init[0], x[0])), float>),
       "Result type of binary_op must match scan accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -673,9 +672,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 exclusive_scan(Group g, V x, T init, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
-  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value ||
-                    (std::is_same<T, half>::value &&
-                     std::is_same<decltype(binary_op(init, x)), float>::value),
+  static_assert(detail::is_same_v<decltype(binary_op(init, x)), T> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(init, x)), float>),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   typename Group::linear_id_type local_linear_id =
@@ -711,9 +710,9 @@ exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
                BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(*first, *first)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(*first, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(*first, *first)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(*first, *first)), float>),
       "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
@@ -761,10 +760,10 @@ exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(*first, *first)),
-                   typename OutPtr::element_type>::value ||
-          (std::is_same<typename OutPtr::element_type, half>::value &&
-           std::is_same<decltype(binary_op(*first, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(*first, *first)),
+                        typename OutPtr::element_type> ||
+          (detail::is_same_v<typename OutPtr::element_type, half> &&
+           detail::is_same_v<decltype(binary_op(*first, *first)), float>),
       "Result type of binary_op must match scan accumulation type.");
   return exclusive_scan(g, first, last, result,
                         sycl::detail::identity<typename OutPtr::element_type,
@@ -779,12 +778,11 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 inclusive_scan(Group g, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
-  static_assert(
-      std::is_same<decltype(binary_op(x[0], x[0])),
-                   typename T::element_type>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(x[0], x[0])), float>::value),
-      "Result type of binary_op must match scan accumulation type.");
+  static_assert(detail::is_same_v<decltype(binary_op(x[0], x[0])),
+                                  typename T::element_type> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(x[0], x[0])), float>),
+                "Result type of binary_op must match scan accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
     result[s] = inclusive_scan(g, x[s], binary_op);
@@ -799,9 +797,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 inclusive_scan(Group, T x, BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
-  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value ||
-                    (std::is_same<T, half>::value &&
-                     std::is_same<decltype(binary_op(x, x)), float>::value),
+  static_assert(detail::is_same_v<decltype(binary_op(x, x)), T> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(x, x)), float>),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<T, __spv::GroupOperation::InclusiveScan,
@@ -822,9 +820,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
                     T>
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
   // FIXME: Do not special-case for half precision
-  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value ||
-                    (std::is_same<T, half>::value &&
-                     std::is_same<decltype(binary_op(init, x)), float>::value),
+  static_assert(detail::is_same_v<decltype(binary_op(init, x)), T> ||
+                    (detail::is_same_v<T, half> &&
+                     detail::is_same_v<decltype(binary_op(init, x)), float>),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   if (sycl::detail::get_local_linear_id(g) == 0) {
@@ -848,9 +846,9 @@ detail::enable_if_t<(detail::is_generic_group<Group>::value &&
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init[0], x[0])), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
+      detail::is_same_v<decltype(binary_op(init[0], x[0])), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init[0], x[0])), float>),
       "Result type of binary_op must match scan accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -875,9 +873,9 @@ inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                BinaryOperation binary_op, T init) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init, *first)), T>::value ||
-          (std::is_same<T, half>::value &&
-           std::is_same<decltype(binary_op(init, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(init, *first)), T> ||
+          (detail::is_same_v<T, half> &&
+           detail::is_same_v<decltype(binary_op(init, *first)), float>),
       "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
@@ -924,10 +922,10 @@ inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                BinaryOperation binary_op) {
   // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(*first, *first)),
-                   typename OutPtr::element_type>::value ||
-          (std::is_same<typename OutPtr::element_type, half>::value &&
-           std::is_same<decltype(binary_op(*first, *first)), float>::value),
+      detail::is_same_v<decltype(binary_op(*first, *first)),
+                        typename OutPtr::element_type> ||
+          (detail::is_same_v<typename OutPtr::element_type, half> &&
+           detail::is_same_v<decltype(binary_op(*first, *first)), float>),
       "Result type of binary_op must match scan accumulation type.");
   return inclusive_scan(g, first, last, result, binary_op,
                         sycl::detail::identity<typename OutPtr::element_type,
diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
index a11d3d89fb513..ae8e87459c58d 100644
--- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
@@ -36,38 +36,38 @@ using cl::sycl::detail::remove_AS;
 
 template <typename T, class BinaryOperation>
 using IsReduPlus = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::plus<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::plus<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::plus<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::plus<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduMultiplies = detail::bool_constant<
-    std::is_same<BinaryOperation, std::multiplies<T>>::value ||
-    std::is_same<BinaryOperation, std::multiplies<void>>::value>;
+    detail::is_same_v<BinaryOperation, std::multiplies<T>> ||
+    detail::is_same_v<BinaryOperation, std::multiplies<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduMinimum = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::minimum<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::minimum<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::minimum<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::minimum<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduMaximum = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::maximum<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::maximum<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::maximum<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::maximum<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduBitOR = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::bit_or<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::bit_or<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_or<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_or<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduBitXOR = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::bit_xor<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::bit_xor<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_xor<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_xor<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduBitAND = detail::bool_constant<
-    std::is_same<BinaryOperation, ONEAPI::bit_and<T>>::value ||
-    std::is_same<BinaryOperation, ONEAPI::bit_and<void>>::value>;
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_and<T>> ||
+    detail::is_same_v<BinaryOperation, ONEAPI::bit_and<void>>>;
 
 template <typename T, class BinaryOperation>
 using IsReduOptForFastAtomicFetch =
@@ -81,13 +81,12 @@ using IsReduOptForFastAtomicFetch =
                            IsReduBitAND<T, BinaryOperation>::value)>;
 
 template <typename T, class BinaryOperation>
-using IsReduOptForFastReduce = detail::bool_constant<
-    (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
-     std::is_same<T, half>::value || std::is_same<T, float>::value ||
-     std::is_same<T, double>::value) &&
-    (IsReduPlus<T, BinaryOperation>::value ||
-     IsReduMinimum<T, BinaryOperation>::value ||
-     IsReduMaximum<T, BinaryOperation>::value)>;
+using IsReduOptForFastReduce = detail::bool_constant<(
+    is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
+    detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
+    detail::is_same_v<T, double>) && (IsReduPlus<T, BinaryOperation>::value ||
+                                     IsReduMinimum<T, BinaryOperation>::value ||
+                                     IsReduMaximum<T, BinaryOperation>::value)>;
 
 // Identity = 0
 template <typename T, class BinaryOperation>
@@ -97,18 +96,16 @@ using IsZeroIdentityOp = bool_constant<
      (IsReduPlus<T, BinaryOperation>::value ||
       IsReduBitOR<T, BinaryOperation>::value ||
       IsReduBitXOR<T, BinaryOperation>::value)) ||
-    ((std::is_same<T, half>::value || std::is_same<T, float>::value ||
-      std::is_same<T, double>::value) &&
-     IsReduPlus<T, BinaryOperation>::value)>;
+    ((detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
+      detail::is_same_v<T, double>) && IsReduPlus<T, BinaryOperation>::value)>;
 
 // Identity = 1
 template <typename T, class BinaryOperation>
-using IsOneIdentityOp = bool_constant<
-    (is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
-     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
-     std::is_same<T, half>::value || std::is_same<T, float>::value ||
-     std::is_same<T, double>::value) &&
-    IsReduMultiplies<T, BinaryOperation>::value>;
+using IsOneIdentityOp = bool_constant<(
+    is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
+    is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
+    detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
+    detail::is_same_v<T, double>) && IsReduMultiplies<T, BinaryOperation>::value>;
 
 // Identity = ~0
 template <typename T, class BinaryOperation>
@@ -119,21 +116,19 @@ using IsOnesIdentityOp = bool_constant<
 
 // Identity = <max possible value>
 template <typename T, class BinaryOperation>
-using IsMinimumIdentityOp = bool_constant<
-    (is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
-     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
-     std::is_same<T, half>::value || std::is_same<T, float>::value ||
-     std::is_same<T, double>::value) &&
-    IsReduMinimum<T, BinaryOperation>::value>;
+using IsMinimumIdentityOp = bool_constant<(
+    is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
+    is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
+    detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
+    detail::is_same_v<T, double>)&&IsReduMinimum<T, BinaryOperation>::value>;
 
 // Identity = <min possible value>
 template <typename T, class BinaryOperation>
-using IsMaximumIdentityOp = bool_constant<
-    (is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
-     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
-     std::is_same<T, half>::value || std::is_same<T, float>::value ||
-     std::is_same<T, double>::value) &&
-    IsReduMaximum<T, BinaryOperation>::value>;
+using IsMaximumIdentityOp = bool_constant<(
+    is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
+    is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
+    detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
+    detail::is_same_v<T, double>) && IsReduMaximum<T, BinaryOperation>::value>;
 
 template <typename T, class BinaryOperation>
 using IsKnownIdentityOp =
@@ -232,8 +227,7 @@ class reducer<T, BinaryOperation,
   }
 
   template <typename _T = T>
-  enable_if_t<std::is_same<_T, T>::value &&
-                  IsReduPlus<T, BinaryOperation>::value,
+  enable_if_t<detail::is_same_v<_T, T> && IsReduPlus<T, BinaryOperation>::value,
               reducer &>
   operator+=(const _T &Partial) {
     combine(Partial);
@@ -241,7 +235,7 @@ class reducer<T, BinaryOperation,
   }
 
   template <typename _T = T>
-  enable_if_t<std::is_same<_T, T>::value &&
+  enable_if_t<detail::is_same_v<_T, T> &&
                   IsReduMultiplies<T, BinaryOperation>::value,
               reducer &>
   operator*=(const _T &Partial) {
@@ -250,7 +244,7 @@ class reducer<T, BinaryOperation,
   }
 
   template <typename _T = T>
-  enable_if_t<std::is_same<_T, T>::value &&
+  enable_if_t<detail::is_same_v<_T, T> &&
                   IsReduBitOR<T, BinaryOperation>::value,
               reducer &>
   operator|=(const _T &Partial) {
@@ -259,7 +253,7 @@ class reducer<T, BinaryOperation,
   }
 
   template <typename _T = T>
-  enable_if_t<std::is_same<_T, T>::value &&
+  enable_if_t<detail::is_same_v<_T, T> &&
                   IsReduBitXOR<T, BinaryOperation>::value,
               reducer &>
   operator^=(const _T &Partial) {
@@ -268,7 +262,7 @@ class reducer<T, BinaryOperation,
   }
 
   template <typename _T = T>
-  enable_if_t<std::is_same<_T, T>::value &&
+  enable_if_t<detail::is_same_v<_T, T> &&
                   IsReduBitAND<T, BinaryOperation>::value,
               reducer &>
   operator&=(const _T &Partial) {
@@ -278,7 +272,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic ADD operation: *ReduVarPtr += MValue;
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduPlus<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -288,7 +282,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic BITWISE OR operation: *ReduVarPtr |= MValue;
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduBitOR<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -298,7 +292,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic BITWISE XOR operation: *ReduVarPtr ^= MValue;
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduBitXOR<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -308,7 +302,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic BITWISE AND operation: *ReduVarPtr &= MValue;
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduBitAND<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -318,7 +312,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic MIN operation: *ReduVarPtr = ONEAPI::minimum(*ReduVarPtr, MValue);
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduMinimum<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -328,7 +322,7 @@ class reducer<T, BinaryOperation,
 
   /// Atomic MAX operation: *ReduVarPtr = ONEAPI::maximum(*ReduVarPtr, MValue);
   template <typename _T = T, class _BinaryOperation = BinaryOperation>
-  enable_if_t<std::is_same<typename remove_AS<_T>::type, T>::value &&
+  enable_if_t<detail::is_same_v<typename remove_AS<_T>::type, T> &&
               (is_geninteger32bit<T>::value || is_geninteger64bit<T>::value) &&
               IsReduMaximum<T, _BinaryOperation>::value>
   atomic_combine(_T *ReduVarPtr) const {
@@ -554,14 +548,14 @@ class reduction_impl {
   }
 
   template <typename AccT>
-  enable_if_t<std::is_same<AccT, rw_accessor_type>::value ||
-                  std::is_same<AccT, accessor_type>::value,
+  enable_if_t<detail::is_same_v<AccT, rw_accessor_type> ||
+                  detail::is_same_v<AccT, accessor_type>,
               result_type *> static inline getOutPointer(const AccT &OutAcc) {
     return OutAcc.get_pointer().get();
   }
 
   template <typename ResT>
-  enable_if_t<std::is_same<ResT, result_type>::value,
+  enable_if_t<detail::is_same_v<ResT, result_type>,
               result_type *> static inline getOutPointer(ResT *OutPtr) {
     return OutPtr;
   }
diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
index 981ec3acf779a..e9183369077b9 100644
--- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp
@@ -37,12 +37,12 @@ using SelectBlockT = select_cl_scalar_integral_unsigned_t<T>;
 
 template <typename T, access::address_space Space>
 using AcceptableForGlobalLoadStore =
-    bool_constant<!std::is_same<void, SelectBlockT<T>>::value &&
+    bool_constant<!detail::is_same_v<void, SelectBlockT<T>> &&
                   Space == access::address_space::global_space>;
 
 template <typename T, access::address_space Space>
 using AcceptableForLocalLoadStore =
-    bool_constant<!std::is_same<void, SelectBlockT<T>>::value &&
+    bool_constant<!detail::is_same_v<void, SelectBlockT<T>> &&
                   Space == access::address_space::local_space>;
 
 #ifdef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 7f8e62409b661..7ed32b07a64f0 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -402,12 +402,12 @@ class image_accessor
   constexpr static bool IsImageAccessAnyRead =
       (IsImageAccessReadOnly || AccessMode == access::mode::read_write);
 
-  static_assert(std::is_same<DataT, cl_int4>::value ||
-                      std::is_same<DataT, cl_uint4>::value ||
-                      std::is_same<DataT, cl_float4>::value ||
-                      std::is_same<DataT, cl_half4>::value,
-                  "The data type of an image accessor must be only cl_int4, "
-                  "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
+  static_assert(detail::is_same_v<DataT, cl_int4> ||
+                    detail::is_same_v<DataT, cl_uint4> ||
+                    detail::is_same_v<DataT, cl_float4> ||
+                    detail::is_same_v<DataT, cl_half4>,
+                "The data type of an image accessor must be only cl_int4, "
+                "cl_uint4, cl_float4 or cl_half4 from SYCL namespace");
 
   static_assert(IsImageAcc || IsHostImageAcc || IsImageArrayAcc,
                 "Expected image type");
@@ -672,7 +672,7 @@ class __image_array_slice__ {
   getAdjustedCoords(const CoordT &Coords) const {
     CoordElemType LastCoord = 0;
 
-    if (std::is_same<float, CoordElemType>::value) {
+    if (detail::is_same_v<float, CoordElemType>) {
       sycl::vec<int, Dimensions + 1> Size = MBaseAcc.getRangeInternal();
       LastCoord =
           MIdx / static_cast<float>(Size.template swizzle<Dimensions>());
@@ -819,7 +819,7 @@ class accessor :
   }
 
   template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
-    return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
+    return detail::is_same_v<T, DataT> && (Dims > 0) && (Dims == Dimensions);
   }
 
   static access::mode getAdjustedMode(const PropertyListT &PropertyList) {
@@ -839,9 +839,8 @@ class accessor :
 #if __cplusplus > 201402L
 
   template <typename TagT> static constexpr bool IsValidTag() {
-    return std::is_same<TagT, mode_tag_t<AccessMode>>::value ||
-           std::is_same<TagT,
-                        mode_target_tag_t<AccessMode, AccessTarget>>::value;
+    return detail::is_same_v<TagT, mode_tag_t<AccessMode>> ||
+           detail::is_same_v<TagT, mode_target_tag_t<AccessMode, AccessTarget>>;
   }
 
 #endif
@@ -959,7 +958,7 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename detail::enable_if_t<
                 detail::IsRunTimePropertyListT<PropertyListT>::value &&
-                std::is_same<T, DataT>::value && Dims == 0 &&
+                detail::is_same_v<T, DataT> && Dims == 0 &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
   accessor(buffer<T, 1, AllocatorT> &BufferRef,
@@ -984,7 +983,7 @@ class accessor :
             typename... PropTypes,
             typename detail::enable_if_t<
                 detail::IsCxPropertyList<PropertyListT>::value &&
-                std::is_same<T, DataT>::value && Dims == 0 &&
+                detail::is_same_v<T, DataT> && Dims == 0 &&
                 ((!IsPlaceH && IsHostBuf) ||
                  (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))> * = nullptr>
   accessor(
@@ -1009,7 +1008,7 @@ class accessor :
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = typename detail::enable_if_t<
                 detail::IsRunTimePropertyListT<PropertyListT>::value &&
-                std::is_same<T, DataT>::value && (Dims == 0) &&
+                detail::is_same_v<T, DataT> && (Dims == 0) &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
            const property_list &PropertyList = {})
@@ -1034,7 +1033,7 @@ class accessor :
             typename... PropTypes,
             typename = typename detail::enable_if_t<
                 detail::IsCxPropertyList<PropertyListT>::value &&
-                std::is_same<T, DataT>::value && (Dims == 0) &&
+                detail::is_same_v<T, DataT> && (Dims == 0) &&
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf || IsHostBuf))>>
   accessor(
       buffer<T, 1, AllocatorT> &BufferRef, handler &CommandGroupHandler,
@@ -2051,13 +2050,13 @@ class host_accessor
   constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
 
   template <typename T, int Dims> static constexpr bool IsSameAsBuffer() {
-    return std::is_same<T, DataT>::value && (Dims > 0) && (Dims == Dimensions);
+    return detail::is_same_v<T, DataT> && (Dims > 0) && (Dims == Dimensions);
   }
 
 #if __cplusplus > 201402L
 
   template <typename TagT> static constexpr bool IsValidTag() {
-    return std::is_same<TagT, mode_tag_t<AccessMode>>::value;
+    return detail::is_same_v<TagT, mode_tag_t<AccessMode>>;
   }
 
 #endif
@@ -2098,7 +2097,7 @@ class host_accessor
 
   template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
             typename = typename detail::enable_if_t<
-                std::is_same<T, DataT>::value && Dims == 0>>
+                detail::is_same_v<T, DataT> && Dims == 0>>
   host_accessor(buffer<T, 1, AllocatorT> &BufferRef,
                 const property_list &PropertyList = {})
       : AccessorT(BufferRef, PropertyList) {}
diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp
index 74ea543d7d0e7..fe0f0cbb9ad22 100644
--- a/sycl/include/CL/sycl/atomic.hpp
+++ b/sycl/include/CL/sycl/atomic.hpp
@@ -20,7 +20,7 @@
 #include <type_traits>
 
 #define __SYCL_STATIC_ASSERT_NOT_FLOAT(T)                                      \
-  static_assert(!std::is_same<T, float>::value,                                \
+  static_assert(!detail::is_same_v<T, float>,                                  \
                 "SYCL atomic function not available for float type")
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -38,11 +38,10 @@ using memory_order = cl::sycl::memory_order;
 
 template <typename T> struct IsValidAtomicType {
   static constexpr bool value =
-      (std::is_same<T, int>::value || std::is_same<T, unsigned int>::value ||
-       std::is_same<T, long>::value || std::is_same<T, unsigned long>::value ||
-       std::is_same<T, long long>::value ||
-       std::is_same<T, unsigned long long>::value ||
-       std::is_same<T, float>::value);
+      (detail::is_same_v<T, int> || detail::is_same_v<T, unsigned int> ||
+       detail::is_same_v<T, long> || detail::is_same_v<T, unsigned long> ||
+       detail::is_same_v<T, long long> ||
+       detail::is_same_v<T, unsigned long long> || detail::is_same_v<T, float>);
 };
 
 template <cl::sycl::access::address_space AS> struct IsValidAtomicAddressSpace {
@@ -224,13 +223,13 @@ class atomic {
 
 #ifdef __SYCL_DEVICE_ONLY__
   template <typename T2 = T>
-  detail::enable_if_t<!std::is_same<cl_float, T2>::value, T>
+  detail::enable_if_t<!detail::is_same_v<cl_float, T2>, T>
   load(memory_order Order = memory_order::relaxed) const {
     return __spirv_AtomicLoad(Ptr, SpirvScope,
                               detail::getSPIRVMemorySemanticsMask(Order));
   }
   template <typename T2 = T>
-  detail::enable_if_t<std::is_same<cl_float, T2>::value, T>
+  detail::enable_if_t<detail::is_same_v<cl_float, T2>, T>
   load(memory_order Order = memory_order::relaxed) const {
     auto *TmpPtr =
         reinterpret_cast<typename multi_ptr<cl_int, addressSpace>::pointer_t>(
diff --git a/sycl/include/CL/sycl/backend/level_zero.hpp b/sycl/include/CL/sycl/backend/level_zero.hpp
index 934b6d8c16802..54757af4ffddc 100644
--- a/sycl/include/CL/sycl/backend/level_zero.hpp
+++ b/sycl/include/CL/sycl/backend/level_zero.hpp
@@ -59,14 +59,14 @@ queue make_queue(const context &Context, pi_native_handle InteropHandle);
 
 // Construction of SYCL platform.
 template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, platform>::value> * = nullptr>
+                          detail::is_same_v<T, platform>> * = nullptr>
 T make(typename interop<backend::level_zero, T>::type Interop) {
   return make_platform(reinterpret_cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL device.
-template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, device>::value> * = nullptr>
+template <typename T, typename detail::enable_if_t<detail::is_same_v<T, device>>
+                          * = nullptr>
 T make(const platform &Platform,
        typename interop<backend::level_zero, T>::type Interop) {
   return make_device(Platform, reinterpret_cast<pi_native_handle>(Interop));
@@ -74,15 +74,15 @@ T make(const platform &Platform,
 
 // Construction of SYCL program.
 template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, program>::value> * = nullptr>
+                          detail::is_same_v<T, program>> * = nullptr>
 T make(const context &Context,
        typename interop<backend::level_zero, T>::type Interop) {
   return make_program(Context, reinterpret_cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL queue.
-template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, queue>::value> * = nullptr>
+template <typename T,
+          typename detail::enable_if_t<detail::is_same_v<T, queue>> * = nullptr>
 T make(const context &Context,
        typename interop<backend::level_zero, T>::type Interop) {
   return make_queue(Context, reinterpret_cast<pi_native_handle>(Interop));
diff --git a/sycl/include/CL/sycl/backend/opencl.hpp b/sycl/include/CL/sycl/backend/opencl.hpp
index abd406260eb80..47d54bbf81725 100644
--- a/sycl/include/CL/sycl/backend/opencl.hpp
+++ b/sycl/include/CL/sycl/backend/opencl.hpp
@@ -67,36 +67,36 @@ __SYCL_EXPORT queue make_queue(const context &Context,
 
 // Construction of SYCL platform.
 template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, platform>::value> * = nullptr>
+                          detail::is_same_v<T, platform>> * = nullptr>
 T make(typename interop<backend::opencl, T>::type Interop) {
   return make_platform(detail::pi::cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL device.
-template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, device>::value> * = nullptr>
+template <typename T, typename detail::enable_if_t<detail::is_same_v<T, device>>
+                          * = nullptr>
 T make(typename interop<backend::opencl, T>::type Interop) {
   return make_device(detail::pi::cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL context.
 template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, context>::value> * = nullptr>
+                          detail::is_same_v<T, context>> * = nullptr>
 T make(typename interop<backend::opencl, T>::type Interop) {
   return make_context(detail::pi::cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL program.
 template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, program>::value> * = nullptr>
+                          detail::is_same_v<T, program>> * = nullptr>
 T make(const context &Context,
        typename interop<backend::opencl, T>::type Interop) {
   return make_program(Context, detail::pi::cast<pi_native_handle>(Interop));
 }
 
 // Construction of SYCL queue.
-template <typename T, typename detail::enable_if_t<
-                          std::is_same<T, queue>::value> * = nullptr>
+template <typename T,
+          typename detail::enable_if_t<detail::is_same_v<T, queue>> * = nullptr>
 T make(const context &Context,
        typename interop<backend::opencl, T>::type Interop) {
   return make_queue(Context, detail::pi::cast<pi_native_handle>(Interop));
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index 8951e4c786f08..b76db7128f696 100755
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -57,7 +57,7 @@ class buffer {
                           std::input_iterator_tag>::value>;
   template <typename ItA, typename ItB>
   using EnableIfSameNonConstIterators = typename detail::enable_if_t<
-      std::is_same<ItA, ItB>::value && !std::is_const<ItA>::value, ItA>;
+      detail::is_same_v<ItA, ItB> && !std::is_const<ItA>::value, ItA>;
 
   buffer(const range<dimensions> &bufferRange,
          const property_list &propList = {})
diff --git a/sycl/include/CL/sycl/detail/accessor_impl.hpp b/sycl/include/CL/sycl/detail/accessor_impl.hpp
index 792e848e636b5..772f9bf63d57a 100644
--- a/sycl/include/CL/sycl/detail/accessor_impl.hpp
+++ b/sycl/include/CL/sycl/detail/accessor_impl.hpp
@@ -225,27 +225,24 @@ constexpr access::mode deduceAccessMode() {
   // property_list = {} is not properly detected by deduction guide,
   // when parameter is passed without curly braces: access(buffer, noinit)
   // thus simplest approach is to check 2 last arguments for being a tag
-  if constexpr (std::is_same<MayBeTag1,
-                             mode_tag_t<access::mode::read>>::value ||
-                std::is_same<MayBeTag2,
-                             mode_tag_t<access::mode::read>>::value) {
+  if constexpr (detail::is_same_v<MayBeTag1, mode_tag_t<access::mode::read>> ||
+                detail::is_same_v<MayBeTag2, mode_tag_t<access::mode::read>>) {
     return access::mode::read;
   }
 
-  if constexpr (std::is_same<MayBeTag1,
-                             mode_tag_t<access::mode::write>>::value ||
-                std::is_same<MayBeTag2,
-                             mode_tag_t<access::mode::write>>::value) {
+  if constexpr (detail::is_same_v<MayBeTag1, mode_tag_t<access::mode::write>> ||
+                detail::is_same_v<MayBeTag2, mode_tag_t<access::mode::write>>) {
     return access::mode::write;
   }
 
-  if constexpr (
-      std::is_same<MayBeTag1,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value ||
-      std::is_same<MayBeTag2,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value) {
+  if constexpr (detail::is_same_v<
+                    MayBeTag1,
+                    mode_target_tag_t<access::mode::read,
+                                      access::target::constant_buffer>> ||
+                detail::is_same_v<
+                    MayBeTag2,
+                    mode_target_tag_t<access::mode::read,
+                                      access::target::constant_buffer>>) {
     return access::mode::read;
   }
 
@@ -254,13 +251,14 @@ constexpr access::mode deduceAccessMode() {
 
 template <typename MayBeTag1, typename MayBeTag2>
 constexpr access::target deduceAccessTarget(access::target defaultTarget) {
-  if constexpr (
-      std::is_same<MayBeTag1,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value ||
-      std::is_same<MayBeTag2,
-                   mode_target_tag_t<access::mode::read,
-                                     access::target::constant_buffer>>::value) {
+  if constexpr (detail::is_same_v<
+                    MayBeTag1,
+                    mode_target_tag_t<access::mode::read,
+                                      access::target::constant_buffer>> ||
+                detail::is_same_v<
+                    MayBeTag2,
+                    mode_target_tag_t<access::mode::read,
+                                      access::target::constant_buffer>>) {
     return access::target::constant_buffer;
   }
 
diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp
index 7e14e0aed035f..760693cb6cb0d 100644
--- a/sycl/include/CL/sycl/detail/cg_types.hpp
+++ b/sycl/include/CL/sycl/detail/cg_types.hpp
@@ -195,13 +195,13 @@ class HostKernel : public HostKernelBase {
   char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
 
   template <class ArgT = KernelArgType>
-  typename detail::enable_if_t<std::is_same<ArgT, void>::value>
+  typename detail::enable_if_t<detail::is_same_v<ArgT, void>>
   runOnHost(const NDRDescT &) {
     MKernel();
   }
 
   template <class ArgT = KernelArgType>
-  typename detail::enable_if_t<std::is_same<ArgT, sycl::id<Dims>>::value>
+  typename detail::enable_if_t<detail::is_same_v<ArgT, sycl::id<Dims>>>
   runOnHost(const NDRDescT &NDRDesc) {
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
@@ -221,7 +221,7 @@ class HostKernel : public HostKernelBase {
 
   template <class ArgT = KernelArgType>
   typename detail::enable_if_t<
-      std::is_same<ArgT, item<Dims, /*Offset=*/false>>::value>
+      detail::is_same_v<ArgT, item<Dims, /*Offset=*/false>>>
   runOnHost(const NDRDescT &NDRDesc) {
     sycl::id<Dims> ID;
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
@@ -240,7 +240,7 @@ class HostKernel : public HostKernelBase {
 
   template <class ArgT = KernelArgType>
   typename detail::enable_if_t<
-      std::is_same<ArgT, item<Dims, /*Offset=*/true>>::value>
+      detail::is_same_v<ArgT, item<Dims, /*Offset=*/true>>>
   runOnHost(const NDRDescT &NDRDesc) {
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
@@ -260,7 +260,7 @@ class HostKernel : public HostKernelBase {
   }
 
   template <class ArgT = KernelArgType>
-  typename detail::enable_if_t<std::is_same<ArgT, nd_item<Dims>>::value>
+  typename detail::enable_if_t<detail::is_same_v<ArgT, nd_item<Dims>>>
   runOnHost(const NDRDescT &NDRDesc) {
     sycl::range<Dims> GroupSize(InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
@@ -305,7 +305,7 @@ class HostKernel : public HostKernelBase {
   }
 
   template <typename ArgT = KernelArgType>
-  enable_if_t<std::is_same<ArgT, cl::sycl::group<Dims>>::value>
+  enable_if_t<detail::is_same_v<ArgT, cl::sycl::group<Dims>>>
   runOnHost(const NDRDescT &NDRDesc) {
     sycl::range<Dims> NGroups(InitializedVal<Dims, range>::template get<0>());
 
diff --git a/sycl/include/CL/sycl/detail/generic_type_traits.hpp b/sycl/include/CL/sycl/detail/generic_type_traits.hpp
index 6e3966028724d..7a10758e61303 100644
--- a/sycl/include/CL/sycl/detail/generic_type_traits.hpp
+++ b/sycl/include/CL/sycl/detail/generic_type_traits.hpp
@@ -296,7 +296,7 @@ template <typename T> class TryToGetPointerT {
 public:
   using type = decltype(check(T()));
   static constexpr bool value =
-      std::is_pointer<T>::value || !std::is_same<T, type>::value;
+      std::is_pointer<T>::value || !detail::is_same_v<T, type>;
 };
 
 // Try to get element_type, otherwise T
@@ -306,7 +306,7 @@ template <typename T> class TryToGetElementType {
 
 public:
   using type = decltype(check(T()));
-  static constexpr bool value = !std::is_same<T, type>::value;
+  static constexpr bool value = !detail::is_same_v<T, type>;
 };
 
 // Try to get vector_t, otherwise T
@@ -316,7 +316,7 @@ template <typename T> class TryToGetVectorT {
 
 public:
   using type = decltype(check(T()));
-  static constexpr bool value = !std::is_same<T, type>::value;
+  static constexpr bool value = !detail::is_same_v<T, type>;
 };
 
 // Try to get pointer_t (if pointer_t indicates on the type with_remainder
@@ -395,7 +395,7 @@ using select_cl_scalar_t = conditional_t<
         std::is_floating_point<T>::value, select_cl_scalar_float_t<T>,
         // half is a special case: it is implemented differently on host and
         // device and therefore, might lower to different types
-        conditional_t<std::is_same<T, half>::value,
+        conditional_t<detail::is_same_v<T, half>,
                       cl::sycl::detail::half_impl::BIsRepresentationT, T>>>;
 
 // select_cl_vector_or_scalar does cl_* type selection for element type of
@@ -410,7 +410,7 @@ struct select_cl_vector_or_scalar<
       // select_cl_scalar_t returns _Float16, so, we try to instantiate vec
       // class with _Float16 DataType, which is not expected there
       // So, leave vector<half, N> as-is
-      vec<conditional_t<std::is_same<typename T::element_type, half>::value,
+      vec<conditional_t<detail::is_same_v<typename T::element_type, half>,
                         typename T::element_type,
                         select_cl_scalar_t<typename T::element_type>>,
           T::get_count()>;
diff --git a/sycl/include/CL/sycl/detail/stl_type_traits.hpp b/sycl/include/CL/sycl/detail/stl_type_traits.hpp
index 03403b7ba063e..ccf261969f77c 100644
--- a/sycl/include/CL/sycl/detail/stl_type_traits.hpp
+++ b/sycl/include/CL/sycl/detail/stl_type_traits.hpp
@@ -36,6 +36,9 @@ using remove_reference_t = typename std::remove_reference<T>::type;
 
 template <typename T> using add_pointer_t = typename std::add_pointer<T>::type;
 
+template <typename T, typename U>
+constexpr bool is_same_v = std::is_same<T, U>::value;
+
 // C++17
 template <bool V> using bool_constant = std::integral_constant<bool, V>;
 
diff --git a/sycl/include/CL/sycl/detail/sycl_mem_obj_allocator.hpp b/sycl/include/CL/sycl/detail/sycl_mem_obj_allocator.hpp
index 827a54ff9f6af..1b75c835bdded 100644
--- a/sycl/include/CL/sycl/detail/sycl_mem_obj_allocator.hpp
+++ b/sycl/include/CL/sycl/detail/sycl_mem_obj_allocator.hpp
@@ -37,11 +37,11 @@ class SYCLMemObjAllocatorHolder : public SYCLMemObjAllocator {
 
   template <typename T>
   using EnableIfDefaultAllocator =
-      enable_if_t<std::is_same<T, sycl_memory_object_allocator>::value>;
+      enable_if_t<detail::is_same_v<T, sycl_memory_object_allocator>>;
 
   template <typename T>
   using EnableIfNonDefaultAllocator =
-      enable_if_t<!std::is_same<T, sycl_memory_object_allocator>::value>;
+      enable_if_t<!detail::is_same_v<T, sycl_memory_object_allocator>>;
 
 public:
   SYCLMemObjAllocatorHolder(AllocatorT Allocator)
diff --git a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
index 2dd5ba8d427ca..1ba9beaeec39c 100644
--- a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
+++ b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
@@ -55,11 +55,11 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
 
   template <typename T>
   using EnableIfDefaultAllocator =
-      enable_if_t<std::is_same<T, sycl_memory_object_allocator>::value>;
+      enable_if_t<detail::is_same_v<T, sycl_memory_object_allocator>>;
 
   template <typename T>
   using EnableIfNonDefaultAllocator =
-      enable_if_t<!std::is_same<T, sycl_memory_object_allocator>::value>;
+      enable_if_t<!detail::is_same_v<T, sycl_memory_object_allocator>>;
 
 public:
   SYCLMemObjT(const size_t SizeInBytes, const property_list &Props,
diff --git a/sycl/include/CL/sycl/detail/type_list.hpp b/sycl/include/CL/sycl/detail/type_list.hpp
index 2caeabe384e58..fa224023203e0 100644
--- a/sycl/include/CL/sycl/detail/type_list.hpp
+++ b/sycl/include/CL/sycl/detail/type_list.hpp
@@ -25,9 +25,8 @@ template <typename... T> struct type_list;
 using empty_type_list = type_list<>;
 
 template <typename T>
-struct is_empty_type_list
-    : conditional_t<std::is_same<T, empty_type_list>::value, std::true_type,
-                    std::false_type> {};
+struct is_empty_type_list : conditional_t<detail::is_same_v<T, empty_type_list>,
+                                          std::true_type, std::false_type> {};
 
 template <> struct type_list<> {};
 
@@ -52,7 +51,7 @@ struct type_list<type_list<H, T...>, T2...> {
 // is_contained
 template <typename T, typename L>
 struct is_contained
-    : conditional_t<std::is_same<remove_cv_t<T>, head_t<L>>::value,
+    : conditional_t<detail::is_same_v<remove_cv_t<T>, head_t<L>>,
                     std::true_type, is_contained<T, tail_t<L>>> {};
 
 template <typename T>
diff --git a/sycl/include/CL/sycl/detail/type_traits.hpp b/sycl/include/CL/sycl/detail/type_traits.hpp
index ee2a3bd9a327d..a41bc966a7ede 100644
--- a/sycl/include/CL/sycl/detail/type_traits.hpp
+++ b/sycl/include/CL/sycl/detail/type_traits.hpp
@@ -209,8 +209,8 @@ struct is_vector_arithmetic
 
 // is_bool
 template <typename T>
-struct is_scalar_bool
-    : bool_constant<std::is_same<remove_cv_t<T>, bool>::value> {};
+struct is_scalar_bool : bool_constant<detail::is_same_v<remove_cv_t<T>, bool>> {
+};
 
 template <typename T>
 struct is_vector_bool
@@ -299,7 +299,7 @@ template <typename T, int N> struct make_larger_impl<vec<T, N>, vec<T, N>> {
   using base_type = vector_element_t<vec<T, N>>;
   using upper_type = typename make_larger_impl<base_type, base_type>::type;
   using new_type = vec<upper_type, N>;
-  static constexpr bool found = !std::is_same<upper_type, void>::value;
+  static constexpr bool found = !detail::is_same_v<upper_type, void>;
   using type = conditional_t<found, new_type, void>;
 };
 
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 4bdffaacc6570..b7a739dab7d5a 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -161,8 +161,8 @@ template <int Dims> struct NotIntMsg<id<Dims>> {
 
 #if __SYCL_ID_QUERIES_FIT_IN_INT__
 template <typename T, typename ValT>
-typename detail::enable_if_t<std::is_same<ValT, size_t>::value ||
-                             std::is_same<ValT, unsigned long long>::value>
+typename detail::enable_if_t<detail::is_same_v<ValT, size_t> ||
+                             detail::is_same_v<ValT, unsigned long long>>
 checkValueRangeImpl(ValT V) {
   static constexpr size_t Limit =
       static_cast<size_t>((std::numeric_limits<int>::max)());
@@ -172,8 +172,8 @@ checkValueRangeImpl(ValT V) {
 #endif
 
 template <int Dims, typename T>
-typename detail::enable_if_t<std::is_same<T, range<Dims>>::value ||
-                             std::is_same<T, id<Dims>>::value>
+typename detail::enable_if_t<detail::is_same_v<T, range<Dims>> ||
+                             detail::is_same_v<T, id<Dims>>>
 checkValueRange(const T &V) {
 #if __SYCL_ID_QUERIES_FIT_IN_INT__
   for (size_t Dim = 0; Dim < Dims; ++Dim)
@@ -210,7 +210,7 @@ void checkValueRange(const range<Dims> &R, const id<Dims> &O) {
 }
 
 template <int Dims, typename T>
-typename detail::enable_if_t<std::is_same<T, nd_range<Dims>>::value>
+typename detail::enable_if_t<detail::is_same_v<T, nd_range<Dims>>>
 checkValueRange(const T &V) {
 #if __SYCL_ID_QUERIES_FIT_IN_INT__
   checkValueRange<Dims>(V.get_global_range());
@@ -430,7 +430,7 @@ class __SYCL_EXPORT handler {
   template <typename T> void setArgHelper(int ArgIndex, T &&Arg) {
     auto StoredArg = static_cast<void *>(storePlainArg(Arg));
 
-    if (!std::is_same<cl_mem, T>::value && std::is_pointer<T>::value) {
+    if (!detail::is_same_v<cl_mem, T> && std::is_pointer<T>::value) {
       MArgs.emplace_back(detail::kernel_param_kind_t::kind_pointer, StoredArg,
                          sizeof(T), ArgIndex);
     } else {
@@ -852,7 +852,8 @@ class __SYCL_EXPORT handler {
       typename detail::remove_cv_t<detail::remove_reference_t<T>>;
 
   template <typename U, typename T>
-  using is_same_type = std::is_same<remove_cv_ref_t<U>, remove_cv_ref_t<T>>;
+  static constexpr bool is_same_type_v =
+      detail::is_same_v<remove_cv_ref_t<U>, remove_cv_ref_t<T>>;
 
   template <typename T> struct ShouldEnableSetArg {
     static constexpr bool value =
@@ -860,10 +861,10 @@ class __SYCL_EXPORT handler {
 #if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707
             && std::is_standard_layout<detail::remove_reference_t<T>>::value
 #endif
-        || is_same_type<sampler, T>::value // Sampler
-        || (!is_same_type<cl_mem, T>::value &&
+        || is_same_type_v<sampler, T> // Sampler
+        || (!is_same_type_v<cl_mem, T> &&
             std::is_pointer<remove_cv_ref_t<T>>::value) // USM
-        || is_same_type<cl_mem, T>::value;              // Interop
+        || is_same_type_v<cl_mem, T>;                   // Interop
   };
 
   /// Sets argument for OpenCL interoperability kernels.
diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp
index c3a298c9bfbc7..a5267ad9aedc4 100644
--- a/sycl/include/CL/sycl/multi_ptr.hpp
+++ b/sycl/include/CL/sycl/multi_ptr.hpp
@@ -29,9 +29,9 @@ class accessor;
 template <typename ElementType, access::address_space Space> class multi_ptr {
 public:
   using element_type =
-      detail::conditional_t<std::is_same<ElementType, half>::value,
-                    cl::sycl::detail::half_impl::BIsRepresentationT,
-                    ElementType>;
+      detail::conditional_t<detail::is_same_v<ElementType, half>,
+                            cl::sycl::detail::half_impl::BIsRepresentationT,
+                            ElementType>;
   using difference_type = std::ptrdiff_t;
 
   // Implementation defined pointer and reference types that correspond to
@@ -156,15 +156,14 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
 
   // Only if Space == global_space || global_device_space and element type is
   // const
-  template <
-      int dimensions, access::mode Mode, access::placeholder isPlaceholder,
-      typename PropertyListT, access::address_space _Space = Space,
-      typename ET = ElementType,
-      typename = typename detail::enable_if_t<
-          _Space == Space &&
-          (Space == access::address_space::global_space ||
-           Space == access::address_space::global_device_space) &&
-          std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
+  template <int dimensions, access::mode Mode,
+            access::placeholder isPlaceholder, typename PropertyListT,
+            access::address_space _Space = Space, typename ET = ElementType,
+            typename = typename detail::enable_if_t<
+                _Space == Space &&
+                (Space == access::address_space::global_space ||
+                 Space == access::address_space::global_device_space) &&
+                std::is_const<ET>::value && detail::is_same_v<ET, ElementType>>>
   multi_ptr(
       accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
                access::target::global_buffer, isPlaceholder, PropertyListT>
@@ -178,7 +177,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
       typename ET = ElementType,
       typename = typename detail::enable_if_t<
           _Space == Space && Space == access::address_space::local_space &&
-          std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
+          std::is_const<ET>::value && detail::is_same_v<ET, ElementType>>>
   multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
                      access::target::local, isPlaceholder, PropertyListT>
                 Accessor)
@@ -191,7 +190,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
       typename ET = ElementType,
       typename = typename detail::enable_if_t<
           _Space == Space && Space == access::address_space::constant_space &&
-          std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
+          std::is_const<ET>::value && detail::is_same_v<ET, ElementType>>>
   multi_ptr(
       accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
                access::target::constant_buffer, isPlaceholder, PropertyListT>
@@ -206,7 +205,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   //   multi_ptr<ElementType, Space> -> multi_ptr<const ElementTYpe, Space>
   template <typename ET = ElementType>
   multi_ptr(typename detail::enable_if_t<
-            std::is_const<ET>::value && std::is_same<ET, ElementType>::value,
+            std::is_const<ET>::value && detail::is_same_v<ET, ElementType>,
             const multi_ptr<typename detail::remove_const_t<ET>, Space>> &ETP)
       : m_Pointer(ETP.get()) {}
 
@@ -220,7 +219,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   // Only available when ElementType is not const-qualified
   template <typename ET = ElementType>
   operator multi_ptr<
-      typename detail::enable_if_t<std::is_same<ET, ElementType>::value &&
+      typename detail::enable_if_t<detail::is_same_v<ET, ElementType> &&
                                        !std::is_const<ET>::value,
                                    void>::type,
       Space>() const {
@@ -232,7 +231,7 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   // Only available when ElementType is const-qualified
   template <typename ET = ElementType>
   operator multi_ptr<
-      typename detail::enable_if_t<std::is_same<ET, ElementType>::value &&
+      typename detail::enable_if_t<detail::is_same_v<ET, ElementType> &&
                                        std::is_const<ET>::value,
                                    const void>::type,
       Space>() const {
diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp
index f2704b7491034..abf57bef7be4d 100644
--- a/sycl/include/CL/sycl/stream.hpp
+++ b/sycl/include/CL/sycl/stream.hpp
@@ -55,9 +55,9 @@ constexpr size_t MAX_ARRAY_SIZE =
 
 template <class F, class T = void>
 using EnableIfFP =
-    typename detail::enable_if_t<std::is_same<F, float>::value ||
-                                     std::is_same<F, double>::value ||
-                                     std::is_same<F, half>::value,
+    typename detail::enable_if_t<detail::is_same_v<F, float> ||
+                                     detail::is_same_v<F, double> ||
+                                     detail::is_same_v<F, half>,
                                  T>;
 
 using GlobalBufAccessorT = accessor<char, 1, cl::sycl::access::mode::read_write,
@@ -186,7 +186,7 @@ inline unsigned append(char *Dst, const char *Src) {
 
 template <typename T>
 inline typename detail::enable_if_t<
-    std::is_same<T, float>::value || std::is_same<T, double>::value, unsigned>
+    detail::is_same_v<T, float> || detail::is_same_v<T, double>, unsigned>
 checkForInfNan(char *Buf, T Val) {
   if (isnan(Val))
     return append(Buf, "nan");
@@ -199,7 +199,7 @@ checkForInfNan(char *Buf, T Val) {
 }
 
 template <typename T>
-inline typename detail::enable_if_t<std::is_same<T, half>::value, unsigned>
+inline typename detail::enable_if_t<detail::is_same_v<T, half>, unsigned>
 checkForInfNan(char *Buf, T Val) {
   if (Val != Val)
     return append(Buf, "nan");
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
index 5f1b895193cd9..c65dcd238e77f 100644
--- a/sycl/include/CL/sycl/types.hpp
+++ b/sycl/include/CL/sycl/types.hpp
@@ -242,13 +242,14 @@ using is_float_to_float =
     std::integral_constant<bool, detail::is_floating_point<T>::value &&
                                      detail::is_floating_point<R>::value>;
 template <typename T>
-using is_standard_type = std::integral_constant<
-    bool, detail::is_sgentype<T>::value && !std::is_same<T, long long>::value &&
-              !std::is_same<T, unsigned long long>::value>;
+using is_standard_type =
+    std::integral_constant<bool, detail::is_sgentype<T>::value &&
+                                     !detail::is_same_v<T, long long> &&
+                                     !detail::is_same_v<T, unsigned long long>>;
 
 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
           typename OpenCLR>
-detail::enable_if_t<std::is_same<T, R>::value, R> convertImpl(T Value) {
+detail::enable_if_t<detail::is_same_v<T, R>, R> convertImpl(T Value) {
   return Value;
 }
 
@@ -258,7 +259,7 @@ detail::enable_if_t<std::is_same<T, R>::value, R> convertImpl(T Value) {
 // implemented for host that takes care of the precision requirements.
 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
           typename OpenCLR>
-detail::enable_if_t<!std::is_same<T, R>::value &&
+detail::enable_if_t<!detail::is_same_v<T, R> &&
                         (is_int_to_int<T, R>::value ||
                          is_int_to_float<T, R>::value ||
                          is_float_to_float<T, R>::value),
@@ -320,7 +321,7 @@ using Rtn = detail::bool_constant<Mode == rounding_mode::rtn>;
 template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
           typename OpenCLR>
 detail::enable_if_t<
-    !std::is_same<T, R>::value && std::is_same<OpenCLT, OpenCLR>::value, R>
+    !detail::is_same_v<T, R> && detail::is_same_v<OpenCLT, OpenCLR>, R>
 convertImpl(T Value) {
   return static_cast<R>(Value);
 }
@@ -330,10 +331,10 @@ convertImpl(T Value) {
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
   detail::enable_if_t<is_sint_to_sint<T, R>::value &&                          \
-                          !std::is_same<OpenCLT, OpenCLR>::value &&            \
-                          (std::is_same<OpenCLR, DestType>::value ||           \
-                           (std::is_same<OpenCLR, signed char>::value &&       \
-                            std::is_same<DestType, char>::value)),             \
+                          !detail::is_same_v<OpenCLT, OpenCLR> &&              \
+                          (detail::is_same_v<OpenCLR, DestType> ||             \
+                           (detail::is_same_v<OpenCLR, signed char> &&         \
+                            detail::is_same_v<DestType, char>)),               \
                       R>                                                       \
   convertImpl(T Value) {                                                       \
     OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
@@ -352,8 +353,8 @@ __SYCL_GENERATE_CONVERT_IMPL(long)
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
   detail::enable_if_t<is_uint_to_uint<T, R>::value &&                          \
-                          !std::is_same<OpenCLT, OpenCLR>::value &&            \
-                          std::is_same<OpenCLR, DestType>::value,              \
+                          !detail::is_same_v<OpenCLT, OpenCLR> &&              \
+                          detail::is_same_v<OpenCLR, DestType>,                \
                       R>                                                       \
   convertImpl(T Value) {                                                       \
     OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
@@ -383,9 +384,9 @@ convertImpl(T Value) {
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
   detail::enable_if_t<is_sint_to_float<T, R>::value &&                         \
-                          (std::is_same<OpenCLR, DestType>::value ||           \
-                           (std::is_same<OpenCLR, _Float16>::value &&          \
-                            std::is_same<DestType, half>::value)),             \
+                          (detail::is_same_v<OpenCLR, DestType> ||             \
+                           (detail::is_same_v<OpenCLR, _Float16> &&            \
+                            detail::is_same_v<DestType, half>)),               \
                       R>                                                       \
   convertImpl(T Value) {                                                       \
     OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
@@ -403,9 +404,9 @@ __SYCL_GENERATE_CONVERT_IMPL(SToF, double)
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
   detail::enable_if_t<is_uint_to_float<T, R>::value &&                         \
-                          (std::is_same<OpenCLR, DestType>::value ||           \
-                           (std::is_same<OpenCLR, _Float16>::value &&          \
-                            std::is_same<DestType, half>::value)),             \
+                          (detail::is_same_v<OpenCLR, DestType> ||             \
+                           (detail::is_same_v<OpenCLR, _Float16> &&            \
+                            detail::is_same_v<DestType, half>)),               \
                       R>                                                       \
   convertImpl(T Value) {                                                       \
     OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
@@ -424,10 +425,10 @@ __SYCL_GENERATE_CONVERT_IMPL(UToF, double)
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
   detail::enable_if_t<is_float_to_float<T, R>::value &&                        \
-                          !std::is_same<OpenCLT, OpenCLR>::value &&            \
-                          (std::is_same<OpenCLR, DestType>::value ||           \
-                           (std::is_same<OpenCLR, _Float16>::value &&          \
-                            std::is_same<DestType, half>::value)) &&           \
+                          !detail::is_same_v<OpenCLT, OpenCLR> &&              \
+                          (detail::is_same_v<OpenCLR, DestType> ||             \
+                           (detail::is_same_v<OpenCLR, _Float16> &&            \
+                            detail::is_same_v<DestType, half>)) &&             \
                           RoundingModeCondition<roundingMode>::value,          \
                       R>                                                       \
   convertImpl(T Value) {                                                       \
@@ -454,12 +455,14 @@ __SYCL_GENERATE_CONVERT_IMPL_FOR_ROUNDING_MODE(rtn, Rtn)
                                      RoundingModeCondition)                    \
   template <typename T, typename R, rounding_mode roundingMode,                \
             typename OpenCLT, typename OpenCLR>                                \
-  detail::enable_if_t<is_float_to_int<T, R>::value &&                          \
-                          (std::is_same<OpenCLR, DestType>::value ||           \
-                           std::is_same<OpenCLR, signed char>::value &&        \
-                               std::is_same<DestType, char>::value) &&         \
-                          RoundingModeCondition<roundingMode>::value,          \
-                      R>                                                       \
+  detail::enable_if_t<                                                         \
+      is_float_to_int<T, R>::value &&                                          \
+          (detail::is_same_v<OpenCLR, DestType> ||                             \
+           detail::is_same_v<OpenCLR, signed char> &&                          \
+               detail::is_same_v<                                              \
+                   DestType,                                                   \
+                   char>)&&RoundingModeCondition<roundingMode>::value,         \
+      R>                                                                       \
   convertImpl(T Value) {                                                       \
     OpenCLT OpValue = cl::sycl::detail::convertDataToType<T, OpenCLT>(Value);  \
     return __spirv_Convert##SPIRVOp##_R##DestType##_##RoundingMode(OpValue);   \
@@ -496,7 +499,7 @@ template <typename T, typename R, rounding_mode roundingMode, typename OpenCLT,
 detail::enable_if_t<
     ((!is_standard_type<T>::value && !is_standard_type<OpenCLT>::value) ||
      (!is_standard_type<R>::value && !is_standard_type<OpenCLR>::value)) &&
-        !std::is_same<OpenCLT, OpenCLR>::value,
+        !detail::is_same_v<OpenCLT, OpenCLR>,
     R>
 convertImpl(T Value) {
   return static_cast<R>(Value);
@@ -661,7 +664,7 @@ template <typename Type, int NumElements> class vec {
 
   // W/o this, things like "vec<char,*> = vec<signed char, *>" doesn't work.
   template <typename Ty = DataT>
-  typename detail::enable_if_t<!std::is_same<Ty, rel_t>::value &&
+  typename detail::enable_if_t<!detail::is_same_v<Ty, rel_t> &&
                                    std::is_convertible<Ty, rel_t>::value,
                                vec &>
   operator=(const vec<rel_t, NumElements> &Rhs) {
@@ -672,15 +675,15 @@ template <typename Type, int NumElements> class vec {
 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
   template <typename T = void>
   using EnableIfNotHostHalf = typename detail::enable_if_t<
-      !std::is_same<DataT, cl::sycl::detail::half_impl::half>::value ||
-          !std::is_same<cl::sycl::detail::half_impl::StorageT,
-                        cl::sycl::detail::host_half_impl::half>::value,
+      !detail::is_same_v<DataT, cl::sycl::detail::half_impl::half> ||
+          !detail::is_same_v<cl::sycl::detail::half_impl::StorageT,
+                             cl::sycl::detail::host_half_impl::half>,
       T>;
   template <typename T = void>
   using EnableIfHostHalf = typename detail::enable_if_t<
-      std::is_same<DataT, cl::sycl::detail::half_impl::half>::value &&
-          std::is_same<cl::sycl::detail::half_impl::StorageT,
-                       cl::sycl::detail::host_half_impl::half>::value,
+      detail::is_same_v<DataT, cl::sycl::detail::half_impl::half> &&
+          detail::is_same_v<cl::sycl::detail::half_impl::StorageT,
+                            cl::sycl::detail::host_half_impl::half>,
       T>;
 
   template <typename Ty = DataT>
@@ -691,7 +694,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename detail::enable_if_t<
       std::is_fundamental<Ty>::value ||
-          std::is_same<typename detail::remove_const_t<Ty>, half>::value,
+          detail::is_same_v<typename detail::remove_const_t<Ty>, half>,
       vec &>
   operator=(const EnableIfNotHostHalf<Ty> &Rhs) {
     m_Data = (DataType)Rhs;
@@ -707,7 +710,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename detail::enable_if_t<
       std::is_fundamental<Ty>::value ||
-          std::is_same<typename detail::remove_const_t<Ty>, half>::value,
+          detail::is_same_v<typename detail::remove_const_t<Ty>, half>,
       vec &>
   operator=(const EnableIfHostHalf<Ty> &Rhs) {
     for (int i = 0; i < NumElements; ++i) {
@@ -725,7 +728,7 @@ template <typename Type, int NumElements> class vec {
   template <typename Ty = DataT>
   typename detail::enable_if_t<
       std::is_fundamental<Ty>::value ||
-          std::is_same<typename detail::remove_const_t<Ty>, half>::value,
+          detail::is_same_v<typename detail::remove_const_t<Ty>, half>,
       vec &>
   operator=(const DataT &Rhs) {
     for (int i = 0; i < NumElements; ++i) {
@@ -792,8 +795,8 @@ template <typename Type, int NumElements> class vec {
 #ifdef __SYCL_DEVICE_ONLY__
   template <typename vector_t_ = vector_t,
             typename = typename detail::enable_if_t<
-                std::is_same<vector_t_, vector_t>::value &&
-                !std::is_same<vector_t_, DataT>::value>>
+                detail::is_same_v<vector_t_, vector_t> &&
+                !detail::is_same_v<vector_t_, DataT>>>
   vec(vector_t openclVector) : m_Data(openclVector) {}
   operator vector_t() const { return m_Data; }
 #endif
@@ -948,7 +951,7 @@ template <typename Type, int NumElements> class vec {
   typename detail::enable_if_t<                                                \
       std::is_convertible<DataT, T>::value &&                                  \
           (std::is_fundamental<T>::value ||                                    \
-           std::is_same<typename detail::remove_const_t<T>, half>::value),     \
+           detail::is_same_v<typename detail::remove_const_t<T>, half>),       \
       vec>                                                                     \
   operator BINOP(const T &Rhs) const {                                         \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
@@ -976,7 +979,7 @@ template <typename Type, int NumElements> class vec {
   typename detail::enable_if_t<                                                \
       std::is_convertible<DataT, T>::value &&                                  \
           (std::is_fundamental<T>::value ||                                    \
-           std::is_same<typename detail::remove_const_t<T>, half>::value),     \
+           detail::is_same_v<typename detail::remove_const_t<T>, half>),       \
       vec>                                                                     \
   operator BINOP(const T &Rhs) const {                                         \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
@@ -1034,7 +1037,7 @@ template <typename Type, int NumElements> class vec {
   template <typename T>                                                        \
   typename detail::enable_if_t<std::is_convertible<T, DataT>::value &&         \
                                    (std::is_fundamental<T>::value ||           \
-                                    std::is_same<T, half>::value),             \
+                                    detail::is_same_v<T, half>),               \
                                vec<rel_t, NumElements>>                        \
   operator RELLOGOP(const T &Rhs) const {                                      \
     return *this RELLOGOP vec(static_cast<const DataT &>(Rhs));                \
@@ -1051,7 +1054,7 @@ template <typename Type, int NumElements> class vec {
   template <typename T>                                                        \
   typename detail::enable_if_t<std::is_convertible<T, DataT>::value &&         \
                                    (std::is_fundamental<T>::value ||           \
-                                    std::is_same<T, half>::value),             \
+                                    detail::is_same_v<T, half>),               \
                                vec<rel_t, NumElements>>                        \
   operator RELLOGOP(const T &Rhs) const {                                      \
     return *this RELLOGOP vec(static_cast<const DataT &>(Rhs));                \
@@ -1325,8 +1328,8 @@ template <typename Type, int NumElements> class vec {
 #ifdef __cpp_deduction_guides
 // all compilers supporting deduction guides also support fold expressions
 template <class T, class... U,
-          class = detail::enable_if_t<(std::is_same<T, U>::value && ...)>>
-vec(T, U...)->vec<T, sizeof...(U) + 1>;
+          class = detail::enable_if_t<(detail::is_same_v<T, U> && ...)>>
+vec(T, U...) -> vec<T, sizeof...(U) + 1>;
 #endif
 
 namespace detail {
@@ -1379,13 +1382,13 @@ class SwizzleOp {
   using EnableIfScalarType = typename detail::enable_if_t<
       std::is_convertible<DataT, T>::value &&
       (std::is_fundamental<T>::value ||
-       std::is_same<typename detail::remove_const_t<T>, half>::value)>;
+       detail::is_same_v<typename detail::remove_const_t<T>, half>)>;
 
   template <typename T>
   using EnableIfNoScalarType = typename detail::enable_if_t<
       !std::is_convertible<DataT, T>::value ||
       !(std::is_fundamental<T>::value ||
-        std::is_same<typename detail::remove_const_t<T>, half>::value)>;
+        detail::is_same_v<typename detail::remove_const_t<T>, half>)>;
 
   template <int... Indices>
   using Swizzle =
@@ -1848,7 +1851,7 @@ class SwizzleOp {
 
   template <int IdxNum = getNumElements()>
   CommonDataT getValue(EnableIfOneIndex<IdxNum, size_t> Index) const {
-    if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
+    if (detail::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
       std::array<int, getNumElements()> Idxs{Indexes...};
       return m_Vector->getValue(Idxs[Index]);
     }
@@ -1859,7 +1862,7 @@ class SwizzleOp {
 
   template <int IdxNum = getNumElements()>
   DataT getValue(EnableIfMultipleIndexes<IdxNum, size_t> Index) const {
-    if (std::is_same<OperationCurrentT<DataT>, GetOp<DataT>>::value) {
+    if (detail::is_same_v<OperationCurrentT<DataT>, GetOp<DataT>>) {
       std::array<int, getNumElements()> Idxs{Indexes...};
       return m_Vector->getValue(Idxs[Index]);
     }
@@ -1903,7 +1906,7 @@ class SwizzleOp {
   template <typename T, int Num>                                               \
   typename detail::enable_if_t<                                                \
       std::is_fundamental<T>::value ||                                         \
-          std::is_same<typename detail::remove_const_t<T>, half>::value,       \
+          detail::is_same_v<typename detail::remove_const_t<T>, half>,         \
       vec<T, Num>>                                                             \
   operator BINOP(const T &Lhs, const vec<T, Num> &Rhs) {                       \
     return vec<T, Num>(Lhs) BINOP Rhs;                                         \
@@ -1915,7 +1918,7 @@ class SwizzleOp {
   typename detail::enable_if_t<                                                \
       std::is_convertible<T, T1>::value &&                                     \
           (std::is_fundamental<T>::value ||                                    \
-           std::is_same<typename detail::remove_const_t<T>, half>::value),     \
+           detail::is_same_v<typename detail::remove_const_t<T>, half>),       \
       vec<T1, Num>>                                                            \
   operator BINOP(                                                              \
       const T &Lhs,                                                            \
@@ -1958,7 +1961,7 @@ __SYCL_BINOP(<<)
   typename detail::enable_if_t<                                                \
       std::is_convertible<T, DataT>::value &&                                  \
           (std::is_fundamental<T>::value ||                                    \
-           std::is_same<typename detail::remove_const_t<T>, half>::value),     \
+           detail::is_same_v<typename detail::remove_const_t<T>, half>),       \
       vec<detail::rel_t<DataT>, Num>>                                          \
   operator RELLOGOP(const T &Lhs, const vec<DataT, Num> &Rhs) {                \
     return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs;                      \
@@ -1970,7 +1973,7 @@ __SYCL_BINOP(<<)
   typename detail::enable_if_t<                                                \
       std::is_convertible<T, T1>::value &&                                     \
           (std::is_fundamental<T>::value ||                                    \
-           std::is_same<typename detail::remove_const_t<T>, half>::value),     \
+           detail::is_same_v<typename detail::remove_const_t<T>, half>),       \
       vec<detail::rel_t<T1>, Num>>                                             \
   operator RELLOGOP(                                                           \
       const T &Lhs,                                                            \
diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index 346b9f33dafae..2a4dd254b7ffe 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -688,7 +688,7 @@ void Scheduler::GraphBuilder::markModifiedIfWrite(MemObjRecord *Record,
 
 template <typename T>
 typename detail::enable_if_t<
-    std::is_same<typename std::remove_cv_t<T>, Requirement>::value,
+    detail::is_same_v<typename std::remove_cv_t<T>, Requirement>,
     EmptyCommand *>
 Scheduler::GraphBuilder::addEmptyCmd(Command *Cmd, const std::vector<T *> &Reqs,
                                      const QueueImplPtr &Queue,
diff --git a/sycl/source/detail/scheduler/scheduler.hpp b/sycl/source/detail/scheduler/scheduler.hpp
index 2665365dde815..57bac1e832d14 100644
--- a/sycl/source/detail/scheduler/scheduler.hpp
+++ b/sycl/source/detail/scheduler/scheduler.hpp
@@ -583,7 +583,7 @@ class Scheduler {
 
     template <typename T>
     typename detail::enable_if_t<
-        std::is_same<typename std::remove_cv_t<T>, Requirement>::value,
+        detail::is_same_v<typename std::remove_cv_t<T>, Requirement>,
         EmptyCommand *>
     addEmptyCmd(Command *Cmd, const std::vector<T *> &Req,
                 const QueueImplPtr &Queue, Command::BlockReason Reason);

From 7e69de474b72462117c05c9ed489558a479592fc Mon Sep 17 00:00:00 2001
From: Artem Gindinson <artem.gindinson@intel.com>
Date: Tue, 17 Nov 2020 09:42:02 +0300
Subject: [PATCH 2/3] Use a namespace alias for sycl::detail

Signed-off-by: Artem Gindinson <artem.gindinson@intel.com>
---
 .../sycl/INTEL/esimd/detail/esimd_intrin.hpp  |  2 +-
 .../sycl/INTEL/esimd/detail/esimd_types.hpp   |  7 ++-
 .../CL/sycl/INTEL/esimd/detail/esimd_util.hpp | 50 ++++++++-----------
 sycl/include/CL/sycl/INTEL/esimd/esimd.hpp    |  2 +-
 .../CL/sycl/INTEL/esimd/esimd_math.hpp        | 15 +++---
 5 files changed, 34 insertions(+), 42 deletions(-)

diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
index 111fd37c3e024..292f4a1b672b5 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_intrin.hpp
@@ -179,7 +179,7 @@ readRegion(const vector_type_t<BT, BN> &Base, std::pair<T, U> Region) {
     return readRegion<BT1, BN1>(Base1, Region.first);
   else {
     static_assert(T::Is_2D);
-    static_assert(sycl::detail::is_same_v<ElemTy, BT1>);
+    static_assert(csd::is_same_v<ElemTy, BT1>);
     // To read a 2D region, we need the parent region
     // Read full rows with non-trivial vertical and horizontal stride = 1.
     constexpr int M = T::Size_y * PaTy::Size_x;
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 aae6535aff487..590ff96d5b67f 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_types.hpp
@@ -209,7 +209,7 @@ template <typename U> constexpr bool is_type() { return false; }
 template <typename U, typename T, typename... Ts> constexpr bool is_type() {
   using UU = typename detail::remove_const_t<U>;
   using TT = typename detail::remove_const_t<T>;
-  return sycl::detail::is_same_v<UU, TT> || is_type<UU, Ts...>();
+  return csd::is_same_v<UU, TT> || is_type<UU, Ts...>();
 }
 
 // calculates the number of elements in "To" type
@@ -228,13 +228,12 @@ struct bitcast_helper {
 template <typename ToEltTy, typename FromEltTy, int FromN,
           typename = csd::enable_if_t<is_vectorizable<ToEltTy>::value>>
 ESIMD_INLINE typename detail::conditional_t<
-    sycl::detail::is_same_v<FromEltTy, ToEltTy>,
-    vector_type_t<FromEltTy, FromN>,
+    csd::is_same_v<FromEltTy, ToEltTy>, vector_type_t<FromEltTy, FromN>,
     vector_type_t<ToEltTy,
                   bitcast_helper<ToEltTy, FromEltTy, FromN>::nToElems()>>
 bitcast(vector_type_t<FromEltTy, FromN> Val) {
   // Noop.
-  if constexpr (sycl::detail::is_same_v<FromEltTy, ToEltTy>)
+  if constexpr (csd::is_same_v<FromEltTy, ToEltTy>)
     return Val;
 
   // Bitcast
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 5bd1c5db2d62b..1cfc26ddf6651 100755
--- a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_util.hpp
@@ -100,11 +100,10 @@ struct is_esimd_scalar
 template <typename T>
 struct is_dword_type
     : std::integral_constant<
-          bool,
-          sycl::detail::is_same_v<int,
-                                  typename sycl::detail::remove_const_t<T>> ||
-              sycl::detail::is_same_v<
-                  unsigned int, typename sycl::detail::remove_const_t<T>>> {};
+          bool, csd::is_same_v<int, typename sycl::detail::remove_const_t<T>> ||
+                    csd::is_same_v<unsigned int,
+                                   typename sycl::detail::remove_const_t<T>>> {
+};
 
 template <typename T, int N>
 struct is_dword_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -120,10 +119,9 @@ template <typename T>
 struct is_word_type
     : std::integral_constant<
           bool,
-          sycl::detail::is_same_v<short,
-                                  typename sycl::detail::remove_const_t<T>> ||
-              sycl::detail::is_same_v<
-                  unsigned short, typename sycl::detail::remove_const_t<T>>> {};
+          csd::is_same_v<short, typename sycl::detail::remove_const_t<T>> ||
+              csd::is_same_v<unsigned short,
+                             typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_word_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -138,10 +136,9 @@ template <typename T>
 struct is_byte_type
     : std::integral_constant<
           bool,
-          sycl::detail::is_same_v<char,
-                                  typename sycl::detail::remove_const_t<T>> ||
-              sycl::detail::is_same_v<
-                  unsigned char, typename sycl::detail::remove_const_t<T>>> {};
+          csd::is_same_v<char, typename sycl::detail::remove_const_t<T>> ||
+              csd::is_same_v<unsigned char,
+                             typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_byte_type<sycl::INTEL::gpu::vector_type<T, N>> {
@@ -155,34 +152,31 @@ template <typename T, int N> struct is_byte_type<sycl::INTEL::gpu::simd<T, N>> {
 template <typename T>
 struct is_fp_type
     : std::integral_constant<
-          bool, sycl::detail::is_same_v<
-                    float, typename sycl::detail::remove_const_t<T>>> {};
+          bool,
+          csd::is_same_v<float, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_df_type
     : std::integral_constant<
-          bool, sycl::detail::is_same_v<
-                    double, typename sycl::detail::remove_const_t<T>>> {};
+          bool,
+          csd::is_same_v<double, typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_fp_or_dword_type
     : std::integral_constant<
           bool,
-          sycl::detail::is_same_v<float,
-                                  typename sycl::detail::remove_const_t<T>> ||
-              sycl::detail::is_same_v<
-                  int, typename sycl::detail::remove_const_t<T>> ||
-              sycl::detail::is_same_v<
-                  unsigned int, typename sycl::detail::remove_const_t<T>>> {};
+          csd::is_same_v<float, typename sycl::detail::remove_const_t<T>> ||
+              csd::is_same_v<int, typename sycl::detail::remove_const_t<T>> ||
+              csd::is_same_v<unsigned int,
+                             typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T>
 struct is_qword_type
     : std::integral_constant<
-          bool, sycl::detail::is_same_v<
-                    long long, typename sycl::detail::remove_const_t<T>> ||
-                    sycl::detail::is_same_v<
-                        unsigned long long,
-                        typename sycl::detail::remove_const_t<T>>> {};
+          bool,
+          csd::is_same_v<long long, typename sycl::detail::remove_const_t<T>> ||
+              csd::is_same_v<unsigned long long,
+                             typename sycl::detail::remove_const_t<T>>> {};
 
 template <typename T, int N>
 struct is_qword_type<sycl::INTEL::gpu::vector_type<T, N>> {
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
index b3437c7c9dcf6..af145bf01eaa8 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd.hpp
@@ -403,7 +403,7 @@ template <typename Ty, int N> class simd {
         Base1 = __esimd_wrregion<ElemTy, BN1, M,
                                  /*VS*/ 0, M, Stride>(Base1, Val, Offset);
       } else {
-        static_assert(sycl::detail::is_same_v<ElemTy, BT>);
+        static_assert(csd::is_same_v<ElemTy, BT>);
         // Read columns with non-trivial horizontal stride.
         constexpr int M = TR::length;
         constexpr int VS = PaTy::Size_x * TR::Stride_y;
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
index dc1bf201ac3d9..5930dad8c30b4 100644
--- a/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
+++ b/sycl/include/CL/sycl/INTEL/esimd/esimd_math.hpp
@@ -73,8 +73,8 @@ __esimd_abs_common_internal(T1 src0, int flag = GENX_NOSAT) {
 
 template <typename T0, typename T1, int SZ>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
-    !sycl::detail::is_same_v<typename sycl::detail::remove_const_t<T0>,
-                             typename sycl::detail::remove_const_t<T1>>,
+    !csd::is_same_v<typename sycl::detail::remove_const_t<T0>,
+                    typename sycl::detail::remove_const_t<T1>>,
     simd<T0, SZ>>
 esimd_abs(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
   return detail::__esimd_abs_common_internal<T0, T1, SZ>(src0, flag);
@@ -82,8 +82,8 @@ esimd_abs(simd<T1, SZ> src0, int flag = GENX_NOSAT) {
 
 template <typename T0, typename T1>
 ESIMD_NODEBUG ESIMD_INLINE typename sycl::detail::enable_if_t<
-    !sycl::detail::is_same_v<typename sycl::detail::remove_const_t<T0>,
-                             typename sycl::detail::remove_const_t<T1>> &&
+    !csd::is_same_v<typename sycl::detail::remove_const_t<T0>,
+                    typename sycl::detail::remove_const_t<T1>> &&
         detail::is_esimd_scalar<T0>::value &&
         detail::is_esimd_scalar<T1>::value,
     typename sycl::detail::remove_const_t<T0>>
@@ -1905,12 +1905,11 @@ ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
 
 template <typename T0, typename T1, int SZ, typename BinaryOperation>
 ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
-  if constexpr (sycl::detail::is_same_v<remove_cvref_t<BinaryOperation>,
-                                        std::plus<>>) {
+  if constexpr (csd::is_same_v<remove_cvref_t<BinaryOperation>, std::plus<>>) {
     T0 retv = detail::esimd_sum<T0>(v);
     return retv;
-  } else if constexpr (sycl::detail::is_same_v<remove_cvref_t<BinaryOperation>,
-                                               std::multiplies<>>) {
+  } else if constexpr (csd::is_same_v<remove_cvref_t<BinaryOperation>,
+                                      std::multiplies<>>) {
     T0 retv = detail::esimd_prod<T0>(v);
     return retv;
   }

From b4fd5f5849b25914fcfe9f16cf0de18c9d5b8ba1 Mon Sep 17 00:00:00 2001
From: Artem Gindinson <artem.gindinson@intel.com>
Date: Tue, 17 Nov 2020 10:41:24 +0300
Subject: [PATCH 3/3] Apply clang-format, as weird as it may seem

Signed-off-by: Artem Gindinson <artem.gindinson@intel.com>
---
 sycl/include/CL/sycl/ONEAPI/reduction.hpp | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/ONEAPI/reduction.hpp b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
index ae8e87459c58d..92aee66f27395 100644
--- a/sycl/include/CL/sycl/ONEAPI/reduction.hpp
+++ b/sycl/include/CL/sycl/ONEAPI/reduction.hpp
@@ -84,9 +84,9 @@ template <typename T, class BinaryOperation>
 using IsReduOptForFastReduce = detail::bool_constant<(
     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
     detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
-    detail::is_same_v<T, double>) && (IsReduPlus<T, BinaryOperation>::value ||
-                                     IsReduMinimum<T, BinaryOperation>::value ||
-                                     IsReduMaximum<T, BinaryOperation>::value)>;
+    detail::is_same_v<T, double>)&&(IsReduPlus<T, BinaryOperation>::value ||
+                                    IsReduMinimum<T, BinaryOperation>::value ||
+                                    IsReduMaximum<T, BinaryOperation>::value)>;
 
 // Identity = 0
 template <typename T, class BinaryOperation>
@@ -97,7 +97,7 @@ using IsZeroIdentityOp = bool_constant<
       IsReduBitOR<T, BinaryOperation>::value ||
       IsReduBitXOR<T, BinaryOperation>::value)) ||
     ((detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
-      detail::is_same_v<T, double>) && IsReduPlus<T, BinaryOperation>::value)>;
+      detail::is_same_v<T, double>)&&IsReduPlus<T, BinaryOperation>::value)>;
 
 // Identity = 1
 template <typename T, class BinaryOperation>
@@ -105,7 +105,7 @@ using IsOneIdentityOp = bool_constant<(
     is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
     detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
-    detail::is_same_v<T, double>) && IsReduMultiplies<T, BinaryOperation>::value>;
+    detail::is_same_v<T, double>)&&IsReduMultiplies<T, BinaryOperation>::value>;
 
 // Identity = ~0
 template <typename T, class BinaryOperation>
@@ -128,7 +128,7 @@ using IsMaximumIdentityOp = bool_constant<(
     is_geninteger8bit<T>::value || is_geninteger16bit<T>::value ||
     is_geninteger32bit<T>::value || is_geninteger64bit<T>::value ||
     detail::is_same_v<T, half> || detail::is_same_v<T, float> ||
-    detail::is_same_v<T, double>) && IsReduMaximum<T, BinaryOperation>::value>;
+    detail::is_same_v<T, double>)&&IsReduMaximum<T, BinaryOperation>::value>;
 
 template <typename T, class BinaryOperation>
 using IsKnownIdentityOp =