diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp index a5208d77976d..8f7e34eaea8c 100644 --- a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp +++ b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp @@ -25,6 +25,10 @@ using ::mlir::LLVM::linearize; using namespace mlir::triton::gpu; +// XXX(Keren): A temporary knob to control the use of legacy MMA conversion +// because LinearLayout seems to have some performance issues. +constexpr bool useLegacyMMAConversion = false; + struct ConvertLayoutOpConversion : public ConvertOpToLLVMPattern { public: @@ -378,6 +382,9 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion /*accumNumReplicates=*/1)) { return false; } + if (useLegacyMMAConversion) { + return false; + } return true; } if (isa(layout)) {