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< 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..64ef519179a69 --- /dev/null +++ b/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 +#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"); + + // 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; +}