[SYCL] Add structure to pass references to ranges

sycl/include/sycl/detail/ranges_ref_view.hpp

@@ -0,0 +1,57 @@
+//==---- ranges_ref_view.hpp --- SYCL iteration with reference to ranges ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/nd_range.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+class NDRDescT;
+
+// The structure to keep dimension and references to ranges unified for
+// all dimensions.
+class ranges_ref_view {
+
+public:
+  ranges_ref_view() = default;
+  ranges_ref_view(const ranges_ref_view &Desc) = default;
+  ranges_ref_view(ranges_ref_view &&Desc) = default;
+  ranges_ref_view &operator=(const ranges_ref_view &Desc) = default;
+  ranges_ref_view &operator=(ranges_ref_view &&Desc) = default;
+
+  template <int Dims_>
+  ranges_ref_view(sycl::range<Dims_> &GlobalSizes,
+                  sycl::range<Dims_> &LocalSizes)
+      : GlobalSize(&(GlobalSizes[0])), LocalSize(&(LocalSizes[0])),
+        Dims{size_t(Dims_)} {}
+
+  // to support usage in sycl::ext::oneapi::experimental::submit_with_event()
+  template <int Dims_>
+  ranges_ref_view(sycl::nd_range<Dims_> &ExecutionRange)
+      : GlobalSize(&ExecutionRange.globalSize[0]),
+        LocalSize(&ExecutionRange.localSize[0]),
+        GlobalOffset(&ExecutionRange.offset[0]), Dims{size_t(Dims_)} {}
+
+  template <int Dims_>
+  ranges_ref_view(sycl::range<Dims_> &Range)
+      : GlobalSize(&(Range[0])), Dims{size_t(Dims_)} {}
+
+  sycl::detail::NDRDescT toNDRDescT() const;
+
+  const size_t *GlobalSize = nullptr;
+  const size_t *LocalSize = nullptr;
+  const size_t *GlobalOffset = nullptr;
+  size_t Dims = 0;
+};
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/nd_range.hpp

