[SYCL] Fix SIGKILL after setting large WG sizes

[dm-vodopyanov]

This patch fixes SIGKILL (out of memory error) caused by large number
of global work size. Now, if work group size wasn't specified in the
kernel source or IL, WGSize sets to {1, 1, 1}.

[dm-vodopyanov]

/summary:run

[alexbatashev]

LGTM

[kbobrovs]

@bader, @dm-vodopyanov
I don't agree with the fix, as users frequently let RT figure out the WG size. Which means performance will likely suffer.

[bader]

If I understand correctly, this patch fixes functional problem. In general it's better to have slowly working application than broken application.

Do you have better solution in mind?

[kbobrovs]

In general it's better to have slowly working application than broken application.

it is hard to disagree. But the fixes can different. To me it seems the problem is not in the removed code, but in the device/kernel property APIs reporting incorrect values.
If some urgent W/A is needed, then I'd suggest to add hard WG size limit - e.g. 64. Or smaller value if 64 still causes the problem.
Then continue to investigate.

[bader]

If some urgent W/A is needed, then I'd suggest to add hard WG size limit - e.g. 64. Or smaller value if 64 still causes the problem.
Then continue to investigate.

That's exactly what Dmitry did, isn't it? He used 1 WG size limit instead of 64, which seems to be most portable value.

[dm-vodopyanov]

@kbobrovs I agree that this is controversial patch and performance may be lower. It reproduces only on OpenCL FPGA Emu on the specific machine. The minimum value from WGSize1D and MaxWGSizes[2] (in our case it is WGSize1D) is correct but we are running out of memory. It could happen on any machine someday. I find 1 the safest way to get rid of the problem. OpenCL FPGA Emu developers don't expect such huge number in WG size when calling clEnqueueNDRangeKernel. Anyway, I can contact with them again and double-check can OpenCL runtime be improved to handle such situation.

[kbobrovs]

That's exactly what Dmitry did, isn't it? He used 1 WG size limit instead of 64, which seems to be most portable value.

It definitely isn't. And it will likely make a difference for performance.

I agree that this is controversial patch and performance may be lower. It reproduces only on OpenCL FPGA Emu on the specific machine.

It seems that are we sacrificing performance on e.g. GPU hacking around potential FPGA Emu bug. Can at least device be checked if it is accelerator before using {1,1,1}? A TODO be added as well?

[alexbatashev]

We can move responsibility of picking the right WG size from SYCL RT to particular plugins. IIRC, OpenCL spec allows for such behavior. Other plugins may implement more platform-specific logic to achieve better performance.