[NVPTX] Fix 64 bits rotations with large shift values #89399
Merged
Conversation
This file contains 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
Artem-B
approved these changes
Apr 19, 2024
| declare i64 @llvm.fshr.i64(i64, i64, i64) | ||
|
|
||
| ; SM35: rotl64 | ||
| define i64 @rotl64(i64 %a, i64 %n) { |
There was a problem hiding this comment.
This test should probably be converted to use llvm/utils/update_llc_test_checks.py
We do care about the arguments and the exact instruction sequences here.
There was a problem hiding this comment.
^^^ we still want to improve the test.
There was a problem hiding this comment.
Oh right! Updated with the script now, it was confusing for a little bit because it doesn't understand -march so I had to swap to --mtriple.
npmiller
force-pushed
the
fix-rot-shifts
branch
2 times, most recently
from
38eca33
to
138f196
Compare
|
@Artem-B I don't have commit permissions, would you mind landing this? I'll try to follow up with the suggestions when I have time. |
ROTL and ROTR can take a shift amount larger than the element size, in which case the effective shift amount should be the shift amount modulo the element size. This patch adds the modulo step when the shift amount isn't known at compile time. Without it the existing implementation would end up shifting beyond the type size and give incorrect results.
Artem-B
approved these changes
May 1, 2024
Comment on lines
+78
to
+86
| ; SM35-NEXT: ld.param.u64 %rd1, [rotate64_param_0]; | ||
| ; SM35-NEXT: { | ||
| ; SM35-NEXT: .reg .b32 %dummy; | ||
| ; SM35-NEXT: mov.b64 {%dummy,%r1}, %rd1; | ||
| ; SM35-NEXT: } | ||
| ; SM35-NEXT: { | ||
| ; SM35-NEXT: .reg .b32 %dummy; | ||
| ; SM35-NEXT: mov.b64 {%r2,%dummy}, %rd1; | ||
| ; SM35-NEXT: } |
There was a problem hiding this comment.
Looks like a minor optimization opportunity for the future -- this could be done as mov.b64 {%r2, %r1}, %rd1.
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.
ROTL and ROTR can take a shift amount larger than the element size, in which case the effective shift amount should be the shift amount modulo the element size.
This patch adds the modulo step when the shift amount isn't known at compile time. Without it the existing implementation would end up shifting beyond the type size and give incorrect results.
cc @Artem-B