[optimize-dot-operands]: Fuse load and trans operations - part 2

third_party/intel/backend/compiler.py

@@ -280,6 +280,7 @@ def make_ttgir(mod, metadata, opt, properties):
         intel.passes.ttgpuir.add_accelerate_matmul(pm)
         intel.passes.ttgpuir.add_materialize_block_pointer(pm)
         intel.passes.ttgpuir.add_remove_layout_conversions(pm)
+        intel.passes.ttgpuir.add_optimize_dot_operands(pm)
         intel.passes.ttgpuir.add_pipeline(pm, opt.num_stages, XPUBackend.get_split_barrier_scope(opt))
 
         if (opt.reduce_variable_liveness):

third_party/intel/include/Utils/Utility.h

@@ -29,6 +29,9 @@ bool isConstant(Value val, int64_t expected);
 
 Value getFinalValue(Value value);
 
+// Erase the operations in \p operations.
+void eraseOperations(SmallPtrSetImpl<Operation *> &operations);
+
 } // namespace mlir::triton::intel
 
 #endif // TRITON_INTEL_UTILS_UTILITY_H

third_party/intel/lib/Dialect/Triton/Transforms/TensorDescToBlockPointer.cpp

@@ -95,7 +95,9 @@ struct TritonIntelTensorDescToBlockPointer
           .Default([&](auto) { return WalkResult::advance(); });
     });
 
-    finalize();
+    if (!cleanUp.empty())
+      tt::intel::eraseOperations(cleanUp);
+
     assert(succeeded(verify(moduleOp)) && "Module verification failed");
   }
 
@@ -267,31 +269,6 @@ struct TritonIntelTensorDescToBlockPointer
     return success();
   }
 
-  void finalize() {
-    // Cleanup unused operations.
-    bool erasedOperation;
-    do {
-      erasedOperation = false;
-      SmallPtrSet<Operation *, 8> erased;
-      for (Operation *op : cleanUp) {
-        if (!op->getUsers().empty() || !op->getRegions().empty())
-          continue;
-
-        erased.insert(op);
-        op->erase();
-        erasedOperation = true;
-      }
-      cleanUp.remove_if([&](Operation *op) { return erased.contains(op); });
-    } while (erasedOperation);
-
-    // Remove operations that contain a region.
-    for (Operation *op : cleanUp) {
-      if (!op->getUsers().empty())
-        continue;
-      op->erase();
-    }
-  }
-
 private:
   SmallPtrSet<Operation *, 8> cleanUp;
 };