From a85d53940ab355ab8b73a444f417be67e7d461e2 Mon Sep 17 00:00:00 2001
From: "Romanov, Vlad" <vlad.romanov@intel.com>
Date: Thu, 17 Nov 2022 11:20:46 -0800
Subject: [PATCH] [SYCL] Disable sort over sub-group for CUDA and HIP BEs

CUDA and HIP BE do not support intel::reqd_sub_group_size, so it's not
possible to control sub-group size which is needed to get the size of
temporary memory for sort algorithm.
---
 SYCL/GroupAlgorithm/SYCL2020/sort.cpp | 14 ++++++++++----
 1 file changed, 10 insertions(+), 4 deletions(-)

diff --git a/SYCL/GroupAlgorithm/SYCL2020/sort.cpp b/SYCL/GroupAlgorithm/SYCL2020/sort.cpp
index 72614090a7..aff683c9a3 100644
--- a/SYCL/GroupAlgorithm/SYCL2020/sort.cpp
+++ b/SYCL/GroupAlgorithm/SYCL2020/sort.cpp
@@ -21,7 +21,7 @@
 //
 // TODO: Test global memory for temporary storage
 // TODO: Consider using USM instead of buffers
-//
+// TODO: Add support for sorting over workgroup for CUDA and HIP BE
 
 #include <sycl/sycl.hpp>
 
@@ -312,12 +312,18 @@ template <class T> void RunOverType(sycl::queue &Q, size_t DataSize) {
     RunSortOVerGroup<UseGroupT::WorkGroup, 1>(Q, Data, Comparator);
     RunSortOVerGroup<UseGroupT::WorkGroup, 2>(Q, Data, Comparator);
 
-    RunSortOVerGroup<UseGroupT::SubGroup, 1>(Q, Data, Comparator);
-    RunSortOVerGroup<UseGroupT::SubGroup, 2>(Q, Data, Comparator);
-
     RunJointSort<UseGroupT::WorkGroup, 1>(Q, Data, Comparator);
     RunJointSort<UseGroupT::WorkGroup, 2>(Q, Data, Comparator);
 
+    if (Q.get_backend() == sycl::backend::ext_oneapi_cuda ||
+        Q.get_backend() == sycl::backend::ext_oneapi_hip) {
+      std::cout << "Note! Skipping sub group testing on CUDA BE" << std::endl;
+      return;
+    }
+
+    RunSortOVerGroup<UseGroupT::SubGroup, 1>(Q, Data, Comparator);
+    RunSortOVerGroup<UseGroupT::SubGroup, 2>(Q, Data, Comparator);
+
     RunJointSort<UseGroupT::SubGroup, 1>(Q, Data, Comparator);
     RunJointSort<UseGroupT::SubGroup, 2>(Q, Data, Comparator);
   };