@@ -15,6 +15,10 @@
 namespace sycl {
 inline namespace _V1 {
 
+namespace detail {
+class ranges_ref_view;
+}
+
 /// Defines the iteration domain of both the work-groups and the overall
 /// dispatch.
 ///
@@ -65,6 +69,8 @@ template <int Dimensions = 1> class nd_range {
   bool operator!=(const nd_range<Dimensions> &rhs) const {
     return !(*this == rhs);
   }
+
+  friend class sycl::_V1::detail::ranges_ref_view;
 };
 
 } // namespace _V1

sycl/source/detail/queue_impl.cpp

@@ -11,6 +11,7 @@
 #include <detail/queue_impl.hpp>
 #include <sycl/context.hpp>
 #include <sycl/detail/common.hpp>
+#include <sycl/detail/ranges_ref_view.hpp>
 #include <sycl/detail/ur.hpp>
 #include <sycl/device.hpp>
 
@@ -126,6 +127,24 @@ prepareSYCLEventAssociatedWithQueue(detail::queue_impl &QueueImpl) {
   return detail::createSyclObjFromImpl<event>(EventImpl);
 }
 
+sycl::detail::NDRDescT ranges_ref_view::toNDRDescT() const {
+  NDRDescT NDRDesc;
+
+  NDRDesc.Dims = Dims;
+  for (size_t i = 0; i < Dims; ++i) {
+    NDRDesc.GlobalSize[i] = GlobalSize[i];
+  }
+  if (LocalSize)
+    for (size_t i = 0; i < Dims; ++i) {
+      NDRDesc.LocalSize[i] = LocalSize[i];
+    }
+  if (GlobalOffset)
+    for (size_t i = 0; i < Dims; ++i) {
+      NDRDesc.GlobalOffset[i] = GlobalOffset[i];
+    }
+  return NDRDesc;
+}
+
 const std::vector<event> &
 queue_impl::getExtendDependencyList(const std::vector<event> &DepEvents,
                                     std::vector<event> &MutableVec,

sycl/unittests/scheduler/CMakeLists.txt

@@ -22,4 +22,5 @@ add_sycl_unittest(SchedulerTests OBJECT
     AccessorDefaultCtor.cpp
     HostTaskAndBarrier.cpp
     BarrierDependencies.cpp
+    RangesRefViewUsage.cpp
 )

sycl/unittests/scheduler/RangesRefViewUsage.cpp

@@ -0,0 +1,88 @@
+//==---- RangesRefViewUsage.cpp --- Check ranges_ref_view ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#include <detail/cg.hpp>
+#include <sycl/detail/ranges_ref_view.hpp>
+
+#include <gtest/gtest.h>
+
+template <int dims>
+void TestNDRangesRefView(sycl::range<dims> global, sycl::range<dims> local,
+                         sycl::id<dims> offset) {
+  {
+    sycl::nd_range<dims> nd_range{global, local, offset};
+    sycl::detail::ranges_ref_view r{nd_range};
+    ASSERT_EQ(r.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(r.GlobalSize[d], global[d]);
+      ASSERT_EQ(r.LocalSize[d], local[d]);
+      ASSERT_EQ(r.GlobalOffset[d], offset[d]);
+    }
+
+    sycl::detail::NDRDescT NDRDesc = r.toNDRDescT();
+    ASSERT_EQ(NDRDesc.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(NDRDesc.GlobalSize[d], global[d]);
+      ASSERT_EQ(NDRDesc.LocalSize[d], local[d]);
+      ASSERT_EQ(NDRDesc.GlobalOffset[d], offset[d]);
+    }
+  }
+  {
+    sycl::detail::ranges_ref_view r{global, local};
+    ASSERT_EQ(r.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(r.GlobalSize[d], global[d]);
+      ASSERT_EQ(r.LocalSize[d], local[d]);
+    }
+    ASSERT_EQ(r.GlobalOffset, nullptr);
+
+    sycl::detail::NDRDescT NDRDesc = r.toNDRDescT();
+    ASSERT_EQ(NDRDesc.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(NDRDesc.GlobalSize[d], global[d]);
+      ASSERT_EQ(NDRDesc.LocalSize[d], local[d]);
+    }
+    for (int d = dims; d < 3; d++) {
+      ASSERT_EQ(NDRDesc.GlobalSize[d], 0UL);
+      ASSERT_EQ(NDRDesc.LocalSize[d], 0UL);
+    }
+    for (int d = 0; d < 3; d++) {
+      ASSERT_EQ(NDRDesc.GlobalOffset[d], 0UL);
+    }
+  }
+  {
+    sycl::detail::ranges_ref_view r{global};
+    ASSERT_EQ(r.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(r.GlobalSize[d], global[d]);
+    }
+    ASSERT_EQ(r.LocalSize, nullptr);
+    ASSERT_EQ(r.GlobalOffset, nullptr);
+
+    sycl::detail::NDRDescT NDRDesc = r.toNDRDescT();
+    ASSERT_EQ(NDRDesc.Dims, size_t{dims});
+    for (int d = 0; d < dims; d++) {
+      ASSERT_EQ(NDRDesc.GlobalSize[d], global[d]);
+    }
+    for (int d = dims; d < 3; d++) {
+      ASSERT_EQ(NDRDesc.GlobalSize[d], 0UL);
+    }
+    for (int d = 0; d < 3; d++) {
+      ASSERT_EQ(NDRDesc.LocalSize[d], 0UL);
+      ASSERT_EQ(NDRDesc.GlobalOffset[d], 0UL);
+    }
+  }
+}
+
+TEST(RangesRefUsage, RangesRefUsage) {
+  TestNDRangesRefView(sycl::range<1>{1024}, sycl::range<1>{64},
+                      sycl::id<1>{10});
+  TestNDRangesRefView(sycl::range<2>{1024, 512}, sycl::range<2>{64, 32},
+                      sycl::id<2>{10, 5});
+  TestNDRangesRefView(sycl::range<3>{1024, 512, 256},
+                      sycl::range<3>{64, 32, 16}, sycl::id<3>{10, 5, 2});
+}