[SYCL] Revert "Optimize setNDRangeDescriptor functions (#18132)" (#18326)

The performance gain is not significant enough to justify the change.

Signed-off-by: Ptak, Slawomir <slawomir.ptak@intel.com>

Author: slawekptak

sycl/include/sycl/handler.hpp

@@ -3684,13 +3684,11 @@ 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,
-                                 [[maybe_unused]] size_t DefaultValue = 0) {
+  template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
     if constexpr (Dims == 3) {
       return Range;
     } else {
-      sycl::range<3> Res{DefaultValue, DefaultValue, DefaultValue};
+      sycl::range<3> Res{0, 0, 0};
       for (int I = 0; I < Dims; ++I)
         Res[I] = Range[I];
       return Res;
@@ -3711,8 +3709,7 @@ class __SYCL_EXPORT handler {
   template <int Dims>
   void setNDRangeDescriptor(sycl::range<Dims> N,
                             bool SetNumWorkGroups = false) {
-    return setNDRangeDescriptorPadded(padRange(N, SetNumWorkGroups ? 0 : 1),
-                                      SetNumWorkGroups, Dims);
+    return setNDRangeDescriptorPadded(padRange(N), SetNumWorkGroups, Dims);
   }
   template <int Dims>
   void setNDRangeDescriptor(sycl::range<Dims> NumWorkItems,
@@ -3722,10 +3719,9 @@ 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(), 1),
-        padRange(LocalRange, LocalRange[0] ? 1 : 0),
+        padRange(ExecutionRange.get_global_range()),
+        padRange(ExecutionRange.get_local_range()),
         padId(ExecutionRange.get_offset()), Dims);
   }
 

sycl/source/detail/cg.hpp

@@ -63,13 +63,22 @@ class ArgDesc {
 // The structure represents NDRange - global, local sizes, global offset and
 // number of dimensions.
 class NDRDescT {
-  template <int Dims>
-  static sycl::range<3> padRange(sycl::range<Dims> Range,
-                                 [[maybe_unused]] size_t DefaultValue = 0) {
+  // 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) {
     if constexpr (Dims == 3) {
       return Range;
     } else {
-      sycl::range<3> Res{DefaultValue, DefaultValue, DefaultValue};
+      sycl::range<3> Res{0, 0, 0};
       for (int I = 0; I < Dims; ++I)
         Res[I] = Range[I];
       return Res;
@@ -93,36 +102,37 @@ class NDRDescT {
   NDRDescT(NDRDescT &&Desc) = default;
 
   NDRDescT(sycl::range<3> N, bool SetNumWorkGroups, int DimsArg)
-      : Dims{size_t(DimsArg)} {
-    if (SetNumWorkGroups) {
-      NumWorkGroups = N;
-    } else {
-      GlobalSize = N;
-    }
+      : GlobalSize{SetNumWorkGroups ? sycl::range<3>{0, 0, 0} : N},
+        NumWorkGroups{SetNumWorkGroups ? N : sycl::range<3>{0, 0, 0}},
+        Dims{size_t(DimsArg)} {
+    setNDRangeLeftover();
   }
 
   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)} {}
+        Dims{size_t(DimsArg)} {
+    setNDRangeLeftover();
+  }
 
   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(), 1),
-                 padRange(ExecutionRange.get_local_range(),
-                          ExecutionRange.get_local_range()[0] ? 1 : 0),
-                 padId(ExecutionRange.get_offset()), size_t(DimsArg)) {}
+      : NDRDescT(padRange(ExecutionRange.get_global_range()),
+                 padRange(ExecutionRange.get_local_range()),
+                 padId(ExecutionRange.get_offset()), size_t(DimsArg)) {
+    setNDRangeLeftover();
+  }
 
   template <int Dims_>
   NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
       : NDRDescT(ExecutionRange, Dims_) {}
 
   template <int Dims_>
   NDRDescT(sycl::range<Dims_> Range)
-      : NDRDescT(padRange(Range, 1), /*SetNumWorkGroups=*/false, Dims_) {}
+      : NDRDescT(padRange(Range), /*SetNumWorkGroups=*/false, Dims_) {}
 
   void setClusterDimensions(sycl::range<3> N, int Dims) {
     if (this->Dims != size_t(Dims)) {