From a1daaceb56d6b869509417284c30382f7379a7dd Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Wed, 14 Oct 2020 22:43:23 +0300 Subject: [PATCH] [SYCL] Fix SIGKILL after setting large WG sizes 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}`. --- sycl/source/detail/scheduler/commands.cpp | 15 +-------------- 1 file changed, 1 insertion(+), 14 deletions(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 83af59acebdd5..c6646252fdafa 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1620,20 +1620,7 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, RT::PiKernel Kernel, get(Kernel, DeviceImpl.getHandleRef(), DeviceImpl.getPlugin()); if (WGSize[0] == 0) { - // kernel does not request specific workgroup shape - set one - id<3> MaxWGSizes = - get_device_info, cl::sycl::info::device::max_work_item_sizes>:: - get(DeviceImpl.getHandleRef(), DeviceImpl.getPlugin()); - - size_t WGSize1D = get_kernel_device_specific_info< - size_t, cl::sycl::info::kernel_device_specific::work_group_size>:: - get(Kernel, DeviceImpl.getHandleRef(), DeviceImpl.getPlugin()); - - assert(MaxWGSizes[2] != 0); - - // Set default work-group size in the Z-direction to either the max - // number of work-items or the maximum work-group size in the Z-direction. - WGSize = {1, 1, min(WGSize1D, MaxWGSizes[2])}; + WGSize = {1, 1, 1}; } NDR.set(NDR.Dims, nd_range<3>(NDR.NumWorkGroups * WGSize, WGSize)); }