[SYCL] Do not attach reqd_work_group_size info when multiple are detected

[jzc]

No description provided.

[dm-vodopyanov]

Not sure exactly, but can sycl::reqd_work_group_size(WGSIZE, WGSIZE) and sycl::reqd_work_group_size(WGSIZE) be moved to some template or function parameter (+ move dimension to the template parameter), to not duplicate the same code in both kernel_launch_* functions.

[jzc]

Unfortunately, it seems attributes cannot accept parameter packs, so the best we can do is the macro version:

#define TEST(...)                                                              \
  {                                                                            \
    range globalRange(__VA_ARGS__);                                            \
    range localRange(__VA_ARGS__);                                             \
    nd_range NDRange(globalRange, localRange);                                 \
    q.parallel_for(NDRange,                                                    \
                   [=](auto) [[sycl::reqd_work_group_size(__VA_ARGS__)]] {});  \
  }

[AlexeySachkov]

Unfortunately, it seems attributes cannot accept parameter packs

We should create a tracker for that, because it is possible to add such support and we have done it already for other attributes

[jzc]

Created an issue: #13686

[AlexeySachkov]

There should be a comment explaining this. My understanding is that strictly speaking, we expect only one value of reqd_work_group_size metadata because of per-optional-kernel-feature device code split that had supposed to happen before.

However, there is an exception when device code split is disabled, which causes kernels with different reqd_work_group_size requirements to be bundled together. I think that ideally we want to assert here that device code split is disabled and that otherwise MultipleReqdWGSize is false, but I'm not sure if we have access to that knowledge here.

[jzc]

However, there is an exception when device code split is disabled, which causes kernels with different reqd_work_group_size requirements to be bundled together. I think that ideally we want to assert here that device code split is disabled and that otherwise MultipleReqdWGSize is false, but I'm not sure if we have access to that knowledge here.

I agree, but yea, the device code split mode information is not present in this function. If we want to go that far, I think it makes sense to add it as a parameter to the function.

[LU-JOHN]

On line 61 can we add a check for !MultipleReqdWGSize? There is no point in checking again if we already know multiple WG sizes are required.

Also can we call Reqs.ReqdWorkGroupSize.reset() after line 69 to keep all the code related to ReqdWorkGroupSize together?

[LU-JOHN]

Are we allowed to discard a requirement just because they are contradictory? From https://intel.github.io/llvm-docs/design/OptionalDeviceFeatures.html:

For a kernel that is decorated with the [[sycl::reqd_work_group_size(W)]] or [[sycl::reqd_sub_group_size(S)]] attribute, the exception must be thrown if the device does not support the work group size W or the sub-group size S.

We can't honor this requirement if we discard reqd_work_group_size.

[AlexeySachkov]

Are we allowed to discard a requirement just because they are contradictory? From https://intel.github.io/llvm-docs/design/OptionalDeviceFeatures.html:

For a kernel that is decorated with the [[sycl::reqd_work_group_size(W)]] or [[sycl::reqd_sub_group_size(S)]] attribute, the exception must be thrown if the device does not support the work group size W or the sub-group size S.

We can't honor this requirement if we discard reqd_work_group_size.

The situation when we discard that metadata and therefore lose ability to emit that error can only happen when a user explicitly specifies non-standard -fsycl-device-code-split=off. We have not claimed to be fully conformant with the SYCL specification with that flag.

Essentially this is a trade-off between user experience and being conformant. The problem with user experience we had is that we also have a check that local size passed to parallel_for is the same as what is attached as an attribute to a kernel. Since we record an attribute on a per-device-image basis assuming that it is the same for all kernels, this caused false alarms, fully preventing users from launching any kernels. Disabled device code split path is essentially a default for FPGA devices and therefore we decided to go this way.