[NVPTX] Expand EXTLOAD for v8f16 and v8bf16 #72672
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
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
In triton-lang/triton#2483 I've encountered a bug in the NVPTX codegen. Given `load<8 x half>` followed by `fpext to <8 x float>` we get ``` ld.shared.v4.b16 {%f1, %f2, %f3, %f4}, [%r15+8]; ld.shared.v4.b16 {%f5, %f6, %f7, %f8}, [%r15]; ``` Which loads float16 values into float registers without any conversion and the result is simply garbage. This PR brings `v8f16` and `v8bf16` into line with the other vector types and expanding it to load + cvt.
peterbell10
force-pushed
the
nvptx-illegal-extload
branch
from
d1de0e2
to
f9bc0e6
Compare
|
Oh wow that's a bad bug. |
jlebar
approved these changes
Nov 17, 2023
sr-tream
pushed a commit
to sr-tream/llvm-project
that referenced
this pull request
Nov 20, 2023
In triton-lang/triton#2483 I've encountered a bug in the NVPTX codegen. Given `load<8 x half>` followed by `fpext to <8 x float>` we get ``` ld.shared.v4.b16 {%f1, %f2, %f3, %f4}, [%r15+8]; ld.shared.v4.b16 {%f5, %f6, %f7, %f8}, [%r15]; ``` Which loads float16 values into float registers without any conversion and the result is simply garbage. This PR brings `v8f16` and `v8bf16` into line with the other vector types by expanding it to load + cvt. cc @manman-ren @Artem-B @jlebar
zahiraam
pushed a commit
to zahiraam/llvm-project
that referenced
this pull request
Nov 20, 2023
In triton-lang/triton#2483 I've encountered a bug in the NVPTX codegen. Given `load<8 x half>` followed by `fpext to <8 x float>` we get ``` ld.shared.v4.b16 {%f1, %f2, %f3, %f4}, [%r15+8]; ld.shared.v4.b16 {%f5, %f6, %f7, %f8}, [%r15]; ``` Which loads float16 values into float registers without any conversion and the result is simply garbage. This PR brings `v8f16` and `v8bf16` into line with the other vector types by expanding it to load + cvt. cc @manman-ren @Artem-B @jlebar
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.
In triton-lang/triton#2483 I've encountered a bug in the NVPTX codegen. Given
load<8 x half>followed byfpext to <8 x float>we getWhich loads float16 values into float registers without any conversion and the result is simply garbage.
This PR brings
v8f16andv8bf16into line with the other vector types by expanding it to load + cvt.cc @manman-ren @Artem-B @jlebar