[mlir][gpu] Separate the barrier elimination code from transform ops (#71762)

Allows the barrier elimination code to be run from C++ as well. The code
from transforms dialect is copied as-is, the pass and populate functions
have beed added at the end.

Co-authored-by: Eric Eaton <[email protected]>

Author: spaceotter

mlir/include/mlir/Dialect/GPU/Transforms/Passes.h

@@ -165,6 +165,9 @@ void populateGpuDecomposeMemrefsPatterns(RewritePatternSet &patterns);
 /// Pass decomposes memref ops inside `gpu.launch` body.
 std::unique_ptr<Pass> createGpuDecomposeMemrefsPass();
 
+/// Erase barriers that do not enforce conflicting memory side effects.
+void populateGpuEliminateBarriersPatterns(RewritePatternSet &patterns);
+
 /// Generate the code for registering passes.
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/GPU/Transforms/Passes.h.inc"

mlir/include/mlir/Dialect/GPU/Transforms/Passes.td

@@ -37,6 +37,24 @@ def GpuMapParallelLoopsPass
   let dependentDialects = ["mlir::gpu::GPUDialect"];
 }
 
+def GpuEliminateBarriers
+    : Pass<"gpu-eliminate-barriers", "mlir::func::FuncOp"> {
+  let summary = "Erase unnecessary barriers";
+  let description = [{
+    Barrier elimination pass. If a barrier does not enforce any conflicting
+    pair of memory effects, including a pair that is enforced by another
+    barrier, it is unnecessary and can be removed. Adapted from
+    "High-Performance GPU-to-CPU Transpilation and Optimization via High-Level
+    Parallel Constructs" by Moses, Ivanov, Domke, Endo, Doerfert, and Zinenko in
+    PPoPP 2023 and implementation in Polygeist.
+  }];
+  let dependentDialects = [
+    "mlir::gpu::GPUDialect",
+    "mlir::memref::MemRefDialect",
+    "mlir::scf::SCFDialect"
+  ];
+}
+
 def GpuDecomposeMemrefsPass : Pass<"gpu-decompose-memrefs"> {
   let summary = "Decomposes memref index computation into explicit ops.";
   let description = [{

mlir/lib/Dialect/GPU/CMakeLists.txt

@@ -62,6 +62,7 @@ add_mlir_dialect_library(MLIRGPUTransforms
   Transforms/ShuffleRewriter.cpp
   Transforms/SPIRVAttachTarget.cpp
   Transforms/ROCDLAttachTarget.cpp
+  Transforms/EliminateBarriers.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/GPU