Skip to content

Commit

Permalink
[PIPELINE] Including mmav5 dot in shmem pipelining (#5835)
Browse files Browse the repository at this point in the history
There is a bug introduced in: [PIPELINE] Relax requirements for wgmma
operand register pipelining (#5810)
I was missing mmav5 dot support. Thanks for @peterbell10 for catching
that and for @masahi for proposing a fix.
  • Loading branch information
pawelszczerbuk authored Feb 5, 2025
1 parent c7924e8 commit 414abc4
Showing 1 changed file with 3 additions and 2 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -523,8 +523,9 @@ assignMemoryLayouts(scf::ForOp &forOp,
if (isa<mlir::triton::DotOpInterface>(use)) {
LDBG("set shared encoding with dot user: " << *use);
auto dot = dyn_cast<tt::DotOp>(use);
auto warpGroupDot = dyn_cast<ttng::WarpGroupDotOp>(use);
mmav3Shmem = canUseMMAv3Pipelining(&op) && warpGroupDot;
bool isMMAv3v5Dot = isa<ttng::WarpGroupDotOp, ttng::TCGen5MMAOp,
ttng::TCGen5MMAScaledOp>(use);
mmav3Shmem = canUseMMAv3Pipelining(&op) && isMMAv3v5Dot;

loadInfo.usedByDot = true;
loadInfo.isMMAv3Shared = mmav3Shmem;
Expand Down

0 comments on commit 414abc4

Please sign in to comment.