diff --git a/sycl/include/sycl/detail/nd_range_view.hpp b/sycl/include/sycl/detail/nd_range_view.hpp
new file mode 100644
index 0000000000000..35b5a606566a9
--- /dev/null
+++ b/sycl/include/sycl/detail/nd_range_view.hpp
@@ -0,0 +1,56 @@
+//==---- nd_range_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 nd_range_view {
+
+public:
+  nd_range_view() = default;
+  nd_range_view(const nd_range_view &Desc) = default;
+  nd_range_view(nd_range_view &&Desc) = default;
+  nd_range_view &operator=(const nd_range_view &Desc) = default;
+  nd_range_view &operator=(nd_range_view &&Desc) = default;
+
+  template <int Dims_>
+  nd_range_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_>
+  nd_range_view(sycl::nd_range<Dims_> &ExecutionRange)
+      : GlobalSize(&ExecutionRange.globalSize[0]),
+        LocalSize(&ExecutionRange.localSize[0]),
+        Offset(&ExecutionRange.offset[0]), Dims{size_t(Dims_)} {}
+
+  template <int Dims_>
+  nd_range_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 *Offset = nullptr;
+  size_t Dims = 0;
+};
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/nd_range.hpp b/sycl/include/sycl/nd_range.hpp
index e4ff4881be17a..30816b8a4b354 100644
--- a/sycl/include/sycl/nd_range.hpp
+++ b/sycl/include/sycl/nd_range.hpp
@@ -15,6 +15,10 @@
 namespace sycl {
 inline namespace _V1 {
 
+namespace detail {
+class nd_range_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::nd_range_view;
 };
 
 } // namespace _V1
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index a107ee491dfe3..a5ce0d6f25826 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/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/nd_range_view.hpp>
 #include <sycl/detail/ur.hpp>
 #include <sycl/device.hpp>
 
@@ -125,6 +126,24 @@ prepareSYCLEventAssociatedWithQueue(detail::queue_impl &QueueImpl) {
   return detail::createSyclObjFromImpl<event>(EventImpl);
 }
 
+sycl::detail::NDRDescT nd_range_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 (Offset)
+    for (size_t i = 0; i < Dims; ++i) {
+      NDRDesc.GlobalOffset[i] = Offset[i];
+    }
+  return NDRDesc;
+}
+
 const std::vector<event> &
 queue_impl::getExtendDependencyList(const std::vector<event> &DepEvents,
                                     std::vector<event> &MutableVec,
diff --git a/sycl/unittests/scheduler/CMakeLists.txt b/sycl/unittests/scheduler/CMakeLists.txt
index afc0e185eb7c0..9041793ecdaf2 100644
--- a/sycl/unittests/scheduler/CMakeLists.txt
+++ b/sycl/unittests/scheduler/CMakeLists.txt
@@ -22,4 +22,5 @@ add_sycl_unittest(SchedulerTests OBJECT
     AccessorDefaultCtor.cpp
     HostTaskAndBarrier.cpp
     BarrierDependencies.cpp
+    NdRangeViewUsage.cpp
 )
diff --git a/sycl/unittests/scheduler/NdRangeViewUsage.cpp b/sycl/unittests/scheduler/NdRangeViewUsage.cpp
new file mode 100644
index 0000000000000..65bc466ab32eb
--- /dev/null
+++ b/sycl/unittests/scheduler/NdRangeViewUsage.cpp
@@ -0,0 +1,87 @@
+//==---- NdRangeViewUsage.cpp --- Check nd_range_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/nd_range_view.hpp>
+
+#include <gtest/gtest.h>
+
+template <int dims>
+void TestNdRangeView(sycl::range<dims> global, sycl::range<dims> local,
+                     sycl::id<dims> offset) {
+  {
+    sycl::nd_range<dims> nd_range{global, local, offset};
+    sycl::detail::nd_range_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.Offset[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::nd_range_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.Offset, 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::nd_range_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.Offset, 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) {
+  TestNdRangeView(sycl::range<1>{1024}, sycl::range<1>{64}, sycl::id<1>{10});
+  TestNdRangeView(sycl::range<2>{1024, 512}, sycl::range<2>{64, 32},
+                  sycl::id<2>{10, 5});
+  TestNdRangeView(sycl::range<3>{1024, 512, 256}, sycl::range<3>{64, 32, 16},
+                  sycl::id<3>{10, 5, 2});
+}