From 75e5ddf55489afea99b571e8de5aee7a553130ab Mon Sep 17 00:00:00 2001
From: Jokeren <robinho364@gmail.com>
Date: Thu, 22 Aug 2024 20:24:22 -0400
Subject: [PATCH 1/2] Update

---
 lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
index a5208d77976d..6e8e3acbc4b6 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 perf issues.
+constexpr bool useLegacyMMAConversion = false;
+
 struct ConvertLayoutOpConversion
     : public ConvertOpToLLVMPattern<ConvertLayoutOp> {
 public:
@@ -378,6 +382,9 @@ struct ConvertLayoutOpUsingLinearLayoutsConversion
                                       /*accumNumReplicates=*/1)) {
           return false;
         }
+        if (useLegacyMMAConversion) {
+          return false;
+        }
         return true;
       }
       if (isa<BlockedEncodingAttr>(layout)) {

From 23e2ba86c0357cdcb1e121b6ce3ff0ad7fe099ee Mon Sep 17 00:00:00 2001
From: Jokeren <robinho364@gmail.com>
Date: Thu, 22 Aug 2024 22:01:36 -0400
Subject: [PATCH 2/2] Update

---
 lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
index 6e8e3acbc4b6..8f7e34eaea8c 100644
--- a/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
+++ b/lib/Conversion/TritonGPUToLLVM/ConvertLayoutOpToLLVM.cpp
@@ -26,7 +26,7 @@ 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 perf issues.
+// because LinearLayout seems to have some performance issues.
 constexpr bool useLegacyMMAConversion = false;
 
 struct ConvertLayoutOpConversion