[SYCL] Adds regression test for reduction resource leak (#624)

Select variants of reductions currently leak additional resources, such
as auxiliary buffers. This commit adds a regression test to ensure this
leak does not resurface.

Signed-off-by: Steffen Larsen <[email protected]>

Author: steffenlarsen

SYCL/Regression/reduction_resource_leak_dw.cpp

@@ -0,0 +1,32 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: env SYCL_DEVICE_FILTER=level_zero ZE_DEBUG=4 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER
+//
+// CHECK-NOT: LEAK
+
+// Tests that additional resources required by discard_write reductions do not
+// leak.
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int main() {
+  queue Q;
+
+  nd_range<1> NDRange(range<1>{49 * 5}, range<1>{49});
+  std::plus<> BOp;
+
+  buffer<int, 1> OutBuf(1);
+  buffer<int, 1> InBuf(49 * 5);
+  Q.submit([&](handler &CGH) {
+    auto In = InBuf.get_access<access::mode::read>(CGH);
+    auto Out = OutBuf.get_access<access::mode::discard_write>(CGH);
+    auto Redu = ext::oneapi::reduction(Out, 0, BOp);
+    CGH.parallel_for<class DiscardSum>(
+        NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
+          Sum.combine(In[NDIt.get_global_linear_id()]);
+        });
+  });
+  return 0;
+}

SYCL/Regression/reduction_resource_leak_usm.cpp

@@ -0,0 +1,31 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: env SYCL_DEVICE_FILTER=level_zero ZE_DEBUG=4 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER
+//
+// CHECK-NOT: LEAK
+
+// Tests that additional resources required by USM reductions do not leak.
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int main() {
+  queue Q;
+
+  nd_range<1> NDRange(range<1>{49 * 5}, range<1>{49});
+  std::plus<> BOp;
+
+  int *Out = malloc_shared<int>(1, Q);
+  int *In = malloc_shared<int>(49 * 5, Q);
+  Q.submit([&](handler &CGH) {
+     auto Redu = ext::oneapi::reduction(Out, 0, BOp);
+     CGH.parallel_for<class USMSum>(
+         NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
+           Sum.combine(In[NDIt.get_global_linear_id()]);
+         });
+   }).wait();
+  sycl::free(In, Q);
+  sycl::free(Out, Q);
+  return 0;
+}