Skip to content

Comments

[WIP] gfx1250 scales load fusion#816

Open
Hardcode84 wants to merge 14 commits intoiree-org:mainfrom
Hardcode84:scales-fusion-mlir
Open

[WIP] gfx1250 scales load fusion#816
Hardcode84 wants to merge 14 commits intoiree-org:mainfrom
Hardcode84:scales-fusion-mlir

Conversation

@Hardcode84
Copy link
Contributor

No description provided.

@Hardcode84 Hardcode84 requested a review from panditsa February 3, 2026 21:07
// Create fused load: select(lane_id < waveSize/2, ptrA, ptrB).
Location loc = op.getLoc();
OpBuilder::InsertionGuard guard(rewriter);
rewriter.setInsertionPoint(loadA);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

-    rewriter.setInsertionPoint(loadA);
+    Operation *insertionPoint = loadA->isBeforeInBlock(loadB) ? loadB : loadA;
+    rewriter.setInsertionPoint(insertionPoint);

This was a change I needed to make for a larger gemm test.

"::mlir::gpu::GPUDialect",
];
let options = [
Option<"waveSize", "wave-size", "unsigned", /*default=*/"32",
Copy link
Contributor

@tgymnich tgymnich Feb 4, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

it does not make sense to have a default value here. The wave size should always be set intentionally and correctly, there is no real default that always works.

elemOps.push_back(value.getDefiningOp());
value = value.getDefiningOp()->getOperand(0);
}
std::reverse(elemOps.begin(), elemOps.end());
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
std::reverse(elemOps.begin(), elemOps.end());
llvm::reverse(elemOps);

Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Signed-off-by: Ivan Butygin <ivan.butygin@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants