Features to support tilelang vmi fa

include/PTO/IR/PTOInterfaces.td

@@ -76,6 +76,26 @@ def MteOpInterface : OpInterface<"MteOpInterface"> {
     operations and raw micro-instructions may both carry this marker.
   }];
   let cppNamespace = "::mlir::pto";
+  let methods = [
+    InterfaceMethod<"Return the source buffer.",
+      "::mlir::Value", "getSource", (ins), [{}],
+      /*defaultImplementation=*/[{
+        auto operation = $_op.getOperation();
+        if (auto dps = ::llvm::dyn_cast<::mlir::DestinationStyleOpInterface>(
+                operation))
+          return dps.getDpsInputOperand(0)->get();
+        return operation->getOperand(0);
+      }]>,
+    InterfaceMethod<"Return the destination buffer.",
+      "::mlir::Value", "getDestination", (ins), [{}],
+      /*defaultImplementation=*/[{
+        auto operation = $_op.getOperation();
+        if (auto dps = ::llvm::dyn_cast<::mlir::DestinationStyleOpInterface>(
+                operation))
+          return dps.getDpsInitOperand(0)->get();
+        return operation->getOperand(1);
+      }]>
+  ];
 }
 
 def SimtOpInterface : OpInterface<"SimtOpInterface"> {

include/PTO/IR/VMIOps.td

@@ -40,7 +40,8 @@ def PTO_PhysicalVMIPartTypeConstraint : AnyTypeOf<
   "PTO physical vector register or mask type">;
 
 class VMI_Op<string mnemonic, list<Trait> traits = []>
-    : PTO_Op<"vmi." # mnemonic, traits>;
+    : PTO_Op<"vmi." # mnemonic,
+             !listconcat(traits, [MicroOpInterface, VectorMicroOpInterface])>;
 
 //===--- Legacy (old) VMI ops ---===//
 

include/PTO/IR/VPTOOps.td

@@ -309,7 +309,8 @@ def TileBufAddrOp : PTO_Op<"tile_buf_addr", [Pure]> {
   }];
 }
 
-def VecScopeOp : PTO_Op<"vecscope", [SingleBlock, NoTerminator]> {
+def VecScopeOp : PTO_Op<"vecscope", [SingleBlock, NoTerminator,
+                                      VectorMicroOpInterface]> {
   let summary = "Structured region container for one VPTO vector scope";
   let description = [{
     `pto.vecscope` marks a structured vector-scope interval without overloading
@@ -324,7 +325,8 @@ def VecScopeOp : PTO_Op<"vecscope", [SingleBlock, NoTerminator]> {
 }
 
 def StrictVecScopeOp : PTO_Op<"strict_vecscope", [SingleBlock, NoTerminator,
-                                                  IsolatedFromAbove]> {
+                                                   IsolatedFromAbove,
+                                                   VectorMicroOpInterface]> {
   let summary = "Structured VPTO vector scope with explicit captures only";
   let description = [{
     `pto.strict_vecscope` is the strict form of `pto.vecscope`. Values used by

include/PTO/Transforms/Passes.td

@@ -355,7 +355,7 @@ def PTONormalizeUncoveredTileSections
 
     This pass is the shared entry point for inferred-section normalization:
     explicit sections remain in place, while uncovered top-level segments that
-    contain tile operations are identified for later wrapping.
+    contain section-relevant operations are identified for later wrapping.
   }];
 
   let constructor = "mlir::pto::createPTONormalizeUncoveredTileSectionsPass()";

lib/PTO/IR/PTO.cpp

@@ -3589,11 +3589,15 @@ LogicalResult mlir::pto::SyncSetOp::verify() {
     }
     switch (getPipe().getPipe()) {
     case PIPE::PIPE_FIX:
+    case PIPE::PIPE_MTE1:
+    case PIPE::PIPE_MTE2:
     case PIPE::PIPE_MTE3:
+    case PIPE::PIPE_V:
       return success();
     default:
-      return emitOpError()
-             << "A5 sync.set expects pipe to be one of <PIPE_FIX>, <PIPE_MTE3>";
+      return emitOpError() << "A5 sync.set expects pipe to be one of "
+                              "<PIPE_FIX>, <PIPE_MTE1>, <PIPE_MTE2>, "
+                              "<PIPE_MTE3>, <PIPE_V>";
     }
   };
   return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5);