[SYCL] global size zero assertion narrowed - #22957
Conversation
| // 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); |
There was a problem hiding this comment.
From https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_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."
But the modified code doesn't ignore the local size: it checks it. And having a check wrapped in an assertion gives a problem that behavior will be different depending on how SYCL runtime was compiled. We can't assume that release builds will be done with disabled assertions. On the contrary, most if not all Linux distributions build projects with enabled assertions.
There was a problem hiding this comment.
I believe for the case nd_range(0, N) we won't reach this assert. check on line 2295 guarantees that for the all cases except parallel_for_work_group we exit early.
There was a problem hiding this comment.
I added a test and it includes this case, of both being zero.
KseniyaTikhomirova
left a comment
KseniyaTikhomirova
left a comment
There was a problem hiding this comment.
Chris, could you please add a test for this? Since we haven't met this failure/assert before - it means we don't test this scenario.
It would be great to check that we send correct values to L0 and the backend handles zero global size correctly.
| 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 |
There was a problem hiding this comment.
But non-zero LocalSize was coming from pytorch. This statement does not seem to be true.
There was a problem hiding this comment.
Per @KseniyaTikhomirova explanation this holds true due to modified if condition in the beginning of the funciton.
| // 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); |
There was a problem hiding this comment.
Previously we had assertion firing with NDR.LocalSize[0] != 0 which was coming from pytorch. Per my understanding the new assert will also fire. I will check this later today, but I have concerns if this change really fixes the reported issue.
There was a problem hiding this comment.
Tried. Yes, I am getting assertion fired with the modified assertion (line number are different as I applied the change to sycl-rel-7_1 branch):
third_party/torch-xpu-ops/test/xpu/extended/test_ops_xpu.py python3: /home/dvrogozh/git/intel-llvm/sycl/source/detail/scheduler/commands.cpp:2323: void sycl::_V1::detail::adjustNDRangePerKernel(NDRDescT&, ur_kernel_handle_t, const device_impl&): Assertion `NDR.GlobalSize[0] == 0 && NDR.LocalSize[0] == 0' failed.
Fatal Python error: Aborted
If assertion can not be narrowed or we don't fully understand why/when execution may reach this place, maybe it would be better to just remove assert entirely till implementation will be clarified?
There was a problem hiding this comment.
@KseniyaTikhomirova pointed out that I have overlooked the condition change in the very beginning of this functions right before the assert. After properly cherry-picking the PR it works for the pytorch. Sorry for confusion.
dvrogozh
left a comment
dvrogozh
left a comment
There was a problem hiding this comment.
LGTM
Please, after the merge consider to cherry-pick this fix into sycl-rel-7_0 and sycl-rel-7_1. @KornevNikita.
3 participants
The SYCL 2020 spec states that "When the global size is zero, the kernel function is not executed, the local size is ignored, and any dependencies are satisfied"
But we have an assert that is a bit too wide and is tripping up PyTorch. Fix is simply to narrow the assertion.