From e17525af2e2aa5b4d24cbc9af9a45d81a19eca82 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Thu, 13 Feb 2020 13:42:38 -0800
Subject: [PATCH 01/17] [SYCL] Add static members to sub_group class

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/sub_group.hpp      | 6 ++++++
 sycl/include/CL/sycl/intel/sub_group_host.hpp | 6 ++++++
 2 files changed, 12 insertions(+)

diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 12dfb0eb262f7..0c8d0c434b3e2 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -132,6 +132,12 @@ void store(multi_ptr<T, Space> dst, const vec<T, N> &x) {
 namespace intel {
 
 struct sub_group {
+
+  using id_type = id<1>;
+  using range_type = range<1>;
+  using linear_id_type = size_t;
+  static constexpr int dimensions = 1;
+
   /* --- common interface members --- */
 
   id<1> get_local_id() const {
diff --git a/sycl/include/CL/sycl/intel/sub_group_host.hpp b/sycl/include/CL/sycl/intel/sub_group_host.hpp
index 2496d116e28d3..e24fadc616162 100644
--- a/sycl/include/CL/sycl/intel/sub_group_host.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group_host.hpp
@@ -20,6 +20,12 @@ namespace sycl {
 template <typename T, access::address_space Space> class multi_ptr;
 namespace intel {
 struct sub_group {
+
+  typedef id<1> id_type;
+  typedef range<1> range_type;
+  typedef size_t linear_id_type;
+  static constexpr int dimensions = 1;
+
   /* --- common interface members --- */
 
   id<1> get_local_id() const {

From efeb49a8b773b9f615a6395cb249261da44a109b Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Thu, 13 Feb 2020 14:14:06 -0800
Subject: [PATCH 02/17] [SYCL] Enable algorithm support for sub_group

- sub_group member functions marked deprecated, to be removed later.
- SPIR-V helpers expanded to convert SYCL group to SPIR-V scope.

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/detail/spirv.hpp         |  31 ++-
 .../include/CL/sycl/intel/group_algorithm.hpp | 255 ++++++++++++------
 sycl/include/CL/sycl/intel/sub_group.hpp      |  33 ++-
 3 files changed, 226 insertions(+), 93 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp
index 0f437f13e00ed..e3c783c7bf65d 100644
--- a/sycl/include/CL/sycl/detail/spirv.hpp
+++ b/sycl/include/CL/sycl/detail/spirv.hpp
@@ -16,25 +16,46 @@
 #ifdef __SYCL_DEVICE_ONLY__
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+namespace intel {
+struct sub_group;
+} // namespace intel
 namespace detail {
 namespace spirv {
 
+template <typename Group> struct group_scope {};
+
+template <int Dimensions> struct group_scope<group<Dimensions>> {
+  static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Workgroup;
+};
+
+template <> struct group_scope<intel::sub_group> {
+  static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup;
+};
+
+template <typename Group> bool GroupAll(bool pred) {
+  return __spirv_GroupAll(group_scope<Group>::value, pred);
+}
+
+template <typename Group> bool GroupAny(bool pred) {
+  return __spirv_GroupAny(group_scope<Group>::value, pred);
+}
+
 // Broadcast with scalar local index
-template <__spv::Scope::Flag S, typename T, typename IdT>
+template <typename Group, typename T, typename IdT>
 detail::enable_if_t<std::is_integral<IdT>::value, T>
 GroupBroadcast(T x, IdT local_id) {
   using OCLT = detail::ConvertToOpenCLType_t<T>;
   using OCLIdT = detail::ConvertToOpenCLType_t<IdT>;
   OCLT ocl_x = detail::convertDataToType<T, OCLT>(x);
   OCLIdT ocl_id = detail::convertDataToType<IdT, OCLIdT>(local_id);
-  return __spirv_GroupBroadcast(S, ocl_x, ocl_id);
+  return __spirv_GroupBroadcast(group_scope<Group>::value, ocl_x, ocl_id);
 }
 
 // Broadcast with vector local index
-template <__spv::Scope::Flag S, typename T, int Dimensions>
+template <typename Group, typename T, int Dimensions>
 T GroupBroadcast(T x, id<Dimensions> local_id) {
   if (Dimensions == 1) {
-    return GroupBroadcast<S>(x, local_id[0]);
+    return GroupBroadcast<Group>(x, local_id[0]);
   }
   using IdT = vec<size_t, Dimensions>;
   using OCLT = detail::ConvertToOpenCLType_t<T>;
@@ -45,7 +66,7 @@ T GroupBroadcast(T x, id<Dimensions> local_id) {
   }
   OCLT ocl_x = detail::convertDataToType<T, OCLT>(x);
   OCLIdT ocl_id = detail::convertDataToType<IdT, OCLIdT>(vec_id);
-  return __spirv_GroupBroadcast(S, ocl_x, ocl_id);
+  return __spirv_GroupBroadcast(group_scope<Group>::value, ocl_x, ocl_id);
 }
 
 } // namespace spirv
diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index ad8fa67313d91..e2bd2c80ca209 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -14,6 +14,7 @@
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/group.hpp>
 #include <CL/sycl/intel/functional.hpp>
+#include <CL/sycl/intel/sub_group.hpp>
 
 #ifndef __DISABLE_SYCL_INTEL_GROUP_ALGORITHMS__
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -30,6 +31,32 @@ template <> inline size_t get_local_linear_range<group<2>>(group<2> g) {
 template <> inline size_t get_local_linear_range<group<3>>(group<3> g) {
   return g.get_local_range(0) * g.get_local_range(1) * g.get_local_range(2);
 }
+template <> inline
+size_t get_local_linear_range<intel::sub_group>(intel::sub_group g) {
+  return g.get_local_range()[0];
+}
+
+template <typename Group>
+typename Group::linear_id_type get_local_linear_id(Group g);
+
+#ifdef __SYCL_DEVICE_ONLY__
+#define __SYCL_GROUP_GET_LOCAL_LINEAR_ID(D)                                    \
+  template <>                                                                  \
+  group<D>::linear_id_type get_local_linear_id<group<D>>(group<D> g) {         \
+    nd_item<D> it = ::sycl::detail::Builder::getNDItem<D>();                   \
+    return it.get_local_linear_id();                                           \
+  }
+__SYCL_GROUP_GET_LOCAL_LINEAR_ID(1);
+__SYCL_GROUP_GET_LOCAL_LINEAR_ID(2);
+__SYCL_GROUP_GET_LOCAL_LINEAR_ID(3);
+#undef __SYCL_GROUP_GET_LOCAL_LINEAR_ID
+#endif // __SYCL_DEVICE_ONLY__
+
+template <>
+intel::sub_group::linear_id_type
+get_local_linear_id<intel::sub_group>(intel::sub_group g) {
+  return g.get_local_id()[0];
+}
 
 template <int Dimensions>
 id<Dimensions> linear_id_to_id(range<Dimensions>, size_t linear_id);
@@ -55,6 +82,10 @@ template <typename T> struct is_group : std::false_type {};
 template <int Dimensions>
 struct is_group<group<Dimensions>> : std::true_type {};
 
+template <typename T> struct is_sub_group : std::false_type {};
+
+template <> struct is_sub_group<intel::sub_group> : std::true_type {};
+
 template <typename T, class BinaryOperation> struct identity {};
 
 template <typename T, typename V> struct identity<T, intel::plus<V>> {
@@ -72,9 +103,7 @@ template <typename T, typename V> struct identity<T, intel::maximum<V>> {
 template <typename Group, typename Ptr, class Function>
 Function for_each(Group g, Ptr first, Ptr last, Function f) {
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
-  ptrdiff_t offset = it.get_local_linear_id();
+  ptrdiff_t offset = detail::get_local_linear_id(g);
   ptrdiff_t stride = detail::get_local_linear_range(g);
   for (Ptr p = first + offset; p < last; p += stride) {
     f(*p);
@@ -103,10 +132,12 @@ using EnableIfIsPointer =
     cl::sycl::detail::enable_if_t<cl::sycl::detail::is_pointer<Ptr>::value, T>;
 
 template <typename Group> bool all_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
-  return __spirv_GroupAll(__spv::Scope::Workgroup, pred);
+  return detail::spirv::GroupAll<Group>(pred);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
                       PI_INVALID_DEVICE);
@@ -115,8 +146,10 @@ template <typename Group> bool all_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool all_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   return all_of(g, pred(x));
 }
 
@@ -124,8 +157,10 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> all_of(Group g, Ptr first, Ptr last,
                                     Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   bool partial = true;
   detail::for_each(g, first, last, [&](const typename Ptr::element_type &x) {
     partial &= pred(x);
@@ -138,10 +173,12 @@ EnableIfIsPointer<Ptr, bool> all_of(Group g, Ptr first, Ptr last,
 }
 
 template <typename Group> bool any_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
-  return __spirv_GroupAny(__spv::Scope::Workgroup, pred);
+  return detail::spirv::GroupAny<Group>(pred);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
                       PI_INVALID_DEVICE);
@@ -150,8 +187,10 @@ template <typename Group> bool any_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool any_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   return any_of(g, pred(x));
 }
 
@@ -159,8 +198,10 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> any_of(Group g, Ptr first, Ptr last,
                                     Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   bool partial = false;
   detail::for_each(g, first, last, [&](const typename Ptr::element_type &x) {
     partial |= pred(x);
@@ -173,10 +214,12 @@ EnableIfIsPointer<Ptr, bool> any_of(Group g, Ptr first, Ptr last,
 }
 
 template <typename Group> bool none_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
-  return __spirv_GroupAll(__spv::Scope::Workgroup, not pred);
+  return detail::spirv::GroupAll<Group>(not pred);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
                       PI_INVALID_DEVICE);
@@ -185,8 +228,10 @@ template <typename Group> bool none_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool none_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   return none_of(g, pred(x));
 }
 
@@ -194,8 +239,10 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> none_of(Group g, Ptr first, Ptr last,
                                      Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   return not any_of(g, first, last, pred);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
@@ -206,10 +253,12 @@ EnableIfIsPointer<Ptr, bool> none_of(Group g, Ptr first, Ptr last,
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T> broadcast(Group g, T x,
                                         typename Group::id_type local_id) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
-  return detail::spirv::GroupBroadcast<__spv::Scope::Workgroup>(x, local_id);
+  return detail::spirv::GroupBroadcast<Group>(x, local_id);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
                       PI_INVALID_DEVICE);
@@ -219,8 +268,10 @@ EnableIfIsScalarArithmetic<T> broadcast(Group g, T x,
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T> broadcast(Group g, T x,
                                         typename Group::id_type local_id) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -236,8 +287,10 @@ EnableIfIsVectorArithmetic<T> broadcast(Group g, T x,
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T>
 broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
   return broadcast(
       g, x, detail::linear_id_to_id(g.get_local_range(), linear_local_id));
@@ -250,8 +303,10 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T>
 broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -266,8 +321,10 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
 
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T> broadcast(Group g, T x) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
   return broadcast(g, x, 0);
 #else
@@ -278,8 +335,10 @@ EnableIfIsScalarArithmetic<T> broadcast(Group g, T x) {
 
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T> broadcast(Group g, T x) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -294,14 +353,16 @@ EnableIfIsVectorArithmetic<T> broadcast(Group g, T x) {
 
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(x, x)), T>::value,
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::Reduce,
-                      __spv::Scope::Workgroup>(
+                      detail::spirv::group_scope<Group>::value>(
       typename detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
@@ -311,8 +372,10 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
 
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(x[0], x[0])),
                    typename T::element_type>::value,
@@ -327,8 +390,10 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> reduce(Group g, V x, T init,
                                      BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(init, x)), T>::value,
       "Result type of binary_op must match reduction accumulation type.");
@@ -343,8 +408,10 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, V x, T init,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> reduce(Group g, V x, T init,
                                      BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(init[0], x[0])),
                    typename T::element_type>::value,
@@ -364,8 +431,10 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, V x, T init,
 template <typename Group, typename Ptr, class BinaryOperation>
 EnableIfIsPointer<Ptr, typename Ptr::element_type>
 reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(*first, *first)),
                    typename Ptr::element_type>::value,
@@ -386,8 +455,10 @@ reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
 template <typename Group, typename Ptr, typename T, class BinaryOperation>
 EnableIfIsPointer<Ptr, T> reduce(Group g, Ptr first, Ptr last, T init,
                                  BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(
       std::is_same<decltype(binary_op(init, *first)), T>::value,
       "Result type of binary_op must match reduction accumulation type.");
@@ -407,13 +478,15 @@ EnableIfIsPointer<Ptr, T> reduce(Group g, Ptr first, Ptr last, T init,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
-                      __spv::Scope::Workgroup>(
+                      detail::spirv::group_scope<Group>::value>(
       typename detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
@@ -424,8 +497,10 @@ EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, T x,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
                              typename T::element_type>::value,
                 "Result type of binary_op must match scan accumulation type.");
@@ -439,8 +514,10 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, T x,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, V x, T init,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init[0], x[0])),
                              typename T::element_type>::value,
                 "Result type of binary_op must match scan accumulation type.");
@@ -454,18 +531,20 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, V x, T init,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, V x, T init,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      detail::Builder::getNDItem<Group::dimensions>();
-  if (it.get_local_linear_id() == 0) {
+  typename Group::linear_id_type local_linear_id =
+      detail::get_local_linear_id(g);
+  if (local_linear_id == 0) {
     x = binary_op(init, x);
   }
   T scan = exclusive_scan(g, x, binary_op);
-  if (it.get_local_linear_id() == 0) {
+  if (local_linear_id == 0) {
     scan = init;
   }
   return scan;
@@ -480,14 +559,14 @@ template <typename Group, typename InPtr, typename OutPtr, typename T,
 EnableIfIsPointer<InPtr, OutPtr>
 exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
                BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(*first, *first)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
-  ptrdiff_t offset = it.get_local_linear_id();
+  ptrdiff_t offset = detail::get_local_linear_id(g);
   ptrdiff_t stride = detail::get_local_linear_range(g);
   ptrdiff_t N = last - first;
   auto roundup = [=](const ptrdiff_t &v,
@@ -531,8 +610,10 @@ EnableIfIsPointer<InPtr, OutPtr> exclusive_scan(Group g, InPtr first,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> inclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
                              typename T::element_type>::value,
                 "Result type of binary_op must match scan accumulation type.");
@@ -546,13 +627,15 @@ EnableIfIsVectorArithmetic<T> inclusive_scan(Group g, T x,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> inclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::InclusiveScan,
-                      __spv::Scope::Workgroup>(
+                      detail::spirv::group_scope<Group>::value>(
       typename detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",
@@ -563,14 +646,14 @@ EnableIfIsScalarArithmetic<T> inclusive_scan(Group g, T x,
 template <typename Group, typename V, class BinaryOperation, typename T>
 EnableIfIsScalarArithmetic<T>
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      detail::Builder::getNDItem<Group::dimensions>();
-  if (it.get_local_linear_id() == 0) {
+  if (detail::get_local_linear_id(g) == 0) {
     x = binary_op(init, x);
   }
   return inclusive_scan(g, x, binary_op);
@@ -583,8 +666,10 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
 template <typename Group, typename V, class BinaryOperation, typename T>
 EnableIfIsVectorArithmetic<T>
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init[0], x[0])), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
   T result;
@@ -599,14 +684,14 @@ template <typename Group, typename InPtr, typename OutPtr,
 EnableIfIsPointer<InPtr, OutPtr>
 inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, *first)), T>::value,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
-  ptrdiff_t offset = it.get_local_linear_id();
+  ptrdiff_t offset = detail::get_local_linear_id(g);
   ptrdiff_t stride = detail::get_local_linear_range(g);
   ptrdiff_t N = last - first;
   auto roundup = [=](const ptrdiff_t &v,
diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 0c8d0c434b3e2..5c4321d937b0b 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -133,9 +133,9 @@ namespace intel {
 
 struct sub_group {
 
-  using id_type = id<1>;
-  using range_type = range<1>;
-  using linear_id_type = size_t;
+  typedef id<1> id_type;
+  typedef range<1> range_type;
+  typedef size_t linear_id_type;
   static constexpr int dimensions = 1;
 
   /* --- common interface members --- */
@@ -159,10 +159,16 @@ struct sub_group {
 
   /* --- vote / ballot functions --- */
 
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::any_of instead.")]]
+#endif
   bool any(bool predicate) const {
     return __spirv_GroupAny(__spv::Scope::Subgroup, predicate);
   }
 
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::all_of instead.")]]
+#endif
   bool all(bool predicate) const {
     return __spirv_GroupAll(__spv::Scope::Subgroup, predicate);
   }
@@ -174,11 +180,17 @@ struct sub_group {
   /* --- collectives --- */
 
   template <typename T>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::broadcast instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
     return detail::spirv::GroupBroadcast<__spv::Scope::Subgroup>(x, local_id);
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::reduce instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::Reduce,
                         __spv::Scope::Subgroup>(
@@ -186,11 +198,17 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::reduce instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
     return op(init, reduce(x, op));
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
                         __spv::Scope::Subgroup>(
@@ -198,6 +216,9 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
                                                BinaryOperation op) const {
     if (get_local_id().get(0) == 0) {
@@ -211,6 +232,9 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::InclusiveScan,
                         __spv::Scope::Subgroup>(
@@ -218,6 +242,9 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
+#if __cplusplus >= 201402L
+  [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
+#endif
   EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
                                                T init) const {
     if (get_local_id().get(0) == 0) {

From 3cd7c841f04b2930cc091730c6ace82fa8eab7f1 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Thu, 13 Feb 2020 14:31:58 -0800
Subject: [PATCH 03/17] [SYCL] Use algorithms interface in sub_group tests

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/test/sub_group/broadcast.cpp | 2 +-
 sycl/test/sub_group/reduce.cpp    | 4 ++--
 sycl/test/sub_group/scan.cpp      | 8 ++++----
 sycl/test/sub_group/vote.cpp      | 4 ++--
 4 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/sycl/test/sub_group/broadcast.cpp b/sycl/test/sub_group/broadcast.cpp
index 41e73b22fc8a3..399aaf025738d 100644
--- a/sycl/test/sub_group/broadcast.cpp
+++ b/sycl/test/sub_group/broadcast.cpp
@@ -30,7 +30,7 @@ template <typename T> void check(queue &Queue) {
         intel::sub_group SG = NdItem.get_sub_group();
         /*Broadcast GID of element with SGLID == SGID */
         syclacc[NdItem.get_global_id()] =
-            SG.broadcast<T>(NdItem.get_global_id(0), SG.get_group_id());
+            broadcast(SG, T(NdItem.get_global_id(0)), SG.get_group_id());
         if (NdItem.get_global_id(0) == 0)
           sgsizeacc[0] = SG.get_max_local_range()[0];
       });
diff --git a/sycl/test/sub_group/reduce.cpp b/sycl/test/sub_group/reduce.cpp
index 24d97cc276262..6581021f78c70 100644
--- a/sycl/test/sub_group/reduce.cpp
+++ b/sycl/test/sub_group/reduce.cpp
@@ -34,10 +34,10 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
             intel::sub_group sg = NdItem.get_sub_group();
             if (skip_init) {
               acc[NdItem.get_global_id(0)] =
-                  sg.reduce(T(NdItem.get_global_id(0)), op);
+                  reduce(sg, T(NdItem.get_global_id(0)), op);
             } else {
               acc[NdItem.get_global_id(0)] =
-                  sg.reduce(T(NdItem.get_global_id(0)), init, op);
+                  reduce(sg, T(NdItem.get_global_id(0)), init, op);
             }
           });
     });
diff --git a/sycl/test/sub_group/scan.cpp b/sycl/test/sub_group/scan.cpp
index bd3a653232127..27a109c429ba5 100644
--- a/sycl/test/sub_group/scan.cpp
+++ b/sycl/test/sub_group/scan.cpp
@@ -36,14 +36,14 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
             intel::sub_group sg = NdItem.get_sub_group();
             if (skip_init) {
               exacc[NdItem.get_global_id(0)] =
-                  sg.exclusive_scan(T(NdItem.get_global_id(0)), op);
+                  exclusive_scan(sg, T(NdItem.get_global_id(0)), op);
               inacc[NdItem.get_global_id(0)] =
-                  sg.inclusive_scan(T(NdItem.get_global_id(0)), op);
+                  inclusive_scan(sg, T(NdItem.get_global_id(0)), op);
             } else {
               exacc[NdItem.get_global_id(0)] =
-                  sg.exclusive_scan(T(NdItem.get_global_id(0)), init, op);
+                  exclusive_scan(sg, T(NdItem.get_global_id(0)), init, op);
               inacc[NdItem.get_global_id(0)] =
-                  sg.inclusive_scan(T(NdItem.get_global_id(0)), op, init);
+                  inclusive_scan(sg, T(NdItem.get_global_id(0)), op, init);
             }
           });
     });
diff --git a/sycl/test/sub_group/vote.cpp b/sycl/test/sub_group/vote.cpp
index 16d0059d86f4d..6f0b4fc68a435 100644
--- a/sycl/test/sub_group/vote.cpp
+++ b/sycl/test/sub_group/vote.cpp
@@ -51,12 +51,12 @@ void check(queue Queue, const int G, const int L, const int D, const int R) {
       cgh.parallel_for<class subgr>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         /* Set to 1 if any local ID in subgroup devided by D has remainder R */
-        if (SG.any(SG.get_local_id().get(0) % D == R)) {
+        if (any_of(SG, SG.get_local_id().get(0) % D == R)) {
           sganyacc[NdItem.get_global_id()] = 1;
         }
         /* Set to 1 if remainder of division of subgroup local ID by D is less
          * than R for all work items in subgroup */
-        if (SG.all(SG.get_local_id().get(0) % D < R)) {
+        if (all_of(SG, SG.get_local_id().get(0) % D < R)) {
           sgallacc[NdItem.get_global_id()] = 1;
         }
       });

From 56cb923ab56c995bf67cc4fd42ea73a3358817e4 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Tue, 17 Mar 2020 13:18:43 -0700
Subject: [PATCH 04/17] [SYCL] Convert typedef to using syntax

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/sub_group.hpp      | 6 +++---
 sycl/include/CL/sycl/intel/sub_group_host.hpp | 6 +++---
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 5c4321d937b0b..93951a76d053c 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -133,9 +133,9 @@ namespace intel {
 
 struct sub_group {
 
-  typedef id<1> id_type;
-  typedef range<1> range_type;
-  typedef size_t linear_id_type;
+  using id_type = id<1>;
+  using range_type = range<1>;
+  using linear_id_type = size_t;
   static constexpr int dimensions = 1;
 
   /* --- common interface members --- */
diff --git a/sycl/include/CL/sycl/intel/sub_group_host.hpp b/sycl/include/CL/sycl/intel/sub_group_host.hpp
index e24fadc616162..d805c42b8869c 100644
--- a/sycl/include/CL/sycl/intel/sub_group_host.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group_host.hpp
@@ -21,9 +21,9 @@ template <typename T, access::address_space Space> class multi_ptr;
 namespace intel {
 struct sub_group {
 
-  typedef id<1> id_type;
-  typedef range<1> range_type;
-  typedef size_t linear_id_type;
+  using id_type = id<1>;
+  using range_type = range<1>;
+  using linear_id_type = size_t;
   static constexpr int dimensions = 1;
 
   /* --- common interface members --- */

From 97cb465185458398fdb564ecfd7a9f05b6f25ff2 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Tue, 17 Mar 2020 17:00:50 -0400
Subject: [PATCH 05/17] [SYCL] Apply clang-format to sub-group tests

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/test/sub_group/broadcast.cpp | 6 ++++--
 sycl/test/sub_group/reduce.cpp    | 6 ++++--
 sycl/test/sub_group/scan.cpp      | 6 ++++--
 3 files changed, 12 insertions(+), 6 deletions(-)

diff --git a/sycl/test/sub_group/broadcast.cpp b/sycl/test/sub_group/broadcast.cpp
index 399aaf025738d..b67308363b66c 100644
--- a/sycl/test/sub_group/broadcast.cpp
+++ b/sycl/test/sub_group/broadcast.cpp
@@ -15,9 +15,11 @@
 
 #include "helper.hpp"
 #include <CL/sycl.hpp>
-template <typename T> class sycl_subgr;
+template <typename T>
+class sycl_subgr;
 using namespace cl::sycl;
-template <typename T> void check(queue &Queue) {
+template <typename T>
+void check(queue &Queue) {
   const int G = 240, L = 60;
   try {
     nd_range<1> NdRange(G, L);
diff --git a/sycl/test/sub_group/reduce.cpp b/sycl/test/sub_group/reduce.cpp
index 6581021f78c70..27e5baccd27ee 100644
--- a/sycl/test/sub_group/reduce.cpp
+++ b/sycl/test/sub_group/reduce.cpp
@@ -17,7 +17,8 @@
 #include "helper.hpp"
 #include <CL/sycl.hpp>
 
-template <typename T, class BinaryOperation> class sycl_subgr;
+template <typename T, class BinaryOperation>
+class sycl_subgr;
 
 using namespace cl::sycl;
 
@@ -67,7 +68,8 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
   }
 }
 
-template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
+template <typename T>
+void check(queue &Queue, size_t G = 240, size_t L = 60) {
   // limit data range for half to avoid rounding issues
   if (std::is_same<T, cl::sycl::half>::value) {
     G = 64;
diff --git a/sycl/test/sub_group/scan.cpp b/sycl/test/sub_group/scan.cpp
index 27a109c429ba5..3a61dfbcba4d9 100644
--- a/sycl/test/sub_group/scan.cpp
+++ b/sycl/test/sub_group/scan.cpp
@@ -18,7 +18,8 @@
 #include <CL/sycl.hpp>
 #include <limits>
 
-template <typename T, class BinaryOperation> class sycl_subgr;
+template <typename T, class BinaryOperation>
+class sycl_subgr;
 
 using namespace cl::sycl;
 
@@ -75,7 +76,8 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
   }
 }
 
-template <typename T> void check(queue &Queue, size_t G = 120, size_t L = 60) {
+template <typename T>
+void check(queue &Queue, size_t G = 120, size_t L = 60) {
   // limit data range for half to avoid rounding issues
   if (std::is_same<T, cl::sycl::half>::value) {
     G = 64;

From ef8e191acd9f55777af14800ebd4d27f2b2ce6ef Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Tue, 17 Mar 2020 17:03:37 -0400
Subject: [PATCH 06/17] [SYCL] Apply clang-format to group_algorithm.hpp

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/group_algorithm.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index e2bd2c80ca209..c931fc3c1298c 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -31,8 +31,8 @@ template <> inline size_t get_local_linear_range<group<2>>(group<2> g) {
 template <> inline size_t get_local_linear_range<group<3>>(group<3> g) {
   return g.get_local_range(0) * g.get_local_range(1) * g.get_local_range(2);
 }
-template <> inline
-size_t get_local_linear_range<intel::sub_group>(intel::sub_group g) {
+template <>
+inline size_t get_local_linear_range<intel::sub_group>(intel::sub_group g) {
   return g.get_local_range()[0];
 }
 

From 3e05d75b13dc4498d0cf0faa1efc07717dac15ba Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Tue, 17 Mar 2020 17:12:41 -0400
Subject: [PATCH 07/17] [SYCL] Fix leader algorithm for sub_group

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/group_algorithm.hpp | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index c931fc3c1298c..ea721acf36744 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -732,12 +732,12 @@ EnableIfIsPointer<InPtr, OutPtr> inclusive_scan(Group g, InPtr first,
 }
 
 template <typename Group> bool leader(Group g) {
-  static_assert(detail::is_group<Group>::value,
-                "Group algorithms only support the sycl::group class.");
+  static_assert(detail::is_group<Group>::value ||
+                    detail::is_sub_group<Group>::value,
+                "Group algorithms only support the sycl::group and "
+                "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
-  nd_item<Group::dimensions> it =
-      cl::sycl::detail::Builder::getNDItem<Group::dimensions>();
-  typename Group::linear_id_type linear_id = it.get_local_linear_id();
+  typename Group::linear_id_type linear_id = detail::get_local_linear_id(g);
   return (linear_id == 0);
 #else
   throw runtime_error("Group algorithms are not supported on host device.",

From 0dde5377bd47f12b2bfc6aa7d7d9fed9bc513919 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 18 Mar 2020 13:47:00 -0400
Subject: [PATCH 08/17] [SYCL] Mark specialization inline

Multiple definition errors again...

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/group_algorithm.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index ea721acf36744..8c44a28f642c6 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -53,7 +53,7 @@ __SYCL_GROUP_GET_LOCAL_LINEAR_ID(3);
 #endif // __SYCL_DEVICE_ONLY__
 
 template <>
-intel::sub_group::linear_id_type
+inline intel::sub_group::linear_id_type
 get_local_linear_id<intel::sub_group>(intel::sub_group g) {
   return g.get_local_id()[0];
 }

From 163fa5a417f5d363565ed19131dab3b8cff440f6 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 18 Mar 2020 17:48:01 -0400
Subject: [PATCH 09/17] [SYCL] Avoid inline namespace

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/group_algorithm.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index 8c44a28f642c6..adeb7edfb8372 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -43,7 +43,7 @@ typename Group::linear_id_type get_local_linear_id(Group g);
 #define __SYCL_GROUP_GET_LOCAL_LINEAR_ID(D)                                    \
   template <>                                                                  \
   group<D>::linear_id_type get_local_linear_id<group<D>>(group<D> g) {         \
-    nd_item<D> it = ::sycl::detail::Builder::getNDItem<D>();                   \
+    nd_item<D> it = cl::sycl::detail::Builder::getNDItem<D>();                 \
     return it.get_local_linear_id();                                           \
   }
 __SYCL_GROUP_GET_LOCAL_LINEAR_ID(1);

From d0200b0c0045f991617546b5ab4c5b72dd499347 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 18 Mar 2020 18:08:35 -0400
Subject: [PATCH 10/17] [SYCL] Simplify is_group and is_sub_group check

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 .../include/CL/sycl/intel/group_algorithm.hpp | 101 +++++++-----------
 1 file changed, 37 insertions(+), 64 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index adeb7edfb8372..0d1a7004c4edf 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -86,6 +86,11 @@ template <typename T> struct is_sub_group : std::false_type {};
 
 template <> struct is_sub_group<intel::sub_group> : std::true_type {};
 
+template <typename T>
+struct is_generic_group
+    : std::integral_constant<bool,
+                             is_group<T>::value || is_sub_group<T>::value> {};
+
 template <typename T, class BinaryOperation> struct identity {};
 
 template <typename T, typename V> struct identity<T, intel::plus<V>> {
@@ -132,8 +137,7 @@ using EnableIfIsPointer =
     cl::sycl::detail::enable_if_t<cl::sycl::detail::is_pointer<Ptr>::value, T>;
 
 template <typename Group> bool all_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -146,8 +150,7 @@ template <typename Group> bool all_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool all_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   return all_of(g, pred(x));
@@ -157,8 +160,7 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> all_of(Group g, Ptr first, Ptr last,
                                     Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   bool partial = true;
@@ -173,8 +175,7 @@ EnableIfIsPointer<Ptr, bool> all_of(Group g, Ptr first, Ptr last,
 }
 
 template <typename Group> bool any_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -187,8 +188,7 @@ template <typename Group> bool any_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool any_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   return any_of(g, pred(x));
@@ -198,8 +198,7 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> any_of(Group g, Ptr first, Ptr last,
                                     Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   bool partial = false;
@@ -214,8 +213,7 @@ EnableIfIsPointer<Ptr, bool> any_of(Group g, Ptr first, Ptr last,
 }
 
 template <typename Group> bool none_of(Group g, bool pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -228,8 +226,7 @@ template <typename Group> bool none_of(Group g, bool pred) {
 
 template <typename Group, typename T, class Predicate>
 bool none_of(Group g, T x, Predicate pred) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   return none_of(g, pred(x));
@@ -239,8 +236,7 @@ template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> none_of(Group g, Ptr first, Ptr last,
                                      Predicate pred) {
 #ifdef __SYCL_DEVICE_ONLY__
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   return not any_of(g, first, last, pred);
@@ -253,8 +249,7 @@ EnableIfIsPointer<Ptr, bool> none_of(Group g, Ptr first, Ptr last,
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T> broadcast(Group g, T x,
                                         typename Group::id_type local_id) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -268,8 +263,7 @@ EnableIfIsScalarArithmetic<T> broadcast(Group g, T x,
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T> broadcast(Group g, T x,
                                         typename Group::id_type local_id) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -287,8 +281,7 @@ EnableIfIsVectorArithmetic<T> broadcast(Group g, T x,
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T>
 broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -303,8 +296,7 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T>
 broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -321,8 +313,7 @@ broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
 
 template <typename Group, typename T>
 EnableIfIsScalarArithmetic<T> broadcast(Group g, T x) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -335,8 +326,7 @@ EnableIfIsScalarArithmetic<T> broadcast(Group g, T x) {
 
 template <typename Group, typename T>
 EnableIfIsVectorArithmetic<T> broadcast(Group g, T x) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__
@@ -353,8 +343,7 @@ EnableIfIsVectorArithmetic<T> broadcast(Group g, T x) {
 
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -372,8 +361,7 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
 
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -390,8 +378,7 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> reduce(Group g, V x, T init,
                                      BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -408,8 +395,7 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, V x, T init,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> reduce(Group g, V x, T init,
                                      BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -431,8 +417,7 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, V x, T init,
 template <typename Group, typename Ptr, class BinaryOperation>
 EnableIfIsPointer<Ptr, typename Ptr::element_type>
 reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -455,8 +440,7 @@ reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
 template <typename Group, typename Ptr, typename T, class BinaryOperation>
 EnableIfIsPointer<Ptr, T> reduce(Group g, Ptr first, Ptr last, T init,
                                  BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(
@@ -478,8 +462,7 @@ EnableIfIsPointer<Ptr, T> reduce(Group g, Ptr first, Ptr last, T init,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
@@ -497,8 +480,7 @@ EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, T x,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
@@ -514,8 +496,7 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, T x,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, V x, T init,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init[0], x[0])),
@@ -531,8 +512,7 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, V x, T init,
 template <typename Group, typename V, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, V x, T init,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
@@ -559,8 +539,7 @@ template <typename Group, typename InPtr, typename OutPtr, typename T,
 EnableIfIsPointer<InPtr, OutPtr>
 exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
                BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(*first, *first)), T>::value,
@@ -610,8 +589,7 @@ EnableIfIsPointer<InPtr, OutPtr> exclusive_scan(Group g, InPtr first,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsVectorArithmetic<T> inclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
@@ -627,8 +605,7 @@ EnableIfIsVectorArithmetic<T> inclusive_scan(Group g, T x,
 template <typename Group, typename T, class BinaryOperation>
 EnableIfIsScalarArithmetic<T> inclusive_scan(Group g, T x,
                                              BinaryOperation binary_op) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
@@ -646,8 +623,7 @@ EnableIfIsScalarArithmetic<T> inclusive_scan(Group g, T x,
 template <typename Group, typename V, class BinaryOperation, typename T>
 EnableIfIsScalarArithmetic<T>
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
@@ -666,8 +642,7 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
 template <typename Group, typename V, class BinaryOperation, typename T>
 EnableIfIsVectorArithmetic<T>
 inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init[0], x[0])), T>::value,
@@ -684,8 +659,7 @@ template <typename Group, typename InPtr, typename OutPtr,
 EnableIfIsPointer<InPtr, OutPtr>
 inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                BinaryOperation binary_op, T init) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
   static_assert(std::is_same<decltype(binary_op(init, *first)), T>::value,
@@ -732,8 +706,7 @@ EnableIfIsPointer<InPtr, OutPtr> inclusive_scan(Group g, InPtr first,
 }
 
 template <typename Group> bool leader(Group g) {
-  static_assert(detail::is_group<Group>::value ||
-                    detail::is_sub_group<Group>::value,
+  static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
 #ifdef __SYCL_DEVICE_ONLY__

From aa2fb0824094847790282ab5e89840a5bb1459ae Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 18 Mar 2020 18:09:35 -0400
Subject: [PATCH 11/17] [SYCL] Move #ifdef SYCL_DEVICE_ONLY

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/group_algorithm.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index 0d1a7004c4edf..0d53b7bd798ae 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -159,10 +159,10 @@ bool all_of(Group g, T x, Predicate pred) {
 template <typename Group, typename Ptr, class Predicate>
 EnableIfIsPointer<Ptr, bool> all_of(Group g, Ptr first, Ptr last,
                                     Predicate pred) {
-#ifdef __SYCL_DEVICE_ONLY__
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+#ifdef __SYCL_DEVICE_ONLY__
   bool partial = true;
   detail::for_each(g, first, last, [&](const typename Ptr::element_type &x) {
     partial &= pred(x);

From e86d27bf3c16a057486aaaf945f2cc599eea0ea1 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Mon, 23 Mar 2020 17:13:22 -0400
Subject: [PATCH 12/17] [SYCL] Add workaround for half types

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 .../include/CL/sycl/intel/group_algorithm.hpp | 124 +++++++++++++-----
 1 file changed, 93 insertions(+), 31 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/group_algorithm.hpp b/sycl/include/CL/sycl/intel/group_algorithm.hpp
index 0d53b7bd798ae..f4b59a2b4068e 100644
--- a/sycl/include/CL/sycl/intel/group_algorithm.hpp
+++ b/sycl/include/CL/sycl/intel/group_algorithm.hpp
@@ -346,8 +346,11 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(x, x)), T>::value,
+      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),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::Reduce,
@@ -364,9 +367,12 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, T x, BinaryOperation binary_op) {
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // 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,
+                   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 reduction accumulation type.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
@@ -381,8 +387,11 @@ EnableIfIsScalarArithmetic<T> reduce(Group g, V x, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init, x)), T>::value,
+      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),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return binary_op(init, reduce(g, x, binary_op));
@@ -398,9 +407,12 @@ EnableIfIsVectorArithmetic<T> reduce(Group g, V x, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // 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,
+                   typename T::element_type>::value ||
+          (std::is_same<T, half>::value &&
+           std::is_same<decltype(binary_op(init[0], x[0])), float>::value),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T result = init;
@@ -420,9 +432,12 @@ reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // FIXME: Do not special-case for half precision
   static_assert(
       std::is_same<decltype(binary_op(*first, *first)),
-                   typename Ptr::element_type>::value,
+                   typename Ptr::element_type>::value ||
+          (std::is_same<typename Ptr::element_type, half>::value &&
+           std::is_same<decltype(binary_op(*first, *first)), float>::value),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   typename Ptr::element_type partial =
@@ -443,8 +458,11 @@ EnableIfIsPointer<Ptr, T> reduce(Group g, Ptr first, Ptr last, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
+  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same<decltype(binary_op(init, *first)), T>::value,
+      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),
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T partial =
@@ -465,7 +483,10 @@ EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, T x,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
+  // 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),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
@@ -483,9 +504,13 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, T x,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
-                             typename T::element_type>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
     result[s] = exclusive_scan(g, x[s], binary_op);
@@ -499,9 +524,13 @@ EnableIfIsVectorArithmetic<T> exclusive_scan(Group g, V x, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(init[0], x[0])),
-                             typename T::element_type>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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),
+      "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], init[s], binary_op);
@@ -515,7 +544,10 @@ EnableIfIsScalarArithmetic<T> exclusive_scan(Group g, V x, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
+  // 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),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   typename Group::linear_id_type local_linear_id =
@@ -542,8 +574,12 @@ exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(*first, *first)), T>::value,
-                "Result type of binary_op must match scan accumulation 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),
+      "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = detail::get_local_linear_id(g);
   ptrdiff_t stride = detail::get_local_linear_range(g);
@@ -577,9 +613,13 @@ template <typename Group, typename InPtr, typename OutPtr,
 EnableIfIsPointer<InPtr, OutPtr> exclusive_scan(Group g, InPtr first,
                                                 InPtr last, OutPtr result,
                                                 BinaryOperation binary_op) {
-  static_assert(std::is_same<decltype(binary_op(*first, *first)),
-                             typename OutPtr::element_type>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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),
+      "Result type of binary_op must match scan accumulation type.");
   return exclusive_scan(
       g, first, last, result,
       detail::identity<typename OutPtr::element_type, BinaryOperation>::value,
@@ -592,9 +632,13 @@ EnableIfIsVectorArithmetic<T> inclusive_scan(Group g, T x,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(x[0], x[0])),
-                             typename T::element_type>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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.");
   T result;
   for (int s = 0; s < x.get_size(); ++s) {
     result[s] = inclusive_scan(g, x[s], binary_op);
@@ -608,7 +652,10 @@ EnableIfIsScalarArithmetic<T> inclusive_scan(Group g, T x,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(x, x)), T>::value,
+  // 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),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return detail::calc<T, __spv::GroupOperation::InclusiveScan,
@@ -626,7 +673,10 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(init, x)), T>::value,
+  // 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),
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   if (detail::get_local_linear_id(g) == 0) {
@@ -645,8 +695,12 @@ inclusive_scan(Group g, V x, BinaryOperation binary_op, T init) {
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(init[0], x[0])), T>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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),
+      "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, init[s]);
@@ -662,8 +716,12 @@ inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
   static_assert(detail::is_generic_group<Group>::value,
                 "Group algorithms only support the sycl::group and "
                 "intel::sub_group class.");
-  static_assert(std::is_same<decltype(binary_op(init, *first)), T>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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),
+      "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = detail::get_local_linear_id(g);
   ptrdiff_t stride = detail::get_local_linear_range(g);
@@ -697,9 +755,13 @@ template <typename Group, typename InPtr, typename OutPtr,
 EnableIfIsPointer<InPtr, OutPtr> inclusive_scan(Group g, InPtr first,
                                                 InPtr last, OutPtr result,
                                                 BinaryOperation binary_op) {
-  static_assert(std::is_same<decltype(binary_op(*first, *first)),
-                             typename OutPtr::element_type>::value,
-                "Result type of binary_op must match scan accumulation type.");
+  // 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),
+      "Result type of binary_op must match scan accumulation type.");
   return inclusive_scan(
       g, first, last, result, binary_op,
       detail::identity<typename OutPtr::element_type, BinaryOperation>::value);

From 1378729ffcfc0ec35f2420c65f0440de4de18737 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Wed, 25 Mar 2020 07:55:58 -0700
Subject: [PATCH 13/17] [SYCL] Re-run clang-format on sub_group.hpp

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/sub_group.hpp | 29 +++++++++++++++---------
 1 file changed, 18 insertions(+), 11 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 93951a76d053c..68bcf11d2632e 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -162,14 +162,16 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::any_of instead.")]]
 #endif
-  bool any(bool predicate) const {
+  bool
+  any(bool predicate) const {
     return __spirv_GroupAny(__spv::Scope::Subgroup, predicate);
   }
 
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::all_of instead.")]]
 #endif
-  bool all(bool predicate) const {
+  bool
+  all(bool predicate) const {
     return __spirv_GroupAll(__spv::Scope::Subgroup, predicate);
   }
 
@@ -183,7 +185,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::broadcast instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
+  EnableIfIsScalarArithmetic<T>
+  broadcast(T x, id<1> local_id) const {
     return detail::spirv::GroupBroadcast<__spv::Scope::Subgroup>(x, local_id);
   }
 
@@ -191,7 +194,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::reduce instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T>
+  reduce(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::Reduce,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -201,7 +205,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::reduce instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T>
+  reduce(T x, T init, BinaryOperation op) const {
     return op(init, reduce(x, op));
   }
 
@@ -209,7 +214,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T>
+  exclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -219,8 +225,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
-                                               BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T>
+  exclusive_scan(T x, T init, BinaryOperation op) const {
     if (get_local_id().get(0) == 0) {
       x = op(init, x);
     }
@@ -235,7 +241,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T>
+  inclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::InclusiveScan,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -245,8 +252,8 @@ struct sub_group {
 #if __cplusplus >= 201402L
   [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
 #endif
-  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
-                                               T init) const {
+  EnableIfIsScalarArithmetic<T>
+  inclusive_scan(T x, BinaryOperation op, T init) const {
     if (get_local_id().get(0) == 0) {
       x = op(init, x);
     }

From 4723889b62a802f08250257645420450e6ded2a4 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Fri, 27 Mar 2020 07:07:40 -0700
Subject: [PATCH 14/17] [SYCL] Refactor __SYCL_DEPRECATED__ attribute

Takes message as a string.

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/detail/defines.hpp | 6 ++----
 sycl/include/CL/sycl/ordered_queue.hpp  | 4 +---
 2 files changed, 3 insertions(+), 7 deletions(-)

diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp
index 29dc311fe6808..a410517eec94a 100644
--- a/sycl/include/CL/sycl/detail/defines.hpp
+++ b/sycl/include/CL/sycl/detail/defines.hpp
@@ -33,11 +33,9 @@
 #endif
 
 #if __cplusplus >= 201402
-#define __SYCL_DEPRECATED__                                                    \
-  [[deprecated("Replaced by in_order queue property")]]
+#define __SYCL_DEPRECATED__(message) [[deprecated(message)]]
 #elif !defined _MSC_VER
-#define __SYCL_DEPRECATED__                                                    \
-  __attribute__((deprecated("Replaced by in_order queue property")))
+#define __SYCL_DEPRECATED__(message) __attribute__((deprecated(message)))
 #else
 #define __SYCL_DEPRECATED__
 #endif
diff --git a/sycl/include/CL/sycl/ordered_queue.hpp b/sycl/include/CL/sycl/ordered_queue.hpp
index 240d780645e8b..eda3b48e18f66 100644
--- a/sycl/include/CL/sycl/ordered_queue.hpp
+++ b/sycl/include/CL/sycl/ordered_queue.hpp
@@ -28,7 +28,7 @@ namespace detail {
 class queue_impl;
 }
 
-class __SYCL_DEPRECATED__ ordered_queue {
+class __SYCL_DEPRECATED__("Replaced by in_order queue property") ordered_queue {
 
 public:
   explicit ordered_queue(const property_list &propList = {})
@@ -257,8 +257,6 @@ class __SYCL_DEPRECATED__ ordered_queue {
                     const detail::code_location &CodeLoc);
 };
 
-#undef __SYCL_DEPRECATED__
-
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 

From 38f290664da5e14e0634dba6920166e7b421b8ee Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Fri, 27 Mar 2020 07:10:57 -0700
Subject: [PATCH 15/17] [SYCL] Adopt __SYCL_DEPRECATED__ in sub_group.hpp

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/sub_group.hpp | 47 ++++++++++--------------
 1 file changed, 20 insertions(+), 27 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 68bcf11d2632e..4f6baa9be912d 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -160,18 +160,16 @@ struct sub_group {
   /* --- vote / ballot functions --- */
 
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::any_of instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::any_of instead.")
 #endif
-  bool
-  any(bool predicate) const {
+  bool any(bool predicate) const {
     return __spirv_GroupAny(__spv::Scope::Subgroup, predicate);
   }
 
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::all_of instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::all_of instead.")
 #endif
-  bool
-  all(bool predicate) const {
+  bool all(bool predicate) const {
     return __spirv_GroupAll(__spv::Scope::Subgroup, predicate);
   }
 
@@ -183,19 +181,17 @@ struct sub_group {
 
   template <typename T>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::broadcast instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::broadcast instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  broadcast(T x, id<1> local_id) const {
+  EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
     return detail::spirv::GroupBroadcast<__spv::Scope::Subgroup>(x, local_id);
   }
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::reduce instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  reduce(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::Reduce,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -203,19 +199,17 @@ struct sub_group {
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::reduce instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  reduce(T x, T init, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
     return op(init, reduce(x, op));
   }
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  exclusive_scan(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -223,10 +217,10 @@ struct sub_group {
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::exclusive_scan instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  exclusive_scan(T x, T init, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
+                                               BinaryOperation op) const {
     if (get_local_id().get(0) == 0) {
       x = op(init, x);
     }
@@ -239,10 +233,9 @@ struct sub_group {
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  inclusive_scan(T x, BinaryOperation op) const {
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::InclusiveScan,
                         __spv::Scope::Subgroup>(
         typename detail::GroupOpTag<T>::type(), x, op);
@@ -250,10 +243,10 @@ struct sub_group {
 
   template <typename T, class BinaryOperation>
 #if __cplusplus >= 201402L
-  [[deprecated("Use sycl::intel::inclusive_scan instead.")]]
+  __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.")
 #endif
-  EnableIfIsScalarArithmetic<T>
-  inclusive_scan(T x, BinaryOperation op, T init) const {
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
+                                               T init) const {
     if (get_local_id().get(0) == 0) {
       x = op(init, x);
     }

From d6efc0111b1123a2e27d3836c2714deb44653ba7 Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Fri, 27 Mar 2020 11:12:28 -0400
Subject: [PATCH 16/17] [SYCL] Fix __SYCL_DEPRECATED__ for MSVC

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/detail/defines.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/CL/sycl/detail/defines.hpp b/sycl/include/CL/sycl/detail/defines.hpp
index a410517eec94a..c778ce9603fc9 100644
--- a/sycl/include/CL/sycl/detail/defines.hpp
+++ b/sycl/include/CL/sycl/detail/defines.hpp
@@ -37,5 +37,5 @@
 #elif !defined _MSC_VER
 #define __SYCL_DEPRECATED__(message) __attribute__((deprecated(message)))
 #else
-#define __SYCL_DEPRECATED__
+#define __SYCL_DEPRECATED__(message)
 #endif

From f65db727d514db4ffbd341e31f968db81190107a Mon Sep 17 00:00:00 2001
From: John Pennycook <john.pennycook@intel.com>
Date: Fri, 27 Mar 2020 12:42:58 -0400
Subject: [PATCH 17/17] [SYCL] Remove unnecessary C++14 guards

Signed-off-by: John Pennycook <john.pennycook@intel.com>
---
 sycl/include/CL/sycl/intel/sub_group.hpp | 18 ------------------
 1 file changed, 18 deletions(-)

diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 4f6baa9be912d..7d610b7983f50 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -159,16 +159,12 @@ struct sub_group {
 
   /* --- vote / ballot functions --- */
 
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::any_of instead.")
-#endif
   bool any(bool predicate) const {
     return __spirv_GroupAny(__spv::Scope::Subgroup, predicate);
   }
 
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::all_of instead.")
-#endif
   bool all(bool predicate) const {
     return __spirv_GroupAll(__spv::Scope::Subgroup, predicate);
   }
@@ -180,17 +176,13 @@ struct sub_group {
   /* --- collectives --- */
 
   template <typename T>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::broadcast instead.")
-#endif
   EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
     return detail::spirv::GroupBroadcast<__spv::Scope::Subgroup>(x, local_id);
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.")
-#endif
   EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::Reduce,
                         __spv::Scope::Subgroup>(
@@ -198,17 +190,13 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::reduce instead.")
-#endif
   EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
     return op(init, reduce(x, op));
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.")
-#endif
   EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::ExclusiveScan,
                         __spv::Scope::Subgroup>(
@@ -216,9 +204,7 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::exclusive_scan instead.")
-#endif
   EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
                                                BinaryOperation op) const {
     if (get_local_id().get(0) == 0) {
@@ -232,9 +218,7 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.")
-#endif
   EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
     return detail::calc<T, __spv::GroupOperation::InclusiveScan,
                         __spv::Scope::Subgroup>(
@@ -242,9 +226,7 @@ struct sub_group {
   }
 
   template <typename T, class BinaryOperation>
-#if __cplusplus >= 201402L
   __SYCL_DEPRECATED__("Use sycl::intel::inclusive_scan instead.")
-#endif
   EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
                                                T init) const {
     if (get_local_id().get(0) == 0) {