[SYCL] Extend range rounding tests to queue shortcut

Introduce a kernel submission wrapper utility to E2E tests,
and use it in the range rounding tests to increase the test
coverage to the queue kernel submission shortcut function.

The main motivation is to test the handler-less kernel submission
path, once it is introduced (for range based submissions).

Author: slawekptak

sycl/test-e2e/Basic/parallel_for_disable_range_roundup.cpp

@@ -6,52 +6,62 @@
 // RUN: %{build} -sycl-std=2020 -o %t2.out
 // RUN: env SYCL_PARALLEL_FOR_RANGE_ROUNDING_TRACE=1 %{run} %t2.out | FileCheck %s --check-prefix CHECK-ENABLED
 
-#include <sycl/atomic.hpp>
+#include <sycl/atomic_ref.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include "../helpers.hpp"
 
 #include <iostream>
 using namespace sycl;
 
-range<1> Range1 = {0};
-
 void check(const char *msg, size_t v, size_t ref) {
   std::cout << msg << v << std::endl;
   assert(v == ref);
 }
 
-int try_rounding_off(size_t size) {
-  range<1> Size{size};
-  int Counter = 0;
-  {
-    buffer<range<1>, 1> BufRange(&Range1, 1);
-    buffer<int, 1> BufCounter(&Counter, 1);
-    queue myQueue;
-
-    std::cout << "Run parallel_for" << std::endl;
-    myQueue.submit([&](handler &cgh) {
-      auto AccRange = BufRange.get_access<access::mode::read_write>(cgh);
-      auto AccCounter = BufCounter.get_access<access::mode::atomic>(cgh);
-      cgh.parallel_for<class PF_init_item1>(Size, [=](item<1> ITEM) {
-        AccCounter[0].fetch_add(1);
-        AccRange[0] = ITEM.get_range(0);
-      });
-    });
-    myQueue.wait();
-  }
-  check("Size seen by user = ", Range1.get(0), size);
-  check("Counter = ", Counter, size);
-  return 0;
+void try_rounding_off(size_t size, bool useShortcutFunction) {
+  range<1> Range{size};
+  queue Queue;
+  range<1> *RangePtr = malloc_shared<range<1>>(1, Queue);
+  int *CounterPtr = malloc_shared<int>(1, Queue);
+  (*CounterPtr) = 0;
+
+  std::cout << "Run parallel_for" << std::endl;
+  auto KernelFunc = [=](item<1> id) {
+    auto atm = atomic_ref<int, sycl::memory_order::relaxed,
+                          sycl::memory_scope::device>(*CounterPtr);
+    atm.fetch_add(1);
+    (*RangePtr) = id.get_range(0);
+  };
+  command_submit_wrappers::parallel_for_wrapper<class TestKernel>(
+      useShortcutFunction, Queue, Range, KernelFunc);
+
+  Queue.wait();
+
+  auto Context = Queue.get_context();
+
+  check("Size seen by user = ", RangePtr->get(0), size);
+  check("Counter = ", *CounterPtr, size);
+
+  free(RangePtr, Context);
+  free(CounterPtr, Context);
 }
 
 int main() {
   int x;
 
   x = 1500;
-  try_rounding_off(x);
+  try_rounding_off(x, true);
+  try_rounding_off(x, false);
 
   return 0;
 }
 
+// CHECK-DISABLED:  Run parallel_for
+// CHECK-DISABLED-NOT: parallel_for range adjusted at dim 0 from 1500
+// CHECK-DISABLED:  Size seen by user = 1500
+// CHECK-DISABLED-NEXT:  Counter = 1500
 // CHECK-DISABLED:  Run parallel_for
 // CHECK-DISABLED-NOT: parallel_for range adjusted at dim 0 from 1500
 // CHECK-DISABLED:  Size seen by user = 1500
@@ -61,3 +71,7 @@ int main() {
 // CHECK-ENABLED-NEXT: parallel_for range adjusted at dim 0 from 1500
 // CHECK-ENABLED-NEXT:  Size seen by user = 1500
 // CHECK-ENABLED-NEXT:  Counter = 1500
+// CHECK-ENABLED:  Run parallel_for
+// CHECK-ENABLED-NEXT: parallel_for range adjusted at dim 0 from 1500
+// CHECK-ENABLED-NEXT:  Size seen by user = 1500
+// CHECK-ENABLED-NEXT:  Counter = 1500