[hipblaslt] Remove swap and dot2 from TF32-cvt sequence#913
Merged
Conversation
b-shi
changed the title
jonatluu
pushed a commit
to jonatluu/rocm-libraries
that referenced
this pull request
Jul 28, 2025
[ROCm/hipBLAS commit: bcac487]
jonatluu
pushed a commit
to jonatluu/rocm-libraries
that referenced
this pull request
Jul 28, 2025
b-shi
force-pushed
the
users/brianshi/tf32_cvt
branch
from
eccb4f1 to
7f6db7f
Compare
msujon-AMD
approved these changes
Jul 31, 2025
Collaborator
|
@b-shi, do we want to merge it when the PLR PR is ready and is about to be merged? |
Contributor
Author
Either would work, but I would prefer to merge this when I can. |
AlexBrownAMD
approved these changes
Aug 1, 2025
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Aug 1, 2025
[hipblaslt] Remove swap and dot2 from TF32-cvt sequence (#913) Removed `v_swap_b32` and `v_dot2c_f32_bf16` usage in TF32 cvt sequence. - swap can be removed by reordering instruction sequence - dot2 was removed because it does not interleave well with mfmas (there is quite a large penalty issuing dot2 after mfmas) - This will impact perf when PLR [changes](ROCm/rocm-libraries#593) get merged. - Instruction sequence per pack increases by 2, (`24->26`)
b-shi
added a commit
that referenced
this pull request
Aug 7, 2025
Due to changes in #913 `s_mov_b64` was introduced. This requires wait states for all cases (previously was just added for SourceSwap = 1).
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Aug 7, 2025
[hipblaslt] Fix hazard in TF32 cvt sequence Due to changes in ROCm/rocm-libraries#913 `s_mov_b64` was introduced. This requires wait states for all cases (previously was just added for SourceSwap = 1).
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
Removed
v_swap_b32andv_dot2c_f32_bf16usage in TF32 cvt sequence.24->26)