[Blackwell] Fix thrown away load due to wrong wait placement#9636
Open
LunNova wants to merge 1 commit intotriton-lang:mainfrom
Open
[Blackwell] Fix thrown away load due to wrong wait placement#9636LunNova wants to merge 1 commit intotriton-lang:mainfrom
LunNova wants to merge 1 commit intotriton-lang:mainfrom
Conversation
LunNova
force-pushed
the
lunnova/blackwell-wait-placement
branch
from
acc7490 to
b0a7b19
Compare
ThomasRaoux
reviewed
Mar 4, 2026
Comment on lines
+536
to
+560
| // Combine partial reductions into one value per thread. | ||
| // This must happen after tcgen05.wait since tcgen05.ld.red is async and | ||
| // the redval registers are not valid until the wait completes. | ||
| if (redvalVals.size() > 1) { | ||
| auto isMin = *redOp == TMEMLoadReduceModifier::MIN; | ||
| auto applyMinMax = [&](Value lhs, Value rhs) { | ||
| return useNaN | ||
| ? (isMin ? LLVM::MinimumOp::create(rewriter, loc, lhs, rhs) | ||
| : LLVM::MaximumOp::create(rewriter, loc, lhs, rhs)) | ||
| ->getResult(0) | ||
| : (isMin ? LLVM::MinNumOp::create(rewriter, loc, lhs, rhs) | ||
| : LLVM::MaxNumOp::create(rewriter, loc, lhs, rhs)) | ||
| ->getResult(0); | ||
| }; | ||
| // Use tree reduction: pair up elements at each level | ||
| while (redvalVals.size() > 1) { | ||
| SmallVector<Value> reduced; | ||
| assert(redvalVals.size() % 2 == 0 && | ||
| "redvalVals must be a multiple of 2"); | ||
| for (size_t i = 0; i < redvalVals.size(); i += 2) { | ||
| reduced.push_back(applyMinMax(redvalVals[i], redvalVals[i + 1])); | ||
| } | ||
| redvalVals = std::move(reduced); | ||
| } | ||
| } |
Collaborator
There was a problem hiding this comment.
nit: could you move that into a separate function?
lezcano
approved these changes
Mar 4, 2026
Contributor
lezcano
left a comment
lezcano
left a comment
There was a problem hiding this comment.
Right, this makes sense.
See Thomas' review, but otherwise LGTM
A faulty check line with no -COUNT- was masking bad codegen.
LunNova
force-pushed
the
lunnova/blackwell-wait-placement
branch
from
b0a7b19 to
d4fa9a9
Compare
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.
A faulty check line with no -COUNT- was masking bad codegen. I think.
Provisionally, fixes #9635. I am not very familiar with this area but trying my best! Feel free to close, edit with maintainer access or whatever.
New contributor declaration
I am not making a trivial change, such as fixing a typo in a comment.
I have written a PR description following these
rules.
I have run
pre-commit run --from-ref origin/main --to-ref HEAD.This PR does not need a test because
Changing existing broken test & codegen.I have not added any
littests.