[SYCL] Clean up sub-group class API (#12945)

This PR cleans up the API of the sub-group class by either removing
functions that do not appear in the spec or marking them deprecated. The
only exception is the set of load/store functions, which after some
discussion with the spec team, it was suggested to leave them as is for
now.

Author: lbushi25

sycl/include/sycl/sub_group.hpp

@@ -219,8 +219,9 @@ struct sub_group {
 
   /* --- one-input shuffles --- */
   /* indices in [0 , sub_group size) */
-
-  template <typename T> T shuffle(T x, id_type local_id) const {
+  template <typename T>
+  __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.")
+  T shuffle(T x, id_type local_id) const {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::spirv::SubgroupShuffle(x, local_id);
 #else
@@ -231,7 +232,9 @@ struct sub_group {
 #endif
   }
 
-  template <typename T> T shuffle_down(T x, uint32_t delta) const {
+  template <typename T>
+  __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.")
+  T shuffle_down(T x, uint32_t delta) const {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::spirv::SubgroupShuffleDown(x, delta);
 #else
@@ -242,7 +245,9 @@ struct sub_group {
 #endif
   }
 
-  template <typename T> T shuffle_up(T x, uint32_t delta) const {
+  template <typename T>
+  __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.")
+  T shuffle_up(T x, uint32_t delta) const {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::spirv::SubgroupShuffleUp(x, delta);
 #else
@@ -253,7 +258,9 @@ struct sub_group {
 #endif
   }
 
-  template <typename T> T shuffle_xor(T x, id_type value) const {
+  template <typename T>
+  __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.")
+  T shuffle_xor(T x, id_type value) const {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::spirv::SubgroupShuffleXor(x, value);
 #else
@@ -358,7 +365,7 @@ struct sub_group {
     }
     return res;
   }
-#else // __NVPTX__ || __AMDGCN__
+#else  // __NVPTX__ || __AMDGCN__
   template <int N, typename CVT, access::address_space Space,
             access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
   std::enable_if_t<
@@ -608,6 +615,7 @@ struct sub_group {
   }
 
   /* --- synchronization functions --- */
+  __SYCL_DEPRECATED("Sub-group barrier with no arguments is deprecated.")
   void barrier() const {
 #ifdef __SYCL_DEVICE_ONLY__
     __spirv_ControlBarrier(
@@ -636,6 +644,7 @@ struct sub_group {
 #endif
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES__
   /* --- deprecated collective functions --- */
   template <typename T>
   __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
@@ -752,6 +761,7 @@ struct sub_group {
                           "Sub-groups are not supported on host.");
 #endif
   }
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES__
 
   linear_id_type get_group_linear_range() const {
 #ifdef __SYCL_DEVICE_ONLY__