llvm
25c3666d - [SYCL] Improve handling of large reqd_work_group_size values (#10620)

Commit
2 years ago
[SYCL] Improve handling of large reqd_work_group_size values (#10620) When a kernel with a large (unsupported) work-group size is submitted, the backend would fail upon trying build the program, and throw the wrong exception (compile_program_error with `errc::build`, but it should just be a `sycl::exception` with `errc::kernel_not_supported` according to [5.8.1. Kernel attributes](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.attributes)). https://github.com/intel/llvm/blob/b65969014f001f9730349a5caad5c2b85c9bc378/sycl/source/detail/program_manager/program_manager.cpp#L749 This PR adds a check before building to validate the `reqd_work_group_size`. Additionally, another bug in `sycl-post-link` was fixed. The bug occurred when a kernel was decorated with a multi-dimensional required work-group size and two dimensions had the same value. Due to how `sycl-post-link` worked, the attached metadata on the device image would include fewer dimensions than specified. (e.g. If the decorated with `reqd_work_group_size(16, 16)`, then the device image would appear as if decorated by `reqd_work_group_size(16)`.)
Author
Justin Cai
Parents
Loading