[SYCL] global size zero assertion narrowed

sycl/source/detail/scheduler/commands.cpp

@@ -2292,10 +2292,12 @@ std::string_view ExecCGCommand::getTypeString() const {
 // for users who need more control.
 static void adjustNDRangePerKernel(NDRDescT &NDR, ur_kernel_handle_t Kernel,
                                    const device_impl &DeviceImpl) {
-  if (NDR.GlobalSize[0] != 0)
-    return; // GlobalSize is set - no need to adjust
-  // check the prerequisites:
-  assert(NDR.LocalSize[0] == 0);
+  if (NDR.NumWorkGroups[0] == 0)
+    return; // Not parallel_for_work_group -- nothing to fill in.
+  // In pfwg mode NumWorkGroups is the only field the user sets; GlobalSize
+  // and LocalSize must both be zero (see NDRDescT contract in
+  // ndrange_desc.hpp).
+  assert(NDR.GlobalSize[0] == 0 && NDR.LocalSize[0] == 0);
   // TODO might be good to cache this info together with the kernel info to
   // avoid get_kernel_work_group_info on every kernel run
   range<3> WGSize = get_kernel_device_specific_info<

sycl/test-e2e/Basic/parallel_for_zero_range.cpp

@@ -0,0 +1,65 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// The native_cpu UR adapter explicitly rejects zero global work size in
+// urEnqueueKernelLaunch (DIE_NO_IMPLEMENTATION ->
+// UR_RESULT_ERROR_UNSUPPORTED_FEATURE).
+// XFAIL: target-native_cpu
+// XFAIL-TRACKER: CMPLRLLVM-77780
+
+// SYCL 2020 (Work-group data parallel kernels): "When the global size is
+// zero, the kernel function is not executed, the local size is ignored, and
+// any dependencies are satisfied."
+//
+// See intel/llvm#22893
+//
+// A USM-shared sentinel byte is set to 0 before each submit. If a kernel
+// body actually ran, it would flip the byte to 0xFF. After Q.wait(), we
+// assert the byte is still 0.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <cassert>
+
+using namespace sycl;
+
+int main() {
+  queue Q;
+
+  unsigned char *Sentinel = malloc_shared<unsigned char>(1, Q);
+  assert(Sentinel && "USM shared alloc failed");
+
+  // Case 1: parallel_for(range<1>{0}) -- plain empty range.
+  *Sentinel = 0x00;
+  Q.submit([&](handler &cgh) {
+     cgh.parallel_for<class zero_range_pf>(range<1>{0},
+                                           [=](id<1>) { *Sentinel = 0xFF; });
+   }).wait();
+  assert(*Sentinel == 0x00 && "parallel_for(range{0}) unexpectedly launched");
+
+  // Case 2: parallel_for(nd_range<1>{{0}, {32}}) -- the PyTorch shape:
+  // zero global size, non-zero local size. Pre-fix, this tripped the
+  // assertion in adjustNDRangePerKernel.
+  *Sentinel = 0x00;
+  Q.submit([&](handler &cgh) {
+     cgh.parallel_for<class zero_range_ndr>(
+         nd_range<1>{range<1>{0}, range<1>{32}},
+         [=](nd_item<1>) { *Sentinel = 0xFF; });
+   }).wait();
+  assert(*Sentinel == 0x00 &&
+         "parallel_for(nd_range{0, 32}) unexpectedly launched");
+
+  // Case 3: parallel_for(nd_range<1>{{0}, {0}})
+  *Sentinel = 0x00;
+  Q.submit([&](handler &cgh) {
+     cgh.parallel_for<class zero_zero_range_ndr>(
+         nd_range<1>{range<1>{0}, range<1>{0}},
+         [=](nd_item<1>) { *Sentinel = 0xFF; });
+   }).wait();
+  assert(*Sentinel == 0x00 &&
+         "parallel_for(nd_range{0, 0}) unexpectedly launched");
+
+  free(Sentinel, Q);
+  return 0;
+}