[SYCL][HIP] Improve test for barrier and enable it for HIP

SYCL/Basic/barrier_order.cpp

@@ -1,4 +1,3 @@
-// UNSUPPORTED: hip
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
@@ -11,28 +10,57 @@ int main() {
   sycl::device dev{sycl::default_selector_v};
   sycl::queue q{dev};
 
-  int *x = sycl::malloc_shared<int>(1, q);
-  int *y = sycl::malloc_shared<int>(1, q);
+  unsigned long long *x = sycl::malloc_shared<unsigned long long>(1, q);
+
+  int error = 0;
+
+  // test barrier without arguments
   *x = 0;
-  *y = 0;
+  for (int i = 0; i < 64; i++) {
+    q.single_task([=]() {
+      // do some busywork
+      volatile float y = *x;
+      for (int j = 0; j < 100; j++) {
+        y = sycl::cos(y);
+      }
+      // update the value
+      *x *= 2;
+    });
+    q.ext_oneapi_submit_barrier();
+    q.single_task([=]() { *x += 1; });
+    q.ext_oneapi_submit_barrier();
+  }
 
-  q.single_task<class kernel1>([=] { *x = 1; });
+  std::cout << std::bitset<8 * sizeof(unsigned long long)>(*x) << std::endl;
 
-  q.ext_oneapi_submit_barrier();
+  q.wait_and_throw();
+  error |= (*x != (unsigned long long)-1);
 
-  q.single_task<class kernel2>([=] {
-    if (*x == 1) {
-      *y = 2;
-    }
-  });
+  // test barrier when events are passed arguments
+  *x = 0;
+  for (int i = 0; i < 64; i++) {
+    sycl::event e = q.single_task([=]() {
+      // do some busywork
+      volatile float y = *x;
+      for (int j = 0; j < 100; j++) {
+        y = sycl::cos(y);
+      }
+      // update the value
+      *x *= 2;
+    });
+    q.ext_oneapi_submit_barrier({e});
+    e = q.single_task([=]() { *x += 1; });
+    q.ext_oneapi_submit_barrier({e});
+  }
 
   q.wait_and_throw();
+  error |= (*x != (unsigned long long)-1);
+
+  std::cout << std::bitset<8 * sizeof(unsigned long long)>(*x) << std::endl;
 
-  int error = (*x != 1 || *y != 2) ? 1 : 0;
   std::cout << (error ? "failed\n" : "passed\n");
 
   sycl::free(x, q);
-  sycl::free(y, q);
 
   return error;
 }