[SYCL] Move the NDRDescT class to a public header

sycl/include/sycl/queue.hpp

@@ -149,6 +149,101 @@ class __SYCL_EXPORT SubmissionInfo {
       ext::oneapi::experimental::event_mode_enum::none;
 };
 
+// The structure represents NDRange - global, local sizes, global offset and
+// number of dimensions.
+
+// TODO: A lot of tests rely on particular values to be set for dimensions that
+// are not used. To clarify, for example, if a 2D kernel is invoked, in
+// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
+// depending on which constructor is used for no clear reason.
+// Instead, only sensible defaults should be used and tests should be updated
+// to reflect this.
+class NDRDescT {
+
+public:
+  NDRDescT() = default;
+  NDRDescT(const NDRDescT &Desc) = default;
+  NDRDescT(NDRDescT &&Desc) = default;
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
+    if (SetNumWorkGroups) {
+      for (size_t I = 0; I < Dims_; ++I) {
+        NumWorkGroups[I] = N[I];
+      }
+    } else {
+      for (size_t I = 0; I < Dims_; ++I) {
+        GlobalSize[I] = N[I];
+      }
+
+      for (int I = Dims_; I < 3; ++I) {
+        GlobalSize[I] = 1;
+      }
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
+           sycl::id<Dims_> Offset)
+      : Dims{size_t(Dims_)} {
+    for (size_t I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      LocalSize[I] = LocalSizes[I];
+      GlobalOffset[I] = Offset[I];
+    }
+
+    for (int I = Dims_; I < 3; ++I) {
+      LocalSize[I] = LocalSizes[0] ? 1 : 0;
+    }
+
+    for (int I = Dims_; I < 3; ++I) {
+      GlobalSize[I] = 1;
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
+      : Dims{size_t(Dims_)} {
+    for (size_t I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      GlobalOffset[I] = Offset[I];
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
+      : NDRDescT(ExecutionRange.get_global_range(),
+                 ExecutionRange.get_local_range(),
+                 ExecutionRange.get_offset()) {}
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> Range)
+      : NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
+
+  template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
+    if (this->Dims != size_t(Dims_)) {
+      throw std::runtime_error(
+          "Dimensionality of cluster, global and local ranges must be same");
+    }
+
+    for (int I = 0; I < Dims_; ++I)
+      ClusterDimensions[I] = N[I];
+  }
+
+  NDRDescT &operator=(const NDRDescT &Desc) = default;
+  NDRDescT &operator=(NDRDescT &&Desc) = default;
+
+  std::array<size_t, 3> GlobalSize{0, 0, 0};
+  std::array<size_t, 3> LocalSize{0, 0, 0};
+  std::array<size_t, 3> GlobalOffset{0, 0, 0};
+  /// Number of workgroups, used to record the number of workgroups from the
+  /// simplest form of parallel_for_work_group. If set, all other fields must be
+  /// zero
+  std::array<size_t, 3> NumWorkGroups{0, 0, 0};
+  std::array<size_t, 3> ClusterDimensions{1, 1, 1};
+  size_t Dims = 0;
+};
+
 } // namespace v1
 } // namespace detail
 

sycl/source/detail/cg.hpp

@@ -60,101 +60,6 @@ class ArgDesc {
   int MIndex;
 };
 
