[SYCL] Optimize setNDRangeDescriptor functions (intel#18132)

Author: slawekptak

sycl/include/sycl/handler.hpp

@@ -3687,11 +3687,13 @@ class __SYCL_EXPORT handler {
   bool HasAssociatedAccessor(detail::AccessorImplHost *Req,
                              access::target AccessTarget) const;
 
-  template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
+  template <int Dims>
+  static sycl::range<3> padRange(sycl::range<Dims> Range,
+                                 [[maybe_unused]] size_t DefaultValue = 0) {
     if constexpr (Dims == 3) {
       return Range;
     } else {
-      sycl::range<3> Res{0, 0, 0};
+      sycl::range<3> Res{DefaultValue, DefaultValue, DefaultValue};
       for (int I = 0; I < Dims; ++I)
         Res[I] = Range[I];
       return Res;
@@ -3712,7 +3714,8 @@ class __SYCL_EXPORT handler {
   template <int Dims>
   void setNDRangeDescriptor(sycl::range<Dims> N,
                             bool SetNumWorkGroups = false) {
-    return setNDRangeDescriptorPadded(padRange(N), SetNumWorkGroups, Dims);
+    return setNDRangeDescriptorPadded(padRange(N, SetNumWorkGroups ? 0 : 1),
+                                      SetNumWorkGroups, Dims);
   }
   template <int Dims>
   void setNDRangeDescriptor(sycl::range<Dims> NumWorkItems,
@@ -3722,9 +3725,10 @@ class __SYCL_EXPORT handler {
   }
   template <int Dims>
   void setNDRangeDescriptor(sycl::nd_range<Dims> ExecutionRange) {
+    sycl::range<Dims> LocalRange = ExecutionRange.get_local_range();
     return setNDRangeDescriptorPadded(
-        padRange(ExecutionRange.get_global_range()),
-        padRange(ExecutionRange.get_local_range()),
+        padRange(ExecutionRange.get_global_range(), 1),
+        padRange(LocalRange, LocalRange[0] ? 1 : 0),
         padId(ExecutionRange.get_offset()), Dims);
   }
 

sycl/source/detail/cg.hpp

@@ -63,22 +63,13 @@ class ArgDesc {
 // The structure represents NDRange - global, local sizes, global offset and
 // number of dimensions.
 class NDRDescT {
-  // The method initializes all sizes for dimensions greater than the passed one
-  // to the default values, so they will not affect execution.
-  void setNDRangeLeftover() {
-    for (int I = Dims; I < 3; ++I) {
-      GlobalSize[I] = 1;
-      LocalSize[I] = LocalSize[0] ? 1 : 0;
-      GlobalOffset[I] = 0;
-      NumWorkGroups[I] = 0;
-    }
-  }
-
-  template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
+  template <int Dims>
+  static sycl::range<3> padRange(sycl::range<Dims> Range,
+                                 [[maybe_unused]] size_t DefaultValue = 0) {
     if constexpr (Dims == 3) {
       return Range;
     } else {
-      sycl::range<3> Res{0, 0, 0};
+      sycl::range<3> Res{DefaultValue, DefaultValue, DefaultValue};
       for (int I = 0; I < Dims; ++I)
         Res[I] = Range[I];
       return Res;
@@ -102,37 +93,36 @@ class NDRDescT {
   NDRDescT(NDRDescT &&Desc) = default;
 
   NDRDescT(sycl::range<3> N, bool SetNumWorkGroups, int DimsArg)
-      : GlobalSize{SetNumWorkGroups ? sycl::range<3>{0, 0, 0} : N},
-        NumWorkGroups{SetNumWorkGroups ? N : sycl::range<3>{0, 0, 0}},
-        Dims{size_t(DimsArg)} {
-    setNDRangeLeftover();
+      : Dims{size_t(DimsArg)} {
+    if (SetNumWorkGroups) {
+      NumWorkGroups = N;
+    } else {
+      GlobalSize = N;
+    }
   }
 
   NDRDescT(sycl::range<3> NumWorkItems, sycl::range<3> LocalSize,
            sycl::id<3> Offset, int DimsArg)
       : GlobalSize{NumWorkItems}, LocalSize{LocalSize}, GlobalOffset{Offset},
-        Dims{size_t(DimsArg)} {
-    setNDRangeLeftover();
-  }
+        Dims{size_t(DimsArg)} {}
 
   NDRDescT(sycl::range<3> NumWorkItems, sycl::id<3> Offset, int DimsArg)
       : GlobalSize{NumWorkItems}, GlobalOffset{Offset}, Dims{size_t(DimsArg)} {}
 
   template <int Dims_>
   NDRDescT(sycl::nd_range<Dims_> ExecutionRange, int DimsArg)
-      : NDRDescT(padRange(ExecutionRange.get_global_range()),
-                 padRange(ExecutionRange.get_local_range()),
-                 padId(ExecutionRange.get_offset()), size_t(DimsArg)) {
-    setNDRangeLeftover();
-  }
+      : NDRDescT(padRange(ExecutionRange.get_global_range(), 1),
+                 padRange(ExecutionRange.get_local_range(),
+                          ExecutionRange.get_local_range()[0] ? 1 : 0),
+                 padId(ExecutionRange.get_offset()), size_t(DimsArg)) {}
 
   template <int Dims_>
   NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
       : NDRDescT(ExecutionRange, Dims_) {}
 
   template <int Dims_>
   NDRDescT(sycl::range<Dims_> Range)
-      : NDRDescT(padRange(Range), /*SetNumWorkGroups=*/false, Dims_) {}
+      : NDRDescT(padRange(Range, 1), /*SetNumWorkGroups=*/false, Dims_) {}
 
   void setClusterDimensions(sycl::range<3> N, int Dims) {
     if (this->Dims != size_t(Dims)) {