[SYCL] Add structure to pass references to ranges

sycl/include/sycl/ext/oneapi/experimental/rangesref.hpp

@@ -0,0 +1,53 @@
+//==-------- rangesref.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 ext::oneapi::experimental {
+
+// The structure to keep references to ranges and dimension unified for
+// all dimensions.
+class RangesRefT {
+
+public:
+  RangesRefT() = default;
+  RangesRefT(const RangesRefT &Desc) = default;
+  RangesRefT(RangesRefT &&Desc) = default;
+
+  template <int Dims_>
+  RangesRefT(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_>
+  RangesRefT(sycl::nd_range<Dims_> &ExecutionRange)
+      : GlobalSize(&ExecutionRange.globalSize[0]),
+        LocalSize(&ExecutionRange.localSize[0]),
+        GlobalOffset(&ExecutionRange.offset[0]), Dims{size_t(Dims_)} {}
+
+  template <int Dims_>
+  RangesRefT(sycl::range<Dims_> &Range)
+      : GlobalSize(&(Range[0])), Dims{size_t(Dims_)} {}
+
+  RangesRefT &operator=(const RangesRefT &Desc) = default;
+  RangesRefT &operator=(RangesRefT &&Desc) = default;
+
+  const size_t *GlobalSize = nullptr;
+  const size_t *LocalSize = nullptr;
+  const size_t *GlobalOffset = nullptr;
+  size_t Dims = 0;
+};
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/nd_range.hpp

@@ -15,6 +15,10 @@
 namespace sycl {
 inline namespace _V1 {
 
+namespace ext::oneapi::experimental {
+class RangesRefT;
+}
+
 /// 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::ext::oneapi::experimental::RangesRefT;
 };
 
 } // namespace _V1

sycl/unittests/scheduler/CMakeLists.txt

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

sycl/unittests/scheduler/RangesRefUsage.cpp

@@ -0,0 +1,112 @@
+//==---- RangesRefUsage.cpp --- Check RangesRefT --------------------------==//
+//
+// 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 <sycl/ext/oneapi/experimental/rangesref.hpp>
+
+#include <gtest/gtest.h>
+
+TEST(RangesRefUsage, RangesRefUsage) {
+  sycl::ext::oneapi::experimental::RangesRefT r0;
+  ASSERT_EQ(r0.Dims, size_t{0});
+
+  {
+    sycl::range<1> global_range{1024};
+    sycl::range<1> local_range{64};
+    sycl::id<1> offset{10};
+    sycl::nd_range<1> nd_range{global_range, local_range, offset};
+
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{nd_range};
+      ASSERT_EQ(r.Dims, size_t{1});
+      ASSERT_EQ(*r.GlobalSize, global_range[0]);
+      ASSERT_EQ(*r.LocalSize, local_range[0]);
+      ASSERT_EQ(*r.GlobalOffset, offset[0]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range, local_range};
+      ASSERT_EQ(r.Dims, size_t{1});
+      ASSERT_EQ(*r.GlobalSize, global_range[0]);
+      ASSERT_EQ(*r.LocalSize, local_range[0]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range};
+      ASSERT_EQ(r.Dims, size_t{1});
+      ASSERT_EQ(*r.GlobalSize, global_range[0]);
+      ASSERT_EQ(r.LocalSize, nullptr);
+    }
+  }
+  {
+    sycl::range<2> global_range{1024, 512};
+    sycl::range<2> local_range{64, 32};
+    sycl::id<2> offset{10, 20};
+    sycl::nd_range<2> nd_range{global_range, local_range, offset};
+
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{nd_range};
+      ASSERT_EQ(r.Dims, size_t{2});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.LocalSize[0], local_range[0]);
+      ASSERT_EQ(r.LocalSize[1], local_range[1]);
+      ASSERT_EQ(r.GlobalOffset[0], offset[0]);
+      ASSERT_EQ(r.GlobalOffset[1], offset[1]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range, local_range};
+      ASSERT_EQ(r.Dims, size_t{2});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.LocalSize[0], local_range[0]);
+      ASSERT_EQ(r.LocalSize[1], local_range[1]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range};
+      ASSERT_EQ(r.Dims, size_t{2});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.LocalSize, nullptr);
+    }
+  }
+  {
+    sycl::range<3> global_range{1024, 512, 256};
+    sycl::range<3> local_range{64, 32, 16};
+    sycl::id<3> offset{10, 20, 30};
+    sycl::nd_range<3> nd_range{global_range, local_range, offset};
+
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{nd_range};
+      ASSERT_EQ(r.Dims, size_t{3});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.GlobalSize[2], global_range[2]);
+      ASSERT_EQ(r.LocalSize[0], local_range[0]);
+      ASSERT_EQ(r.LocalSize[1], local_range[1]);
+      ASSERT_EQ(r.LocalSize[2], local_range[2]);
+      ASSERT_EQ(r.GlobalOffset[0], offset[0]);
+      ASSERT_EQ(r.GlobalOffset[1], offset[1]);
+      ASSERT_EQ(r.GlobalOffset[2], offset[2]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range, local_range};
+      ASSERT_EQ(r.Dims, size_t{3});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.GlobalSize[2], global_range[2]);
+      ASSERT_EQ(r.LocalSize[0], local_range[0]);
+      ASSERT_EQ(r.LocalSize[1], local_range[1]);
+      ASSERT_EQ(r.LocalSize[2], local_range[2]);
+    }
+    {
+      sycl::ext::oneapi::experimental::RangesRefT r{global_range};
+      ASSERT_EQ(r.Dims, size_t{3});
+      ASSERT_EQ(r.GlobalSize[0], global_range[0]);
+      ASSERT_EQ(r.GlobalSize[1], global_range[1]);
+      ASSERT_EQ(r.GlobalSize[2], global_range[2]);
+      ASSERT_EQ(r.LocalSize, nullptr);
+    }
+  }
+}