[SYCL] Remove two-input sub-group shuffles (#2614)

These functions were deprecated in bea6aa2.

Removing them now allows for a simpler implementation of the interface to
SPIR-V that is more closely tied to supported functionality.  These changes
will make it simpler to implement sub-group shuffles for other backends.

Signed-off-by: John Pennycook <[email protected]>

Author: Pennycook

sycl/include/CL/sycl/ONEAPI/sub_group.hpp

@@ -191,7 +191,7 @@ struct sub_group {
 
   template <typename T> T shuffle_down(T x, uint32_t delta) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return sycl::detail::spirv::SubgroupShuffleDown(x, x, delta);
+    return sycl::detail::spirv::SubgroupShuffleDown(x, delta);
 #else
     (void)x;
     (void)delta;
@@ -202,7 +202,7 @@ struct sub_group {
 
   template <typename T> T shuffle_up(T x, uint32_t delta) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    return sycl::detail::spirv::SubgroupShuffleUp(x, x, delta);
+    return sycl::detail::spirv::SubgroupShuffleUp(x, delta);
 #else
     (void)x;
     (void)delta;
@@ -222,52 +222,6 @@ struct sub_group {
 #endif
   }
 
-  /* --- two-input shuffles --- */
-  /* indices in [0 , 2 * sub_group size) */
-
-  template <typename T>
-  __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.")
-  T shuffle(T x, T y, id_type local_id) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return sycl::detail::spirv::SubgroupShuffleDown(
-        x, y, (local_id - get_local_id()).get(0));
-#else
-    (void)x;
-    (void)y;
-    (void)local_id;
-    throw runtime_error("Sub-groups are not supported on host device.",
-                        PI_INVALID_DEVICE);
-#endif
-  }
-
-  template <typename T>
-  __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.")
-  T shuffle_down(T current, T next, uint32_t delta) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return sycl::detail::spirv::SubgroupShuffleDown(current, next, delta);
-#else
-    (void)current;
-    (void)next;
-    (void)delta;
-    throw runtime_error("Sub-groups are not supported on host device.",
-                        PI_INVALID_DEVICE);
-#endif
-  }
-
-  template <typename T>
-  __SYCL_DEPRECATED("Two-input sub-group shuffles are deprecated.")
-  T shuffle_up(T previous, T current, uint32_t delta) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return sycl::detail::spirv::SubgroupShuffleUp(previous, current, delta);
-#else
-    (void)previous;
-    (void)current;
-    (void)delta;
-    throw runtime_error("Sub-groups are not supported on host device.",
-                        PI_INVALID_DEVICE);
-#endif
-  }
-
   /* --- sub_group load/stores --- */
   /* these can map to SIMD or block read/write hardware where available */
 

sycl/include/CL/sycl/detail/spirv.hpp

@@ -443,16 +443,16 @@ EnableIfNativeShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
 }
 
 template <typename T>
-EnableIfNativeShuffle<T> SubgroupShuffleDown(T x, T y, id<1> local_id) {
+EnableIfNativeShuffle<T> SubgroupShuffleDown(T x, id<1> local_id) {
   using OCLT = detail::ConvertToOpenCLType_t<T>;
   return __spirv_SubgroupShuffleDownINTEL(
-      OCLT(x), OCLT(y), static_cast<uint32_t>(local_id.get(0)));
+      OCLT(x), OCLT(x), static_cast<uint32_t>(local_id.get(0)));
 }
 
 template <typename T>
-EnableIfNativeShuffle<T> SubgroupShuffleUp(T x, T y, id<1> local_id) {
+EnableIfNativeShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   using OCLT = detail::ConvertToOpenCLType_t<T>;
-  return __spirv_SubgroupShuffleUpINTEL(OCLT(x), OCLT(y),
+  return __spirv_SubgroupShuffleUpINTEL(OCLT(x), OCLT(x),
                                         static_cast<uint32_t>(local_id.get(0)));
 }
 
@@ -488,22 +488,20 @@ EnableIfBitcastShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
 }
 
 template <typename T>
