Merge commit '6116bfe6e08c5bfca49869b31c4e32983221d6be'

Author: whitneywhtsang

bench/triton_bench/matmul_ogs_details/_finalize_scatter.py

@@ -59,7 +59,7 @@ void registerTestTritonAMDGPURangeAnalysis();
 
 inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::registerAllPasses();
-  mlir::registerTritonPasses();
+  mlir::triton::registerTritonPasses();
   mlir::triton::gpu::registerTritonGPUPasses();
   mlir::registerTritonNvidiaGPUPasses();
   mlir::test::intel::registerTestAxisInfoPass();

bench/triton_bench/matmul_ogs_details/_finalize_split_k.py

@@ -7,10 +7,6 @@
 #include "llvm/ADT/SetVector.h"
 #include "llvm/Support/raw_ostream.h"
 
-#include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGPU/IR/Dialect.h"
-#include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"
-#include <atomic>
 #include <limits>
 
 namespace mlir {

bin/RegisterTritonDialects.h

@@ -6,12 +6,8 @@
 
 #include "mlir/Support/LLVM.h"
 #include "triton/Analysis/Utility.h"
-#include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/Triton/IR/Utility.h"
-#include "triton/Dialect/TritonGPU/IR/Dialect.h"
 
 #include <optional>
-#include <type_traits>
 
 namespace mlir::triton {
 

include/triton/Analysis/Allocation.h

@@ -2,7 +2,6 @@
 #define TRITON_ANALYSIS_MEMBAR_H
 
 #include "Allocation.h"
-#include "llvm/ADT/SmallPtrSet.h"
 
 #include <set>
 

include/triton/Analysis/AxisInfo.h

@@ -6,18 +6,14 @@
 namespace mlir {
 namespace triton {
 
-std::unique_ptr<Pass> createCombineOpsPass();
-
-std::unique_ptr<Pass> createLoopInvariantCodeMotionPass();
-std::unique_ptr<Pass> createReorderBroadcastPass();
-std::unique_ptr<Pass> createRewriteTensorPointerPass();
-std::unique_ptr<Pass> createLoopUnrollPass();
-
-} // namespace triton
+// Generate the pass class declarations.
+#define GEN_PASS_DECL
+#include "triton/Dialect/Triton/Transforms/Passes.h.inc"
 
 #define GEN_PASS_REGISTRATION
 #include "triton/Dialect/Triton/Transforms/Passes.h.inc"
 
+} // namespace triton
 } // namespace mlir
 
 #endif

include/triton/Analysis/Membar.h

@@ -19,8 +19,6 @@ def TritonCombineOps : Pass</*cli-arg*/"triton-combine", /*Op*/"mlir::ModuleOp">
        => dot(x,y,splat(0))`
   }];
 
-  let constructor = "mlir::triton::createCombineOpsPass()";
-
   let dependentDialects = ["mlir::arith::ArithDialect"];
 }
 
@@ -33,7 +31,7 @@ def TritonReorderBroadcast : Pass</*cli-arg*/"triton-reorder-broadcast", /*Op*/"
     In the event of a match, the broadcast (or splat) operation is delayed
     and performed after the ElementWise operation.
   }];
-  let constructor = "mlir::triton::createReorderBroadcastPass()";
+
   let dependentDialects = ["mlir::triton::TritonDialect"];
 }
 
@@ -45,8 +43,6 @@ def TritonRewriteTensorPointer : Pass</*cli-arg*/"triton-rewrite-tensor-pointer"
     the pointer/mask/other for each load/store.
   }];
 
-  let constructor = "mlir::triton::createRewriteTensorPointerPass()";
-
   let dependentDialects = ["mlir::triton::TritonDialect"];
 }
 
@@ -56,7 +52,7 @@ def TritonLoopUnroll : Pass</*cli-arg*/"triton-loop-unroll", /*Op*/"mlir::Module
     The pass unrolls a scf loop with tt.loop_unroll_factor attribute. The attribute specialises how many iterations
     the loop should be unrolled.
   }];
-  let constructor = "mlir::triton::createLoopUnrollPass()";
+
   let dependentDialects = ["mlir::triton::TritonDialect"];
 }
 
@@ -68,7 +64,7 @@ def TritonLoopInvariantCodeMotion : Pass</*cli-arg*/"triton-licm", /*Op*/"mlir::
     generates a trip-count check. For scf.while loops, it clones the condition
     from the before body.
   }];
-  let constructor = "mlir::triton::createLoopInvariantCodeMotionPass()";
+
   let dependentDialects = ["mlir::triton::TritonDialect"];
 }
 

include/triton/Dialect/Triton/Transforms/Passes.h

@@ -1,10 +1,7 @@
 #include "triton/Analysis/Alias.h"
 
-#include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Support/LLVM.h"
-#include "triton/Analysis/Utility.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
-#include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"
 
 namespace mlir {
 

include/triton/Dialect/Triton/Transforms/Passes.td

@@ -2,10 +2,8 @@
 
 #include <algorithm>
 #include <limits>
-#include <numeric>
 
 #include "mlir/Analysis/Liveness.h"
-#include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Support/LLVM.h"
 #include "triton/Analysis/Alias.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"