[SYCL] Improve handling of large reqd_work_group_size values #10620
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
jzc
changed the title
| uint64_t ReqdWGSizeAllDimsTotal = 1; | ||
| std::vector<uint64_t> ReqdWGSizeVec; | ||
| size_t ReqdWGSizeAllDimsTotal = 1; | ||
| std::vector<size_t> ReqdWGSizeVec; | ||
| int Dims = 0; | ||
| while (!ReqdWGSize.empty()) { | ||
| // The reqd_work_group_size data is stored as uint32_t's, |
There was a problem hiding this comment.
The comment should be updated that only old binaries store reqd_work_group_size as 32-bit.
We should also put TODOs to both runtime and the compiler to remove that extra _size_t property once ABI break is allowed again
sycl/test-e2e/OptionalKernelFeatures/large-reqd-work-group-size.cpp
Outdated
Show resolved
Hide resolved
AlexeySachkov
approved these changes
Aug 9, 2023
maarquitos14
reviewed
Aug 9, 2023
sycl/test-e2e/OptionalKernelFeatures/large-reqd-work-group-size.cpp
Outdated
Show resolved
Hide resolved
sycl/test-e2e/OptionalKernelFeatures/large-reqd-work-group-size.cpp
Outdated
Show resolved
Hide resolved
KornevNikita
reviewed
Aug 9, 2023
KornevNikita
approved these changes
Aug 9, 2023
|
@jzc can you please fix the failure in post-commit? https://github.com/intel/llvm/actions/runs/5819408460/job/15777743360 |
bader
pushed a commit
that referenced
this pull request
Aug 10, 2023
Post commit started to failed on 25c3666 (#10620) for macOS build because: ``` error: no viable overloaded '=' ReqdWorkGroupSize = NewReqdWorkGroupSize; note: candidate function not viable: no known conversion from 'llvm::SmallVector<size_t, 3>' (aka 'SmallVector<unsigned long, 3>') to 'const std::optional<llvm::SmallVector<unsigned long long, 3>>' for 1st argument _LIBCPP_INLINE_VISIBILITY optional& operator=(const optional&) = default;
mdtoguchi
pushed a commit
to mdtoguchi/llvm
that referenced
this pull request
Oct 18, 2023
…0620) 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)`.)
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
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 asycl::exceptionwitherrc::kernel_not_supportedaccording to 5.8.1. Kernel attributes).llvm/sycl/source/detail/program_manager/program_manager.cpp
Line 749 in b659690
This PR adds a check before building to validate the
reqd_work_group_size.Additionally, another bug in
sycl-post-linkwas 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 howsycl-post-linkworked, the attached metadata on the device image would include fewer dimensions than specified. (e.g. If the decorated withreqd_work_group_size(16, 16), then the device image would appear as if decorated byreqd_work_group_size(16).)