-EnableIfBitcastShuffle<T> SubgroupShuffleDown(T x, T y, id<1> local_id) {
+EnableIfBitcastShuffle<T> SubgroupShuffleDown(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
   auto ShuffleX = detail::bit_cast<ShuffleT>(x);
-  auto ShuffleY = detail::bit_cast<ShuffleT>(y);
   ShuffleT Result = __spirv_SubgroupShuffleDownINTEL(
-      ShuffleX, ShuffleY, static_cast<uint32_t>(local_id.get(0)));
+      ShuffleX, ShuffleX, static_cast<uint32_t>(local_id.get(0)));
   return detail::bit_cast<T>(Result);
 }
 
 template <typename T>
-EnableIfBitcastShuffle<T> SubgroupShuffleUp(T x, T y, id<1> local_id) {
+EnableIfBitcastShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   using ShuffleT = ConvertToNativeShuffleType_t<T>;
   auto ShuffleX = detail::bit_cast<ShuffleT>(x);
-  auto ShuffleY = detail::bit_cast<ShuffleT>(y);
   ShuffleT Result = __spirv_SubgroupShuffleUpINTEL(
-      ShuffleX, ShuffleY, static_cast<uint32_t>(local_id.get(0)));
+      ShuffleX, ShuffleX, static_cast<uint32_t>(local_id.get(0)));
   return detail::bit_cast<T>(Result);
 }
 
@@ -550,33 +548,29 @@ EnableIfGenericShuffle<T> SubgroupShuffleXor(T x, id<1> local_id) {
 }
 
 template <typename T>
-EnableIfGenericShuffle<T> SubgroupShuffleDown(T x, T y, id<1> local_id) {
+EnableIfGenericShuffle<T> SubgroupShuffleDown(T x, id<1> local_id) {
   T Result;
   char *XBytes = reinterpret_cast<char *>(&x);
-  char *YBytes = reinterpret_cast<char *>(&y);
   char *ResultBytes = reinterpret_cast<char *>(&Result);
   auto ShuffleBytes = [=](size_t Offset, size_t Size) {
-    uint64_t ShuffleX, ShuffleY, ShuffleResult;
+    uint64_t ShuffleX, ShuffleResult;
     detail::memcpy(&ShuffleX, XBytes + Offset, Size);
-    detail::memcpy(&ShuffleY, YBytes + Offset, Size);
-    ShuffleResult = SubgroupShuffleDown(ShuffleX, ShuffleY, local_id);
+    ShuffleResult = SubgroupShuffleDown(ShuffleX, local_id);
     detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
   };
   GenericCall<T>(ShuffleBytes);
   return Result;
 }
 
 template <typename T>
-EnableIfGenericShuffle<T> SubgroupShuffleUp(T x, T y, id<1> local_id) {
+EnableIfGenericShuffle<T> SubgroupShuffleUp(T x, id<1> local_id) {
   T Result;
   char *XBytes = reinterpret_cast<char *>(&x);
-  char *YBytes = reinterpret_cast<char *>(&y);
   char *ResultBytes = reinterpret_cast<char *>(&Result);
   auto ShuffleBytes = [=](size_t Offset, size_t Size) {
-    uint64_t ShuffleX, ShuffleY, ShuffleResult;
+    uint64_t ShuffleX, ShuffleResult;
     detail::memcpy(&ShuffleX, XBytes + Offset, Size);
-    detail::memcpy(&ShuffleY, YBytes + Offset, Size);
-    ShuffleResult = SubgroupShuffleUp(ShuffleX, ShuffleY, local_id);
+    ShuffleResult = SubgroupShuffleUp(ShuffleX, local_id);
     detail::memcpy(ResultBytes + Offset, &ShuffleResult, Size);
   };
   GenericCall<T>(ShuffleBytes);