Skip to content

Commit 3e149ae

Browse files
committed
remove convert layout op which converts subgroup2d block to dpas
this is being handled by loadstoreoptollvm currently
1 parent e301b43 commit 3e149ae

File tree

1 file changed

+16
-0
lines changed

1 file changed

+16
-0
lines changed

third_party/intel/lib/TritonIntelGPUToLLVM/ConvertLayoutOpToLLVM.cpp

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,7 @@
33
#include "llvm/ADT/TypeSwitch.h"
44

55
#include "intel/include/Analysis/Utility.h"
6+
#include "intel/include/Dialect/TritonIntelGPU/Transforms/Utility.h"
67

78
namespace mlir::triton::gpu {
89
namespace {
@@ -27,6 +28,21 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
2728
auto srcTy = op.getSrc().getType();
2829
auto dstTy = op.getType();
2930

31+
if (auto srcTensorTy = cast<RankedTensorType>(srcTy)) {
32+
if (auto dstTensorTy = cast<RankedTensorType>(dstTy)) {
33+
// TODO: replace this with proper conversion once conversion is removed
34+
// from LoadStoreOpToLLVM.
35+
if (intel::hasSubgroup2DBlockEncoding(srcTensorTy) &&
36+
intel::hasDotDpasEncoding(dstTensorTy)) {
37+
// need to delete the op and do nothing
38+
llvm::errs() << "need to delete op " << op << "\n";
39+
// what if we just delete it
40+
rewriter.replaceOp(op, op.getSrc());
41+
return success();
42+
}
43+
}
44+
}
45+
3046
LinearLayout conversion = minimalCvtLayout(srcTy, dstTy);
3147
LinearLayout srcLayout =
3248
toLinearLayout(srcTy.getShape(), srcTy.getEncoding());

0 commit comments

Comments
 (0)