[SYCL] Add tests for user-defined reductions extension

Spec: intel/llvm#7202
Implementation: intel/llvm#7436

Author: dm-vodopyanov

UserDefinedReductions/user_defined_reductions.cpp

@@ -0,0 +1,68 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <numeric>
+
+#include <sycl/ext/oneapi/experimental/user_defined_reductions.hpp>
+#include <sycl/sycl.hpp>
+
+template <typename T> struct UserDefinedSum {
+  T operator()(T a, T b) { return a + b; }
+};
+
+using namespace sycl;
+
+int main() {
+  const int N = 1024;
+  queue q;
+
+  buffer<int, 1> inputBuf(N);
+  buffer<int> outputBuf{2};
+  {
+    // Initialize buffer on the host with 0, 1, 2, 3, ..., 1023
+    host_accessor a{inputBuf};
+    std::iota(a.begin(), a.end(), 0);
+  }
+  {
+    q.submit([&](sycl::handler &h) {
+      auto inputValues = sycl::accessor(inputBuf, h);
+      accessor outputValues{outputBuf, h, write_only, no_init};
+
+      constexpr size_t group_size = 16;
+
+      // Create enough local memory for the algorithm
+      constexpr size_t temp_memory_size = group_size * sizeof(int);
+      auto scratch = sycl::local_accessor<std::byte, 1>(temp_memory_size, h);
+
+      h.parallel_for(
+          sycl::nd_range<1>(range<1>(16), range<1>(16)), [=](nd_item<1> it) {
+            // Create a handle that associates the group with an allocation it
+            // can use
+            auto handle =
+                sycl::ext::oneapi::experimental::group_with_scratchpad(
+                    it.get_group(), sycl::span(&scratch[0], temp_memory_size));
+
+            int *first = inputValues.get_pointer();
+            int *last = first + 1024;
+            // Pass the handle as the first argument to the group algorithm
+            int sum_joint_reduce =
+                sycl::ext::oneapi::experimental::joint_reduce(
+                    handle, first, last, UserDefinedSum<int>{});
+            outputValues[0] = sum_joint_reduce;
+
+            int sum_reduce_over_group =
+                sycl::ext::oneapi::experimental::reduce_over_group(
+                    handle, inputValues[it.get_global_id(0)], 0,
+                    UserDefinedSum<int>{});
+            outputValues[1] = sum_reduce_over_group;
+          });
+    });
+    q.wait();
+  }
+  assert(outputBuf.get_host_access()[0] == 523776);
+  assert(outputBuf.get_host_access()[1] == 120);
+
+  // TODO: add tests for reduce_over_group overloads ith native binary_op
+
+  return 0;
+}