[Pipeline] Update the pipeline (#928)

- Enable new blocking pass and remove the old one
- Remove the PropagationLayout Pass

include/imex/Dialect/XeTile/Transforms/Passes.h

@@ -40,23 +40,13 @@ std::unique_ptr<mlir::Pass> createXeTileInitDuplicatePass();
 
 std::unique_ptr<mlir::Pass>
 createXeTileBlockingPass(const std::string &device = "pvc");
-std::unique_ptr<mlir::Pass>
-createNewXeTileBlockingPass(const std::string &device = "pvc");
-std::unique_ptr<mlir::Pass> createXeTileBlockAligningPass();
 std::unique_ptr<mlir::Pass> createXeTileWgToSgPass();
 std::unique_ptr<mlir::Pass> createXeTileCanonicalizationPass();
 
 ///
 void populateXeTileInitDuplicatePatterns(imex::XeTypeConverter &converter,
                                          mlir::RewritePatternSet &patterns);
 
-///
-void populateXeTileBlockingPatterns(imex::XeTypeConverter &converter,
-                                    mlir::RewritePatternSet &patterns,
-                                    std::shared_ptr<XeuArchInterface> ptruArch);
-
-#define GEN_PASS_DECL_NEWXETILEBLOCKING
-#define GEN_PASS_DECL_XETILEBLOCKALIGNING
 #define GEN_PASS_DECL_XETILEBLOCKING
 #define GEN_PASS_DECL_XETILECANONICALIZATION
 #define GEN_PASS_DECL_XETILEINITDUPLICATE

include/imex/Dialect/XeTile/Transforms/Passes.td

@@ -34,53 +34,6 @@ def XeTileInitDuplicate : Pass<"xetile-init-duplicate", "::mlir::gpu::GPUModuleO
   let dependentDialects = ["imex::xetile::XeTileDialect"];
 }
 
-def XeTileBlocking : Pass<"xetile-blocking", "::mlir::gpu::GPUModuleOp">{
-  let summary = "transform XeTile large tiles(input) into arrays of smaller "
-                "blocks with appropriate size, such that the operator on each "
-                "of the blocks can be mapped into one hardware instruction.";
-
-  let description = [{
-    This transform pass preprocesses the xetile program by decomposing large XeTile tiles
-    into smaller ones that can be handled by a hardware instruction. This blocked layout
-    is represented by high dimension vectors, inner dimension matches to DPAS size config.
-    This lowers 2D vector to 4D vector.
-  }];
-
-  let constructor = "imex::createXeTileBlockingPass()";
-  let dependentDialects = ["imex::xetile::XeTileDialect",
-                           "mlir::arith::ArithDialect",
-                           "mlir::gpu::GPUDialect",
-                           "mlir::memref::MemRefDialect",
-                           "mlir::vector::VectorDialect"];
-
-  let options = [
-     Option<"device", "device", "std::string",
-            /*default=*/"\"pvc\"",
-            "gpu platform architecture where these ops are running">
- ];
-}
-
-
-// TODO: [block-aligning] remove the following code when upstreaming the pass.
-// The pass is not supposed to be exposed to users. Temporary keep in case we
-// need debug it for down stream development.
-def XeTileBlockAligning: Pass <"xetile-block-aligning", "::mlir::gpu::GPUModuleOp"> {
-  let summary = "optimize the performance for mma.";
-
-  let description = [{
-    This transform is to optimize performance by aligning the block size among operators
-    to reduce in-register data movements. Currently, it mainly focues on the alignment
-    between load and MMA operators.
-  }];
-
-  let constructor = "imex::createXeTileBlockAligningPass()";
-  let dependentDialects = ["imex::xetile::XeTileDialect",
-                           "mlir::arith::ArithDialect",
-                           "mlir::gpu::GPUDialect",
-                           "mlir::memref::MemRefDialect",
-                           "mlir::vector::VectorDialect"];
-}
-
 def XeTileWgToSg : Pass<"xetile-wg-to-sg", "::mlir::gpu::GPUModuleOp">{
   let summary = "Transform WG level XeTile code to SG XeTile";
 
@@ -116,7 +69,7 @@ def XeTileCanonicalization : Pass<"xetile-canonicalization", "::mlir::gpu::GPUMo
   ];
 }
 
-def NewXeTileBlocking : Pass<"new-xetile-blocking", "::mlir::gpu::GPUModuleOp">{
+def XeTileBlocking : Pass<"xetile-blocking", "::mlir::gpu::GPUModuleOp">{
   let summary = "transform XeTile large tiles(input) into arrays of smaller "
                 "blocks with appropriate size, such that the operator on each "
                 "of the blocks can be mapped into one hardware instruction.";
@@ -127,7 +80,7 @@ def NewXeTileBlocking : Pass<"new-xetile-blocking", "::mlir::gpu::GPUModuleOp">{
     the xetile-blocking pass.
   }];
 
-  let constructor = "imex::createNewXeTileBlockingPass()";
+  let constructor = "imex::createXeTileBlockingPass()";
   let dependentDialects = ["imex::xetile::XeTileDialect",
                            "mlir::arith::ArithDialect",
                            "mlir::math::MathDialect",

include/imex/Transforms/Passes.h

@@ -34,7 +34,6 @@ std::unique_ptr<mlir::Pass> createBF16ToGPUPass();
 std::unique_ptr<mlir::Pass> createCastIndexPass();
 std::unique_ptr<mlir::Pass> createRemoveTemporariesPass();
 std::unique_ptr<mlir::Pass> createVectorLinearizePass();
-std::unique_ptr<mlir::Pass> createPropagatePackedLayoutPass();
 std::unique_ptr<mlir::Pass> createRemoveSingleElemVectorPass();
 std::unique_ptr<mlir::Pass>
 createOptimizeTransposePass(const std::string &device = "pvc");

include/imex/Transforms/Passes.td

@@ -159,15 +159,6 @@ def VectorLinearize : Pass<"imex-vector-linearize"> {
   ];
 }
 
-def PropagatePackedLayout : Pass<"imex-propagate-packed-layout"> {
-  let summary = "Propagate packed layout (i.e. VNNI) through ops";
-  let constructor = "imex::createPropagatePackedLayoutPass()";
-
-  let dependentDialects = [
-    "::mlir::vector::VectorDialect"
-  ];
-}
-
 def RemoveSingleElemVector : Pass<"imex-remove-single-elem-vector"> {
   let summary = "Remove <1xT> vectors, and replace them with scalars.";
   let constructor = "imex::createRemoveSingleElemVectorPass()";