Temp no handler range

sycl/include/sycl/detail/nd_range_view.hpp

@@ -0,0 +1,57 @@
+//==---- 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_> &N, bool SetNumWorkGroups = false)
+      : MGlobalSize(&(N[0])), MSetNumWorkGroups(SetNumWorkGroups),
+        MDims{size_t(Dims_)} {}
+
+  template <int Dims_>
+  nd_range_view(sycl::range<Dims_> &GlobalSize, sycl::id<Dims_> &Offset)
+      : MGlobalSize(&(GlobalSize[0])), MOffset(&(Offset[0])),
+        MDims{size_t(Dims_)} {}
+
+  template <int Dims_>
+  nd_range_view(sycl::nd_range<Dims_> &ExecutionRange)
+      : MGlobalSize(&(ExecutionRange.globalSize[0])),
+        MLocalSize(&(ExecutionRange.localSize[0])),
+        MOffset(&(ExecutionRange.offset[0])), MDims{size_t(Dims_)} {}
+
+  sycl::detail::NDRDescT toNDRDescT() const;
+
+  const size_t *MGlobalSize = nullptr;
+  const size_t *MLocalSize = nullptr;
+  const size_t *MOffset = nullptr;
+  bool MSetNumWorkGroups = false;
+  size_t MDims = 0;
+};
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/handler.hpp

@@ -362,6 +362,24 @@ class RoundedRangeKernelWithKH {
   }
 };
 
+template <typename WrapperT, typename TransformedArgType, int Dims,
+          typename KernelType,
+          std::enable_if_t<detail::KernelLambdaHasKernelHandlerArgT<
+              KernelType, TransformedArgType>::value> * = nullptr>
+auto getRangeRoundedKernelLambda(KernelType KernelFunc, range<Dims> UserRange) {
+  return detail::RoundedRangeKernelWithKH<TransformedArgType, Dims, KernelType>{
+      UserRange, KernelFunc};
+}
+
+template <typename WrapperT, typename TransformedArgType, int Dims,
+          typename KernelType,
+          std::enable_if_t<!detail::KernelLambdaHasKernelHandlerArgT<
+              KernelType, TransformedArgType>::value> * = nullptr>
+auto getRangeRoundedKernelLambda(KernelType KernelFunc, range<Dims> UserRange) {
+  return detail::RoundedRangeKernel<TransformedArgType, Dims, KernelType>{
+      UserRange, KernelFunc};
+}
+
 using std::enable_if_t;
 using sycl::detail::queue_impl;
 
@@ -1218,7 +1236,7 @@ class __SYCL_EXPORT handler {
     if (HasRoundedRange) {
       using NameWT = typename detail::get_kernel_wrapper_name_t<NameT>::name;
       auto Wrapper =
-          getRangeRoundedKernelLambda<NameWT, TransformedArgType, Dims>(
+          detail::getRangeRoundedKernelLambda<NameWT, TransformedArgType, Dims>(
               KernelFunc, UserRange);
 
       using KName = std::conditional_t<std::is_same<KernelType, NameT>::value,
@@ -3265,26 +3283,6 @@ class __SYCL_EXPORT handler {
   void GetRangeRoundingSettings(size_t &MinFactor, size_t &GoodFactor,
                                 size_t &MinRange);
 
-  template <typename WrapperT, typename TransformedArgType, int Dims,
-            typename KernelType,
-            std::enable_if_t<detail::KernelLambdaHasKernelHandlerArgT<
-                KernelType, TransformedArgType>::value> * = nullptr>
-  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
-                                   range<Dims> UserRange) {
-    return detail::RoundedRangeKernelWithKH<TransformedArgType, Dims,
-                                            KernelType>{UserRange, KernelFunc};
-  }
-
-  template <typename WrapperT, typename TransformedArgType, int Dims,
-            typename KernelType,
-            std::enable_if_t<!detail::KernelLambdaHasKernelHandlerArgT<
-                KernelType, TransformedArgType>::value> * = nullptr>
-  auto getRangeRoundedKernelLambda(KernelType KernelFunc,
-                                   range<Dims> UserRange) {
-    return detail::RoundedRangeKernel<TransformedArgType, Dims, KernelType>{
-        UserRange, KernelFunc};
-  }
-
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   const std::shared_ptr<detail::context_impl> &getContextImplPtr() const;
 #endif

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