From ad185393302edf84144b66bacac5a05a27975663 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 17 Aug 2026 15:14:55 -0700 Subject: [PATCH 1/5] narrow assert check --- sycl/source/detail/scheduler/commands.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index e71a8d2878db8..b05b5f377eef8 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/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< From 618e56bf8e5d6ce15c8f9ab790d39dcb5c2af9e6 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 18 Aug 2026 14:08:32 -0700 Subject: [PATCH 2/5] new test --- .../Basic/parallel_for_zero_range.cpp | 49 +++++++++++++++++++ 1 file changed, 49 insertions(+) create mode 100644 sycl/test-e2e/Basic/parallel_for_zero_range.cpp diff --git a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp new file mode 100644 index 0000000000000..eded32d47276d --- /dev/null +++ b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp @@ -0,0 +1,49 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// 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 +#include + +#include + +using namespace sycl; + +int main() { + queue Q; + + unsigned char *Sentinel = malloc_shared(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( + 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( + 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"); + + free(Sentinel, Q); + return 0; +} From 13a44d7d6252754b9f5d071b4571d37e3f0c1577 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 18 Aug 2026 14:09:29 -0700 Subject: [PATCH 3/5] new test --- sycl/test-e2e/Basic/parallel_for_zero_range.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp index eded32d47276d..c1901731cf2ee 100644 --- a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp +++ b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp @@ -5,7 +5,7 @@ // zero, the kernel function is not executed, the local size is ignored, and // any dependencies are satisfied." // -// See intel/llvm#22893 +// 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 @@ -27,8 +27,8 @@ int main() { // Case 1: parallel_for(range<1>{0}) -- plain empty range. *Sentinel = 0x00; Q.submit([&](handler &cgh) { - cgh.parallel_for( - range<1>{0}, [=](id<1>) { *Sentinel = 0xFF; }); + cgh.parallel_for(range<1>{0}, + [=](id<1>) { *Sentinel = 0xFF; }); }).wait(); assert(*Sentinel == 0x00 && "parallel_for(range{0}) unexpectedly launched"); From be328836c4f111b3508bd8ea9cd43b9066d0db5f Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 18 Aug 2026 14:14:21 -0700 Subject: [PATCH 4/5] clang-format and 0 0 case --- sycl/test-e2e/Basic/parallel_for_zero_range.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp index c1901731cf2ee..a44ed757e5f65 100644 --- a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp +++ b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp @@ -44,6 +44,16 @@ int main() { 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( + 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; } From 97ad7c882e3fff397234d76525fc11ce84c6ea85 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 19 Aug 2026 12:01:18 -0700 Subject: [PATCH 5/5] native-cpu exclusion --- sycl/test-e2e/Basic/parallel_for_zero_range.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp index a44ed757e5f65..64ef519179a69 100644 --- a/sycl/test-e2e/Basic/parallel_for_zero_range.cpp +++ b/sycl/test-e2e/Basic/parallel_for_zero_range.cpp @@ -1,6 +1,12 @@ // 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."