[SYCL] extends previous Global/LocalSize 0 fix - #23065
Conversation
There was a problem hiding this comment.
Warning
Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.
Pull request overview
Extends the SYCL “zero global/local range” handling to avoid passing nullptr-like launch parameters into downstream layers (notably DeviceASAN), while preserving SYCL 2020 semantics that kernels with empty ranges are not executed.
Changes:
- Refines
adjustNDRangePerKernelearly-return logic to distinguish “global/local explicitly set” vs “both zero”. - Ensures a non-zero local work-group size is available for no-op launches where both global and local are zero, preventing sanitizer failures.
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| if (NDR.GlobalSize[0] != 0) | ||
| return; // GlobalSize is set - no need to adjust | ||
| if (NDR.LocalSize[0] != 0) | ||
| return; // User set LocalSize but GlobalSize is zero (e.g. nd_range with | ||
| // zero global, non-zero local). Per SYCL 2020 the kernel is not | ||
| // executed; leave the range as-is. |
| // Zero global and zero local: either parallel_for_work_group (NumWorkGroups | ||
| // is set) or plain parallel_for with an empty range. Fill in WGSize so that | ||
| // downstream layers (in particular DeviceASAN's preLaunchKernel) always see | ||
| // a non-zero local work size even when the launch is a no-op. |
There was a problem hiding this comment.
🔵 Needs a closer look
Multi-dimensional ranges with zeroes outside dimension 0 remain unsafe.
Review details
Suppressed comments (1)
sycl/source/detail/scheduler/commands.cpp:2300
- These checks only inspect dimension 0, so a 2D/3D empty range whose zero is in another dimension is not fixed. For example,
parallel_for(range<2>{16, 0}, ...)returns here with a null local size, whilend_range<2>{{16, 0}, {4, 0}}passes a zero local component; DeviceASAN subsequently divides by every local component. Check all active dimensions so every zero-sized range gets a non-zero downstream local size.
if (NDR.GlobalSize[0] != 0)
return; // GlobalSize is set - no need to adjust
if (NDR.LocalSize[0] != 0)
- Files reviewed: 1/1 changed files
- Comments generated: 0 new
- Review effort level: Balanced
The SYCL spec allows 0 size global and local ranges. The kernels are not run, but any dependencies are left in place.
This was fixed in #22957 but the fix allows a nullptr to sneak through. By itself not a problem, but when the tests are run with address sanitization, we get failures. This fixes that. Existing tests cover all branches. No need for additional testing.