[NFC][SYCL] Replace #pragma unroll with dim_loop in accessor.hpp #6939
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
The utility was introduced in intel#6560 because "#pragma unroll" doesn't always work and template-based solution is much more reliable. Original PR only changed the loops that resulted in immediate performance difference but other occurrences were missed. This PR updates remaining ones. Note that I've found them by looking into the LLVM IR produced by our device compiler and having the loop really unrolled improves readability of such dumps (and most likely codesize/perf, although not significantly).
|
@intel/llvm-gatekeepers PR is ready. |
whitneywhtsang
added a commit
to whitneywhtsang/llvm
that referenced
this pull request
Oct 28, 2022
…hpp (intel#6939)" This reverts commit fee486e.
whitneywhtsang
added a commit
to whitneywhtsang/llvm
that referenced
this pull request
Nov 3, 2022
…ccessor.hpp (intel#6939)"" This reverts commit 98f50ae.
bader
pushed a commit
that referenced
this pull request
Feb 22, 2023
…backend. (#7948) A performance regression was reported when using `reduce_over_group` with sycl::vec. This was due to a loop over calls to the scalar `reduce_over_group` for each of the `sycl::vec` components that was not unrolled and led to register spills even at -O3. It was initially possible to fix the performance by calling `#pragma unroll` and declare `reduce_over_group` with `__attribute__((always_inline))`. However the `SYCL_UNROLL` macro that calls `#pragma unroll` has been removed in favour of `dim_loop` (#6939). I have used dim_loop to fix the loop unrolling. However, in the cuda backend, just using `dim_loop` in this way actually makes the performance worse. This is because `dim_loop` introduces new non inlined function calls in the cuda backend that lead to register spills. The solution to this coincides with the solution of several user reports that the cuda backend is not aggressive enough with inlining. In this PR I have also therefore increased the inlining threshold multiplier value to 11. See https://reviews.llvm.org/D142232/new/ for the corresponding upstream PR (for the inlining threshold change) that includes much more details on benchmarking dpc++ cuda with this change. In short, for dpc++ cuda, there is no other downside apart from a very small increase in compile time in some cases, but there is a massive benefit to increasing the inlining threshold across a large amount of applications. Testing using opencl cpu backend reveals that this code change has no effect on this backend. This change is required for the cuda backend but should have no performance effect for other backends. fixes #6583. --------- Signed-off-by: JackAKirk <[email protected]> Co-authored-by: JackAKirk <[email protected]>
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.
The utility was introduced in #6560 because "#pragma unroll" doesn't always work and template-based solution is much more reliable. Original PR only changed the loops that resulted in immediate performance difference but other occurrences were missed. This PR updates remaining ones. Note that I've found them by looking into the LLVM IR produced by our device compiler and having the loop really unrolled improves readability of such dumps (and most likely codesize/perf, although not significantly).