Skip to content

Commit 4d76de4

Browse files
[SYCL] Fix SIGKILL after setting large WG sizes (#2641)
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}`.
1 parent e7492fb commit 4d76de4

File tree

1 file changed

+1
-14
lines changed

1 file changed

+1
-14
lines changed

sycl/source/detail/scheduler/commands.cpp

Lines changed: 1 addition & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1620,20 +1620,7 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, RT::PiKernel Kernel,
16201620
get(Kernel, DeviceImpl.getHandleRef(), DeviceImpl.getPlugin());
16211621

16221622
if (WGSize[0] == 0) {
1623-
// kernel does not request specific workgroup shape - set one
1624-
id<3> MaxWGSizes =
1625-
get_device_info<id<3>, cl::sycl::info::device::max_work_item_sizes>::
1626-
get(DeviceImpl.getHandleRef(), DeviceImpl.getPlugin());
1627-
1628-
size_t WGSize1D = get_kernel_device_specific_info<
1629-
size_t, cl::sycl::info::kernel_device_specific::work_group_size>::
1630-
get(Kernel, DeviceImpl.getHandleRef(), DeviceImpl.getPlugin());
1631-
1632-
assert(MaxWGSizes[2] != 0);
1633-
1634-
// Set default work-group size in the Z-direction to either the max
1635-
// number of work-items or the maximum work-group size in the Z-direction.
1636-
WGSize = {1, 1, min(WGSize1D, MaxWGSizes[2])};
1623+
WGSize = {1, 1, 1};
16371624
}
16381625
NDR.set(NDR.Dims, nd_range<3>(NDR.NumWorkGroups * WGSize, WGSize));
16391626
}

0 commit comments

Comments
 (0)