[NFC][SYCL] fix LIT tests

- Remove sub_group/common_ocl.cpp because it duplicates
  sub_group/common.cpp and directly use OpenCL API that causes
  instability on some configurations.
- Fix sub_group/shuffle*.cpp tests to align with  shuffle_xor
  restrictions mentioned in spec: "If the result of the XOR is
  greater than max_sub_group_size then it is considered out-of-range"

Author: vladimirlaz

sycl/test/sub_group/common_ocl.cpp

@@ -8,8 +8,7 @@
 
 #include "helper.hpp"
 #include <CL/sycl.hpp>
-template <typename T, int N>
-class sycl_subgr;
+template <typename T, int N> class sycl_subgr;
 
 using namespace cl::sycl;
 
@@ -66,8 +65,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
         acc_up[NdItem.get_global_id()] = SG.shuffle_up(vwggid, sgid);
         /* Save GID+SGID */
         acc_down[NdItem.get_global_id()] = SG.shuffle_down(vwggid, sgid);
-        /* Save GID XOR SGID */
-        acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(vwggid, sgid);
+        /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
+        acc_xor[NdItem.get_global_id()] =
+            SG.shuffle_xor(vwggid, sgid % SG.get_max_local_range()[0]);
       });
     });
     auto acc = buf.template get_access<access::mode::read_write>();
@@ -81,12 +81,18 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
 
     size_t sg_size = sgsizeacc[0];
     int SGid = 0;
+    int SGLid = 0;
+    int SGBeginGid = 0;
     for (int j = 0; j < G; j++) {
       if (j % L % sg_size == 0) {
         SGid++;
+        SGLid = 0;
+        SGBeginGid = j;
       }
       if (j % L == 0) {
         SGid = 0;
+        SGLid = 0;
+        SGBeginGid = j;
       }
       /*GID of middle element in every subgroup*/
       exit_if_not_equal_vec<T, N>(
@@ -115,17 +121,19 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
           exit_if_not_equal_vec(acc2_up[j], vec<T, N>(j - SGid + sg_size),
                                 "shuffle2_up");
       }
-      /* GID XOR SGID */
-      exit_if_not_equal_vec(acc_xor[j], vec<T, N>(j ^ SGid), "shuffle_xor");
+      /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
+      exit_if_not_equal_vec(acc_xor[j],
+                            vec<T, N>(SGBeginGid + (SGLid ^ (SGid % sg_size))),
+                            "shuffle_xor");
+      SGLid++;
     }
   } catch (exception e) {
     std::cout << "SYCL exception caught: " << e.what();
     exit(1);
   }
 }
 
-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) {
   try {
     nd_range<1> NdRange(G, L);
     buffer<T> buf2(G);
@@ -171,8 +179,9 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
         acc_up[NdItem.get_global_id()] = SG.shuffle_up<T>(wggid, sgid);
         /* Save GID+SGID */
         acc_down[NdItem.get_global_id()] = SG.shuffle_down<T>(wggid, sgid);
-        /* Save GID XOR SGID */
-        acc_xor[NdItem.get_global_id()] = SG.shuffle_xor<T>(wggid, sgid);
+        /* Save GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
+        acc_xor[NdItem.get_global_id()] =
+            SG.shuffle_xor<T>(wggid, sgid % SG.get_max_local_range()[0]);
       });
     });
     auto acc = buf.template get_access<access::mode::read_write>();
@@ -186,13 +195,20 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
 
     size_t sg_size = sgsizeacc[0];
     int SGid = 0;
+    int SGLid = 0;
+    int SGBeginGid = 0;
     for (int j = 0; j < G; j++) {
       if (j % L % sg_size == 0) {
         SGid++;
+        SGLid = 0;
+        SGBeginGid = j;
       }
       if (j % L == 0) {
         SGid = 0;
+        SGLid = 0;
+        SGBeginGid = j;
       }
+
       /*GID of middle element in every subgroup*/
       exit_if_not_equal<T>(acc[j], j / L * L + SGid * sg_size + sg_size / 2,
                            "shuffle");
@@ -215,8 +231,10 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
         if (j % L - SGid + sg_size < L) /* Do not go out  LG*/
           exit_if_not_equal<T>(acc2_up[j], j - SGid + sg_size, "shuffle2_up");
       }
-      /* GID XOR SGID */
-      exit_if_not_equal<T>(acc_xor[j], j ^ SGid, "shuffle_xor");
+      /* Value GID with SGLID = ( SGLID XOR SGID ) % SGMaxSize */
+      exit_if_not_equal<T>(acc_xor[j], SGBeginGid + (SGLid ^ (SGid % sg_size)),
+                           "shuffle_xor");
+      SGLid++;
     }
   } catch (exception e) {
     std::cout << "SYCL exception caught: " << e.what();