-// The structure represents NDRange - global, local sizes, global offset and
-// number of dimensions.
-
-// TODO: A lot of tests rely on particular values to be set for dimensions that
-// are not used. To clarify, for example, if a 2D kernel is invoked, in
-// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
-// depending on which constructor is used for no clear reason.
-// Instead, only sensible defaults should be used and tests should be updated
-// to reflect this.
-class NDRDescT {
-
-public:
-  NDRDescT() = default;
-  NDRDescT(const NDRDescT &Desc) = default;
-  NDRDescT(NDRDescT &&Desc) = default;
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
-    if (SetNumWorkGroups) {
-      for (size_t I = 0; I < Dims_; ++I) {
-        NumWorkGroups[I] = N[I];
-      }
-    } else {
-      for (size_t I = 0; I < Dims_; ++I) {
-        GlobalSize[I] = N[I];
-      }
-
-      for (int I = Dims_; I < 3; ++I) {
-        GlobalSize[I] = 1;
-      }
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
-           sycl::id<Dims_> Offset)
-      : Dims{size_t(Dims_)} {
-    for (size_t I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      LocalSize[I] = LocalSizes[I];
-      GlobalOffset[I] = Offset[I];
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      LocalSize[I] = LocalSizes[0] ? 1 : 0;
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      GlobalSize[I] = 1;
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
-      : Dims{size_t(Dims_)} {
-    for (size_t I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      GlobalOffset[I] = Offset[I];
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
-      : NDRDescT(ExecutionRange.get_global_range(),
-                 ExecutionRange.get_local_range(),
-                 ExecutionRange.get_offset()) {}
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> Range)
-      : NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
-
-  template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
-    if (this->Dims != size_t(Dims_)) {
-      throw std::runtime_error(
-          "Dimensionality of cluster, global and local ranges must be same");
-    }
-
-    for (int I = 0; I < Dims_; ++I)
-      ClusterDimensions[I] = N[I];
-  }
-
-  NDRDescT &operator=(const NDRDescT &Desc) = default;
-  NDRDescT &operator=(NDRDescT &&Desc) = default;
-
-  std::array<size_t, 3> GlobalSize{0, 0, 0};
-  std::array<size_t, 3> LocalSize{0, 0, 0};
-  std::array<size_t, 3> GlobalOffset{0, 0, 0};
-  /// Number of workgroups, used to record the number of workgroups from the
-  /// simplest form of parallel_for_work_group. If set, all other fields must be
-  /// zero
-  std::array<size_t, 3> NumWorkGroups{0, 0, 0};
-  std::array<size_t, 3> ClusterDimensions{1, 1, 1};
-  size_t Dims = 0;
-};
-
 /// Base class for all types of command groups.
 class CG {
 public:
@@ -248,7 +153,7 @@ class CG {
 class CGExecKernel : public CG {
 public:
   /// Stores ND-range description.
-  NDRDescT MNDRDesc;
+  detail::v1::NDRDescT MNDRDesc;
   std::shared_ptr<HostKernelBase> MHostKernel;
   std::shared_ptr<detail::kernel_impl> MSyclKernel;
   std::shared_ptr<detail::kernel_bundle_impl> MKernelBundle;
@@ -265,7 +170,7 @@ class CGExecKernel : public CG {
   bool MKernelUsesClusterLaunch = false;
   size_t MKernelWorkGroupMemorySize = 0;
 
-  CGExecKernel(NDRDescT NDRDesc, std::shared_ptr<HostKernelBase> HKernel,
+  CGExecKernel(detail::v1::NDRDescT NDRDesc, std::shared_ptr<HostKernelBase> HKernel,
                std::shared_ptr<detail::kernel_impl> SyclKernel,
                std::shared_ptr<detail::kernel_bundle_impl> KernelBundle,
                CG::StorageInitHelper CGData, std::vector<ArgDesc> Args,

sycl/source/detail/error_handling/error_handling.cpp

@@ -22,7 +22,7 @@ inline namespace _V1 {
 namespace detail::enqueue_kernel_launch {
 
 void handleOutOfResources(const device_impl &DeviceImpl,
-                          ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) {
+                          ur_kernel_handle_t Kernel, const detail::v1::NDRDescT &NDRDesc) {
   sycl::platform Platform = DeviceImpl.get_platform();
   sycl::backend Backend = Platform.get_backend();
   if (Backend == sycl::backend::ext_oneapi_cuda) {
@@ -75,7 +75,7 @@ void handleOutOfResources(const device_impl &DeviceImpl,
 
 void handleInvalidWorkGroupSize(const device_impl &DeviceImpl,
                                 ur_kernel_handle_t Kernel,
-                                const NDRDescT &NDRDesc) {
+                                const detail::v1::NDRDescT &NDRDesc) {
   sycl::platform Platform = DeviceImpl.get_platform();
 
   // Some of the error handling below is special for particular OpenCL
@@ -349,7 +349,7 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl,
 }
 
 void handleInvalidWorkItemSize(const device_impl &DeviceImpl,
-                               const NDRDescT &NDRDesc) {
+                               const detail::v1::NDRDescT &NDRDesc) {
 
   const AdapterPtr &Adapter = DeviceImpl.getAdapter();
   ur_device_handle_t Device = DeviceImpl.getHandleRef();
@@ -370,7 +370,7 @@ void handleInvalidWorkItemSize(const device_impl &DeviceImpl,
 }
 
 void handleInvalidValue(const device_impl &DeviceImpl,
-                        const NDRDescT &NDRDesc) {
+                        const detail::v1::NDRDescT &NDRDesc) {
   const AdapterPtr &Adapter = DeviceImpl.getAdapter();
   ur_device_handle_t Device = DeviceImpl.getHandleRef();
 
@@ -393,7 +393,7 @@ void handleInvalidValue(const device_impl &DeviceImpl,
 }
 
 void handleErrorOrWarning(ur_result_t Error, const device_impl &DeviceImpl,
-                          ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) {
+                          ur_kernel_handle_t Kernel, const detail::v1::NDRDescT &NDRDesc) {
   assert(Error != UR_RESULT_SUCCESS &&
          "Success is expected to be handled on caller side");
   switch (Error) {

sycl/source/detail/error_handling/error_handling.hpp

@@ -26,7 +26,7 @@ namespace enqueue_kernel_launch {
 /// This function actually never returns and always throws an exception with
 /// error description.
 void handleErrorOrWarning(ur_result_t, const device_impl &, ur_kernel_handle_t,
-                          const NDRDescT &);
+                          const detail::v1::NDRDescT &);
 } // namespace enqueue_kernel_launch
 
 namespace kernel_get_group_info {

sycl/source/detail/graph/graph_impl.cpp

@@ -1573,7 +1573,7 @@ void exec_graph_impl::populateURKernelUpdateStructs(
     std::vector<ur_kernel_arg_mem_obj_properties_t> &MemobjProps,
     std::vector<ur_exp_command_buffer_update_pointer_arg_desc_t> &PtrDescs,
     std::vector<ur_exp_command_buffer_update_value_arg_desc_t> &ValueDescs,
-    sycl::detail::NDRDescT &NDRDesc,
+    sycl::detail::v1::NDRDescT &NDRDesc,
     ur_exp_command_buffer_update_kernel_launch_desc_t &UpdateDesc) const {
   sycl::detail::context_impl &ContextImpl =
       *sycl::detail::getSyclObjImpl(MContext);
@@ -1800,7 +1800,7 @@ void exec_graph_impl::updateURImpl(
       PtrDescsList(NumUpdatableNodes);
   std::vector<std::vector<ur_exp_command_buffer_update_value_arg_desc_t>>
       ValueDescsList(NumUpdatableNodes);
-  std::vector<sycl::detail::NDRDescT> NDRDescList(NumUpdatableNodes);
+  std::vector<sycl::detail::v1::NDRDescT> NDRDescList(NumUpdatableNodes);
   std::vector<ur_exp_command_buffer_update_kernel_launch_desc_t> UpdateDescList(
       NumUpdatableNodes);
   std::vector<FastKernelCacheValPtr> KernelBundleObjList(NumUpdatableNodes);

sycl/source/detail/graph/graph_impl.hpp

@@ -31,7 +31,6 @@ class handler;
 namespace detail {
 class SYCLMemObjT;
 class queue_impl;
-class NDRDescT;
 class ArgDesc;
 class CG;
 } // namespace detail
@@ -869,7 +868,7 @@ class exec_graph_impl {
       std::vector<ur_kernel_arg_mem_obj_properties_t> &MemobjProps,
       std::vector<ur_exp_command_buffer_update_pointer_arg_desc_t> &PtrDescs,
       std::vector<ur_exp_command_buffer_update_value_arg_desc_t> &ValueDescs,
-      sycl::detail::NDRDescT &NDRDesc,
+      sycl::detail::v1::NDRDescT &NDRDesc,
       ur_exp_command_buffer_update_kernel_launch_desc_t &UpdateDesc) const;
 
   /// Execution schedule of nodes in the graph.

sycl/source/detail/graph/node_impl.hpp

@@ -437,7 +437,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
                             "the node was originally created with.");
     }
 
-    NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
+    NDRDesc = sycl::detail::v1::NDRDescT{ExecutionRange};
   }
 
   template <int Dimensions> void updateRange(range<Dimensions> ExecutionRange) {
@@ -458,7 +458,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
                             "the node was originally created with.");
     }
 
-    NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
+    NDRDesc = sycl::detail::v1::NDRDescT{ExecutionRange};
   }
   /// Update this node with the command-group from another node.
   /// @param Other The other node to update, must be of the same node type.

sycl/source/detail/handler_impl.hpp

@@ -155,7 +155,7 @@ class handler_impl {
   std::vector<detail::ArgDesc> MAssociatedAccesors;
 
   /// Struct that encodes global size, local size, ...
-  detail::NDRDescT MNDRDesc;
+  detail::v1::NDRDescT MNDRDesc;
 
   /// Type of the command group, e.g. kernel, fill. Can also encode version.
   /// Use getType and setType methods to access this variable unless

sycl/source/detail/program_manager/program_manager.cpp

@@ -862,7 +862,7 @@ static void setSpecializationConstants(device_image_impl &InputImpl,
 // its ref count incremented.
 ur_program_handle_t ProgramManager::getBuiltURProgram(
     context_impl &ContextImpl, device_impl &DeviceImpl,
-    KernelNameStrRefT KernelName, const NDRDescT &NDRDesc) {
+    KernelNameStrRefT KernelName, const detail::v1::NDRDescT &NDRDesc) {
   device_impl *RootDevImpl;
   ur_bool_t MustBuildOnSubdevice = true;
 
@@ -1124,7 +1124,7 @@ ur_program_handle_t ProgramManager::getBuiltURProgram(
 FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
     context_impl &ContextImpl, device_impl &DeviceImpl,
     KernelNameStrRefT KernelName,
-    KernelNameBasedCacheT *KernelNameBasedCachePtr, const NDRDescT &NDRDesc) {
+    KernelNameBasedCacheT *KernelNameBasedCachePtr, const detail::v1::NDRDescT &NDRDesc) {
   if constexpr (DbgProgMgr > 0) {
     std::cerr << ">>> ProgramManager::getOrCreateKernel(" << &ContextImpl
               << ", " << &DeviceImpl << ", " << KernelName << ")\n";
@@ -3593,7 +3593,7 @@ std::optional<sycl::exception> checkDevSupportJointMatrixMad(
 std::optional<sycl::exception>
 checkDevSupportDeviceRequirements(const device_impl &Dev,
                                   const RTDeviceBinaryImage &Img,
-                                  const NDRDescT &NDRDesc) {
+                                  const detail::v1::NDRDescT &NDRDesc) {
   auto getPropIt = [&Img](const std::string &PropName) {
     auto &PropRange = Img.getDeviceRequirements();
     RTDeviceBinaryImage::PropertyRange::ConstIterator PropIt = std::find_if(

sycl/source/detail/program_manager/program_manager.hpp

@@ -62,7 +62,7 @@ bool doesDevSupportDeviceRequirements(const device_impl &Dev,
 std::optional<sycl::exception>
 checkDevSupportDeviceRequirements(const device_impl &Dev,
                                   const RTDeviceBinaryImage &BinImages,
-                                  const NDRDescT &NDRDesc = {});
+                                  const detail::v1::NDRDescT &NDRDesc = {});
 
 bool doesImageTargetMatchDevice(const RTDeviceBinaryImage &Img,
                                 const device_impl &DevImpl);
@@ -179,7 +179,7 @@ class ProgramManager {
   ur_program_handle_t getBuiltURProgram(context_impl &ContextImpl,
                                         device_impl &DeviceImpl,
                                         KernelNameStrRefT KernelName,
-                                        const NDRDescT &NDRDesc = {});
+                                        const detail::v1::NDRDescT &NDRDesc = {});
 
   /// Builds a program from a given set of images or retrieves that program from
   /// cache.
@@ -201,7 +201,7 @@ class ProgramManager {
   getOrCreateKernel(context_impl &ContextImpl, device_impl &DeviceImpl,
                     KernelNameStrRefT KernelName,
                     KernelNameBasedCacheT *KernelNameBasedCachePtr,
-                    const NDRDescT &NDRDesc = {});
+                    const detail::v1::NDRDescT &NDRDesc = {});
 
   ur_kernel_handle_t getCachedMaterializedKernel(
       KernelNameStrRefT KernelName,