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

[Oleksandr Alex Zinenko]

================
@@ -37,6 +37,20 @@ def GpuMapParallelLoopsPass
   let dependentDialects = ["mlir::gpu::GPUDialect"];
 }
 
+def GpuEliminateBarriers
+    : Pass<"gpu-eliminate-barriers", "mlir::func::FuncOp"> {
+  let summary = "Erase unecessary barriers";
+  let constructor = "mlir::createGpuEliminateBarriersPass()";
+  let description = [{
+    Erase barriers that are unneeded because there are no dependent side effects
+    across the barrier.
----------------
ftynse wrote:

Carrying over the documentation from the pattern may help the reader understand how this is implemented.

```
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.
```

https://github.com/llvm/llvm-project/pull/71762