Skip to content

[OPTIMIZER][BACKEND] Enabled elementwise ops (including casts) between ldmatrix and mma.sync #1595

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ptillet merged 21 commits into from
May 1, 2023
Merged

[OPTIMIZER][BACKEND] Enabled elementwise ops (including casts) between ldmatrix and mma.sync #1595

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
62f094b
.
ptillet Apr 28, 2023
7a78edf
.
ptillet Apr 28, 2023
9b87e3d
.
ptillet Apr 28, 2023
42ae7cd
progress
ptillet Apr 29, 2023
1fd119b
.
ptillet Apr 29, 2023
e140335
.
ptillet Apr 29, 2023
659ea72
Merge branch 'main' into phil/better-float8-pipeline
ptillet Apr 29, 2023
abb7ed3
.
ptillet Apr 29, 2023
90fb2d0
Merge branch 'main' into phil/better-float8-pipeline
ptillet Apr 29, 2023
4adacee
progress
ptillet Apr 29, 2023
cfd484d
.
ptillet Apr 30, 2023
1bf8899
.
ptillet Apr 30, 2023
c781f47
added lit test
ptillet Apr 30, 2023
9408904
more progress
ptillet Apr 30, 2023
1955de1
stronger prefetch test
ptillet Apr 30, 2023
9347cb3
.
ptillet Apr 30, 2023
ed536f7
.
ptillet Apr 30, 2023
4abb910
Merge branch 'main' into phil/better-float8-pipeline
ptillet Apr 30, 2023
24fa9d9
.
ptillet Apr 30, 2023
39b0330
.
ptillet May 1, 2023
560ec0d
.
ptillet May 1, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 7 additions & 7 deletions lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -620,10 +620,10 @@ struct ConvertLayoutOpConversion
// is implemented
SmallVector<Value> reorderedVals;
for (unsigned i = 0; i < vecVals.size(); i += 4) {
reorderedVals.push_back(vecVals[i]);
reorderedVals.push_back(vecVals[i + 2]);
reorderedVals.push_back(vecVals[i + 1]);
reorderedVals.push_back(vecVals[i + 3]);
reorderedVals.push_back(bitcast(vecVals[i], i32_ty));
reorderedVals.push_back(bitcast(vecVals[i + 2], i32_ty));
reorderedVals.push_back(bitcast(vecVals[i + 1], i32_ty));
reorderedVals.push_back(bitcast(vecVals[i + 3], i32_ty));
}

Value view = getTypeConverter()->packLLElements(loc, reorderedVals,
Expand All @@ -642,19 +642,19 @@ struct ConvertLayoutOpConversion
auto loc = op.getLoc();
Value src = op.getSrc();
Value dst = op.getResult();
bool isHMMA = supportMMA(dst, mmaLayout.getVersionMajor());

auto smemObj =
getSharedMemoryObjectFromStruct(loc, adaptor.getSrc(), rewriter);
Value res;

if (!isOuter && mmaLayout.isAmpere() && isHMMA) { // tensor core v2
if (!isOuter && mmaLayout.isAmpere()) { // tensor core v2

res = SharedToDotOperandMMAv2::convertLayout(
dotOperandLayout.getOpIdx(), rewriter, loc, src, dotOperandLayout,
smemObj, getTypeConverter(), tid_val());

} else if (!isOuter && mmaLayout.isVolta() && isHMMA) { // tensor core v1
} else if (!isOuter && mmaLayout.isVolta() &&
supportMMA(dst, mmaLayout.getVersionMajor())) { // tensor core v1
bool isMMAv1Row = dotOperandLayout.getMMAv1IsRow();
auto srcSharedLayout = src.getType()
.cast<RankedTensorType>()
Expand Down
Loading