Introduce -miopen-lowering-step5 pass.

llvm · Jun 6, 2020 · 45d62cc · 45d62cc
1 parent 9488261
commit 45d62cc
Show file tree

Hide file tree

Showing 3 changed files with 22 additions and 1 deletion.
diff --git a/mlir/include/mlir/Dialect/MIOpen/Passes.h b/mlir/include/mlir/Dialect/MIOpen/Passes.h
@@ -36,6 +36,9 @@ std::unique_ptr<Pass> createLowerMIOpenOpsStep3Pass();
 /// Create a pass to convert MIOpen threadwise operations to other dialects.
 std::unique_ptr<Pass> createLowerMIOpenOpsStep4Pass();
 
+/// Create a pass to convert affine / loop to std dialect.
+std::unique_ptr<Pass> createLowerMIOpenOpsStep5Pass();
+
 /// Create a pass to convert transform operations to affine maps.
 std::unique_ptr<Pass> createAffineTransformPass();
 

diff --git a/mlir/include/mlir/Dialect/MIOpen/Passes.td b/mlir/include/mlir/Dialect/MIOpen/Passes.td
@@ -37,7 +37,12 @@ def MIOpenOpsStep3Pass : Pass<"miopen-lowering-step3", "ModuleOp"> {
 }
 
 def MIOpenOpsStep4Pass : Pass<"miopen-lowering-step4", "ModuleOp"> {
-  let summary = "expand threadwise copy and threadwise gemm to all other dialects. Notice GPU dialect will explicitly NOT be used in this pass";
+  let summary = "expand threadwise copy and threadwise gemm to all other dialects.";
   let constructor = "mlir::miopen::createLowerMIOpenOpsStep4Pass()";
 }
+
+def MIOpenOpsStep5Pass : Pass<"miopen-lowering-step5", "ModuleOp"> {
+  let summary = "expand loop / affine dialects to std. Notice GPU dialect will explicitly NOT be used in this pass";
+  let constructor = "mlir::miopen::createLowerMIOpenOpsStep5Pass()";
+}
 #endif // MLIR_DIALECT_MIOPEN_PASSES
diff --git a/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp b/mlir/lib/Dialect/MIOpen/Transforms/LowerMIOpenOps.cpp
@@ -64,6 +64,11 @@ struct LowerMIOpenOpsStep4Pass
     : public MIOpenOpsStep4PassBase<LowerMIOpenOpsStep4Pass> {
   void runOnOperation() override;
 };
+
+struct LowerMIOpenOpsStep5Pass
+    : public MIOpenOpsStep5PassBase<LowerMIOpenOpsStep5Pass> {
+  void runOnOperation() override;
+};
 } // end anonymous namespace
 
 void LowerMIOpenOpsStep1Pass::runOnOperation() {
@@ -94,7 +99,11 @@ void LowerMIOpenOpsStep4Pass::runOnOperation() {
   OwningRewritePatternList patterns;
   patterns.insert<ThreadwiseGemmRewritePattern>(&getContext());
   patterns.insert<ThreadwiseCopyRewritePattern>(&getContext());
+  applyPatternsAndFoldGreedily(getOperation(), patterns);
+}
 
+void LowerMIOpenOpsStep5Pass::runOnOperation() {
+  OwningRewritePatternList patterns;
   populateAffineToStdConversionPatterns(patterns, &getContext());
   populateLoopToStdConversionPatterns(patterns, &getContext());
   applyPatternsAndFoldGreedily(getOperation(), patterns);
@@ -115,3 +124,7 @@ std::unique_ptr<Pass> mlir::miopen::createLowerMIOpenOpsStep3Pass() {
 std::unique_ptr<Pass> mlir::miopen::createLowerMIOpenOpsStep4Pass() {
   return std::make_unique<LowerMIOpenOpsStep4Pass>();
 }
+
+std::unique_ptr<Pass> mlir::miopen::createLowerMIOpenOpsStep5Pass() {
+  return std::make_unique<LowerMIOpenOpsStep5Pass>();
+}