[MLIR] Don't sort operand of commutative ops when comparing two ops as there is a correctness issue

[tomnatan]

Summary

This feature was introduced in D123492.

Doing equivalence on pointers to sort operands of commutative operations is incorrect when checking equivalence of ops in separate regions (where the lhs and rhs operands are marked as equivalent but are not the same value).

It was also discussed in D123492 and D129480 that the correct solution would be to stable sort the operands in canonicalization (based on some numbering in the region maybe), but until that lands, reverting this change will unblock us and other users.

An example of a pass that might not work properly because of this is DuplicateFunctionEliminationPass.

[tomnatan]

Fixed comment

[tomnatan]

Fix/remove failing tests

[tomnatan]

fix flang test

[clementval]

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

flang/test/Fir/commute.fir

11

I would prefer to keep the test as is and mark it as fail until we have another solution in place,

[tomnatan]

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

Only that test or all tests I've removed? how do you mark a test as fails? Just in the name and remove the checks?

[clementval]

flang/test/Fir/commute.fir

11

At least this one.

There is a way to tell lit we expect this to fail. You should add // XFAIL: * after the run line.

[tomnatan]

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

[clementval]

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

If you can share smth it would be nice. Even an edited version where we can see where the problem arise.

[tomnatan]

flang/test/Fir/commute.fir

11

Doesn't this mean that all tests in the fail are expected to fail whereas here we only except one of them?

[clementval]

flang/test/Fir/commute.fir

11

This file is treated as a single test so it would be fine.

[mehdi_amini]

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

I believe the way to expose this in a test upstream is to write a test pass that exposes the issue. That could be a pass that takes every functions in the input module and do pair-wise operation equivalence between them, and returns true/false.

Something I'm unsure about is if the failure would be deterministic: since we use pointer values this seems like it'll depend on the actual allocation...

[tomnatan]

In D154699#4481303, @mehdi_amini wrote:

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

I believe the way to expose this in a test upstream is to write a test pass that exposes the issue. That could be a pass that takes every functions in the input module and do pair-wise operation equivalence between them, and returns true/false.

Something I'm unsure about is if the failure would be deterministic: since we use pointer values this seems like it'll depend on the actual allocation...

Ended up finding a very simple reproducer:

func.func @main(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = func.call @f1(%arg0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  %1 = func.call @f2(%0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  return %1 : tensor<4x8xf32>
}
func.func @f1(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}
func.func @f2(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}

If you run mlir-opt %s --duplicate-function-elimination it will dedup the functions only with this revert.

I can include this test case in duplicate-function-elimination.mlir

[clementval]

In D154699#4481345, @tomnatan wrote:

In D154699#4481303, @mehdi_amini wrote:

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

I believe the way to expose this in a test upstream is to write a test pass that exposes the issue. That could be a pass that takes every functions in the input module and do pair-wise operation equivalence between them, and returns true/false.

Something I'm unsure about is if the failure would be deterministic: since we use pointer values this seems like it'll depend on the actual allocation...

Ended up finding a very simple reproducer:

func.func @main(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = func.call @f1(%arg0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  %1 = func.call @f2(%0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  return %1 : tensor<4x8xf32>
}
func.func @f1(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}
func.func @f2(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}

If you run mlir-opt %s --duplicate-function-elimination it will dedup the functions only with this revert.

I can include this test case in duplicate-function-elimination.mlir

Thanks for the example. It is useful. So if reverted your dedub example works but CSE will not work anymore in example like the flang one. Is that correct?

[tomnatan]

In D154699#4481380, @clementval wrote:

In D154699#4481345, @tomnatan wrote:

In D154699#4481303, @mehdi_amini wrote:

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

I believe the way to expose this in a test upstream is to write a test pass that exposes the issue. That could be a pass that takes every functions in the input module and do pair-wise operation equivalence between them, and returns true/false.

Something I'm unsure about is if the failure would be deterministic: since we use pointer values this seems like it'll depend on the actual allocation...

Ended up finding a very simple reproducer:

func.func @main(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = func.call @f1(%arg0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  %1 = func.call @f2(%0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  return %1 : tensor<4x8xf32>
}
func.func @f1(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}
func.func @f2(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}

If you run mlir-opt %s --duplicate-function-elimination it will dedup the functions only with this revert.

I can include this test case in duplicate-function-elimination.mlir

Thanks for the example. It is useful. So if reverted your dedub example works but CSE will not work anymore in example like the flang one. Is that correct?

Yes exactly, but keeping it as is means the isRegionEquivalentTo and IsEqivalentTo can return incorrect results for regions/ops that are actually identical (not just equivalent).

[clementval]

In D154699#4481394, @tomnatan wrote:

In D154699#4481380, @clementval wrote:

In D154699#4481345, @tomnatan wrote:

In D154699#4481303, @mehdi_amini wrote:

In D154699#4481235, @tomnatan wrote:

In D154699#4481215, @clementval wrote:

Can we have an example where it fails for you? I'm not saying that we should not revert but it would be good to have an real world example because what was discussed in other patch is downstream AFAIK.

The example is an internal Google mlir dump, I'm not sure if I can post it publicly as not all dialects are open sourced. But it's essentially two very big but identical regions for which IsRegionEquivalentTo returns false, but returns true with this change.

I believe the way to expose this in a test upstream is to write a test pass that exposes the issue. That could be a pass that takes every functions in the input module and do pair-wise operation equivalence between them, and returns true/false.

Something I'm unsure about is if the failure would be deterministic: since we use pointer values this seems like it'll depend on the actual allocation...

Ended up finding a very simple reproducer:

func.func @main(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = func.call @f1(%arg0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  %1 = func.call @f2(%0, %arg1) : (tensor<4x8xf32>, tensor<4x8xf32>) -> tensor<4x8xf32>
  return %1 : tensor<4x8xf32>
}
func.func @f1(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}
func.func @f2(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> tensor<4x8xf32> {
  %0 = mhlo.add %arg0, %arg1 : tensor<4x8xf32>
  %1 = mhlo.add %0, %arg1 : tensor<4x8xf32>
  %2 = mhlo.add %1, %1 : tensor<4x8xf32>
  %3 = mhlo.add %2, %1 : tensor<4x8xf32>
  return %3 : tensor<4x8xf32>
}

If you run mlir-opt %s --duplicate-function-elimination it will dedup the functions only with this revert.

I can include this test case in duplicate-function-elimination.mlir

Thanks for the example. It is useful. So if reverted your dedub example works but CSE will not work anymore in example like the flang one. Is that correct?

Yes exactly, but keeping it as is means the isRegionEquivalentTo and IsEqivalentTo can return incorrect results for regions/ops that are actually identical (not just equivalent).

I think there was also a suggestion to have different algo for CSE and isRegionEquivalentTo. But that would require work anyway.

[tomnatan]

change fir test

[tomnatan]

Is there anything else blocking us from submitting this CL? I think we should first fix the correctness issue and then move the sorting of operands to the beginning of CSE as suggested in D129480.

flang/test/Fir/commute.fir

11

I changed the test to have the add op with reversed operands, wouldn't it be better than just having XFAIL? i.e. this test will start failing if we introduce a proper fix for commutative ops, and to fix it we just need to remove the extra add in the result.

[clementval]

flang/test/Fir/commute.fir

11

That's what I was suggesting at first. Keep the test as it was and mark it XFAIL.

[tomnatan]

flang/test/Fir/commute.fir

11

Not sure I follow why we need to mark it as XFAIL as it currently passes because I added the additional add op that isn't csed and when we support commutative sorting it will fail because of the additional add op. Does that makes sense?

[clementval]

flang/test/Fir/commute.fir

11

My suggestion is to not alter this test as we want it to be like that when CSE works for commutative op. Marking it as XFAIL with no other modification is like saying this is a TODO.

[tomnatan]

flang/test/Fir/commute.fir

11

But does this test actually need to depend on the TODO to support this? What's the harm of updating is as I did and later having to revert back when and if commutative is supported properly? And also what about @f2 here? We will be hiding future failures it.

[clementval]

flang/test/Fir/commute.fir

11

What's the point to have a test for commutative op CSE that doesn't work as intended. I would prefer to mark it as XFAIL.

[tomnatan]

revert change in flang test and add XFILE instead

[tomnatan]

flang/test/Fir/commute.fir

11

I see your point, done. Also added https://github.com/llvm/llvm-project/issues/63784 for tracking.

[tomnatan]

Is there anything else blocking us from merging this change? We continue to discuss a proper solution separately.

[clementval]

In D154699#4492973, @tomnatan wrote:

Is there anything else blocking us from merging this change? We continue to discuss a proper solution separately.

@tomnatan Do you have any plan on working to add the proper solution?

[tomnatan]

In D154699#4494467, @clementval wrote:

In D154699#4492973, @tomnatan wrote:

Is there anything else blocking us from merging this change? We continue to discuss a proper solution separately.

@tomnatan Do you have any plan on working to add the proper solution?

I'm afraid I can't commit to taking this on as I won't have time to do this in the near future. But I think we should first fix the correctness issue and then prioritize a solution.

[tomnatan]

@clementval can I go ahead and merge the change?

[clementval]

In D154699#4501335, @tomnatan wrote:

@clementval can I go ahead and merge the change?

Can you change the title of this commit before landing it. It's not a pure revert of D123492.

[tomnatan]

In D154699#4502309, @clementval wrote:

In D154699#4501335, @tomnatan wrote:

@clementval can I go ahead and merge the change?

Can you change the title of this commit before landing it. It's not a pure revert of D123492.

Done.

flang/test/Fir/commute.fir

 // RUN: fir-opt %s | tco | FileCheck %s
+//
+// XFAIL:*
+// See: https://github.com/llvm/llvm-project/issues/63784
 
 // CHECK-LABEL: define i32 @f1(i32 %0, i32 %1)
 func.func @f1(%a : i32, %b : i32) -> i32 {
 
   // CHECK: %[[reg3:.*]] = add i32 %0, %1
   %1 = arith.addi %a, %b : i32
   %2 = arith.addi %b, %a : i32
   // CHECK:  mul i32 %[[reg3]], %[[reg3]]
   %3 = arith.muli %1, %2 : i32
   return %3 : i32
 }
 
 // CHECK-LABEL: define i32 @f2(ptr %0)
 func.func @f2(%a : !fir.ref<i32>) -> i32 {
   %1 = fir.load %a : !fir.ref<i32>

mlir/lib/IR/OperationSupport.cpp

llvm::hash_code hash =
       llvm::hash_combine(op->getName(), op->getDiscardableAttrDictionary(),
                          op->getResultTypes(), op->hashProperties());
 
   //   - Location if required
   if (!(flags & Flags::IgnoreLocations))
     hash = llvm::hash_combine(hash, op->getLoc());
 
   //   - Operands
-  ValueRange operands = op->getOperands();
-  SmallVector<Value> operandStorage;
-  if (op->hasTrait<mlir::OpTrait::IsCommutative>()) {
-    operandStorage.append(operands.begin(), operands.end());
-    llvm::sort(operandStorage, [](Value a, Value b) -> bool {
-      return a.getAsOpaquePointer() < b.getAsOpaquePointer();
-    });
-    operands = operandStorage;
-  }
-  for (Value operand : operands)
+  for (Value operand : op->getOperands())
     hash = llvm::hash_combine(hash, hashOperands(operand));
 
-  //   - Operands
+  //   - Results
   for (Value result : op->getResults())
     hash = llvm::hash_combine(hash, hashResults(result));
   return hash;
 }
 
 /*static*/ bool OperationEquivalence::isRegionEquivalentTo(
     Region *lhs, Region *rhs,
     function_ref<LogicalResult(Value, Value)> checkEquivalent,
if (lhs->getName() != rhs->getName() ||
       lhs->getNumOperands() != rhs->getNumOperands() ||
       lhs->getNumResults() != rhs->getNumResults() ||
       lhs->hashProperties() != rhs->hashProperties())
     return false;
   if (!(flags & IgnoreLocations) && lhs->getLoc() != rhs->getLoc())
     return false;
 
   // 2. Compare operands.
-  ValueRange lhsOperands = lhs->getOperands(), rhsOperands = rhs->getOperands();
-  SmallVector<Value> lhsOperandStorage, rhsOperandStorage;
-  if (lhs->hasTrait<mlir::OpTrait::IsCommutative>()) {
-    auto sortValues = [](ValueRange values) {
-      SmallVector<Value> sortedValues = llvm::to_vector(values);
-      llvm::sort(sortedValues, [](Value a, Value b) {
-        auto aArg = llvm::dyn_cast<BlockArgument>(a);
-        auto bArg = llvm::dyn_cast<BlockArgument>(b);
-
-        // Case 1. Both `a` and `b` are `BlockArgument`s.
-        if (aArg && bArg) {
-          if (aArg.getParentBlock() == bArg.getParentBlock())
-            return aArg.getArgNumber() < bArg.getArgNumber();
-          return aArg.getParentBlock() < bArg.getParentBlock();
-        }
-
-        // Case 2. One of then is a `BlockArgument` and other is not. Treat
-        // `BlockArgument` as lesser.
-        if (aArg && !bArg)
-          return true;
-        if (bArg && !aArg)
-          return false;
-
-        // Case 3. Both are values.
-        return a.getAsOpaquePointer() < b.getAsOpaquePointer();
-      });
-      return sortedValues;
-    };
-    lhsOperandStorage = sortValues(lhsOperands);
-    lhsOperands = lhsOperandStorage;
-    rhsOperandStorage = sortValues(rhsOperands);
-    rhsOperands = rhsOperandStorage;
-  }
-
-  for (auto operandPair : llvm::zip(lhsOperands, rhsOperands)) {
+  for (auto operandPair : llvm::zip(lhs->getOperands(), rhs->getOperands())) {
     Value curArg = std::get<0>(operandPair);
     Value otherArg = std::get<1>(operandPair);
     if (curArg == otherArg)
       continue;
     if (curArg.getType() != otherArg.getType())
       return false;
     if (failed(checkEquivalent(curArg, otherArg)))
       return false;

mlir/test/Dialect/Func/duplicate-function-elimination.mlir

func.func @user(%arg0: f32, %arg1: f32) -> f32 {
   %1 = call @also_add_lr(%arg0, %arg1) : (f32, f32) -> f32
   %2 = call @add_rl(%0, %1) : (f32, f32) -> f32
   %3 = call @also_add_rl(%arg0, %2) : (f32, f32) -> f32
  return %3 : f32
 }
 
 // CHECK:     @add_lr
 // CHECK-NOT: @also_add_lr
-// CHECK-NOT: @add_rl
+// CHECK:     @add_rl
 // CHECK-NOT: @also_add_rl
 // CHECK:     @user
-// CHECK-4:     call @add_lr
+// CHECK-2:     call @add_lr
+// CHECK-2:     call @add_rl
 
 // -----
 
 func.func @ite(%pred: i1, %then: f32, %else: f32) -> f32 {
   %0 = scf.if %pred -> f32 {
     scf.yield %then : f32
   } else {
     scf.yield %else : f32

 // CHECK-NOT: @also_ite
 // CHECK:     @reverse_ite
 // CHECK:     @user
 // CHECK-2:     call @ite
 // CHECK:       call @reverse_ite
 
 // -----
 
-func.func @deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32, %odd: f32) 
+func.func @deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32, %odd: f32)
     -> f32 {
   %0 = scf.if %p0 -> f32 {
     %1 = scf.if %p1 -> f32 {
       %2 = scf.if %p2 -> f32 {
         %3 = scf.if %p3 -> f32 {
           scf.yield %even : f32
         } else {
           scf.yield %odd : f32
%1 = scf.if %p1 -> f32 {
       }
       scf.yield %2 : f32
     }
     scf.yield %1 : f32
   }
   return %0 : f32
 }
 
-func.func @also_deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32, 
+func.func @also_deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32,
     %odd: f32) -> f32 {
   %0 = scf.if %p0 -> f32 {
     %1 = scf.if %p1 -> f32 {
       %2 = scf.if %p2 -> f32 {
         %3 = scf.if %p3 -> f32 {
           scf.yield %even : f32
         } else {
           scf.yield %odd : f32
%1 = scf.if %p1 -> f32 {
       }
       scf.yield %2 : f32
     }
     scf.yield %1 : f32
   }
   return %0 : f32
 }
 
-func.func @reverse_deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32, 
+func.func @reverse_deep_tree(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %even: f32,
     %odd: f32) -> f32 {
   %0 = scf.if %p0 -> f32 {
     %1 = scf.if %p1 -> f32 {
       %2 = scf.if %p2 -> f32 {
         %3 = scf.if %p3 -> f32 {
           scf.yield %odd : f32
         } else {
           scf.yield %even : f32
%1 = scf.if %p1 -> f32 {
       }
       scf.yield %2 : f32
     }
     scf.yield %1 : f32
   }
   return %0 : f32
 }
 
-func.func @user(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %odd: f32, %even: f32) 
+func.func @user(%p0: i1, %p1: i1, %p2: i1, %p3: i1, %odd: f32, %even: f32)
     -> (f32, f32, f32) {
-  %0 = call @deep_tree(%p0, %p1, %p2, %p3, %odd, %even) 
+  %0 = call @deep_tree(%p0, %p1, %p2, %p3, %odd, %even)
       : (i1, i1, i1, i1, f32, f32) -> f32
-  %1 = call @also_deep_tree(%p0, %p1, %p2, %p3, %odd, %even) 
+  %1 = call @also_deep_tree(%p0, %p1, %p2, %p3, %odd, %even)
       : (i1, i1, i1, i1, f32, f32) -> f32
-  %2 = call @reverse_deep_tree(%p0, %p1, %p2, %p3, %odd, %even) 
+  %2 = call @reverse_deep_tree(%p0, %p1, %p2, %p3, %odd, %even)
       : (i1, i1, i1, i1, f32, f32) -> f32
   return %0, %1, %2 : f32, f32, f32
 }
 
 // CHECK:     @deep_tree
 // CHECK-NOT: @also_deep_tree
 // CHECK:     @reverse_deep_tree
 // CHECK:     @user
 // CHECK-2:     call @deep_tree
 // CHECK:       call @reverse_deep_tree

mlir/test/Transforms/cse.mlir

func.func @dont_remove_duplicated_read_op_with_sideeffecting() -> i32 {
   "test.op_with_memwrite"() : () -> ()
   // CHECK: %[[READ_VALUE1:.*]] = "test.op_with_memread"() : () -> i32
   %1 = "test.op_with_memread"() : () -> (i32)
   // CHECK-NEXT: %{{.*}} = arith.addi %[[READ_VALUE0]], %[[READ_VALUE1]] : i32
   %2 = arith.addi %0, %1 : i32
   return %2 : i32
 }
 
-/// This test is checking that identical commutative operation are gracefully
-/// handled but the CSE pass.
-// CHECK-LABEL: func @check_cummutative_cse
-func.func @check_cummutative_cse(%a : i32, %b : i32) -> i32 {
-  // CHECK: %[[ADD1:.*]] = arith.addi %{{.*}}, %{{.*}} : i32
-  %1 = arith.addi %a, %b : i32
-  %2 = arith.addi %b, %a : i32
-  // CHECK-NEXT:  arith.muli %[[ADD1]], %[[ADD1]] : i32
-  %3 = arith.muli %1, %2 : i32
-  return %3 : i32
-}
-
 // Check that an operation with a single region can CSE.
 func.func @cse_single_block_ops(%a : tensor<?x?xf32>, %b : tensor<?x?xf32>)
   -> (tensor<?x?xf32>, tensor<?x?xf32>) {
   %0 = test.cse_of_single_block_op inputs(%a, %b) {
     ^bb0(%arg0 : f32):
     test.region_yield %arg0 : f32
   } : tensor<?x?xf32>, tensor<?x?xf32> -> tensor<?x?xf32>
   %1 = test.cse_of_single_block_op inputs(%a, %b) {
func.func @no_cse_single_block_ops_different_bodies(%a : tensor<?x?xf32>, %b : tensor<?x?xf32>, %c : f32, %d : i1)
   } : tensor<?x?xf32>, tensor<?x?xf32> -> tensor<?x?xf32>
   return %0, %1 : tensor<?x?xf32>, tensor<?x?xf32>
 }
 // CHECK-LABEL: func @no_cse_single_block_ops_different_bodies
 //       CHECK:   %[[OP0:.+]] = test.cse_of_single_block_op
 //       CHECK:   %[[OP1:.+]] = test.cse_of_single_block_op
 //       CHECK:   return %[[OP0]], %[[OP1]]
 
-// Account for commutative ops within regions during CSE.
-func.func @cse_single_block_with_commutative_ops(%a : tensor<?x?xf32>, %b : tensor<?x?xf32>, %c : f32)
-  -> (tensor<?x?xf32>, tensor<?x?xf32>) {
-  %0 = test.cse_of_single_block_op inputs(%a, %b) {
-    ^bb0(%arg0 : f32, %arg1 : f32):
-    %1 = arith.addf %arg0, %arg1 : f32
-    %2 = arith.mulf %1, %c : f32
-    test.region_yield %2 : f32
-  } : tensor<?x?xf32>, tensor<?x?xf32> -> tensor<?x?xf32>
-  %1 = test.cse_of_single_block_op inputs(%a, %b) {
-    ^bb0(%arg0 : f32, %arg1 : f32):
-    %1 = arith.addf %arg1, %arg0 : f32
-    %2 = arith.mulf %c, %1 : f32
-    test.region_yield %2 : f32
-  } : tensor<?x?xf32>, tensor<?x?xf32> -> tensor<?x?xf32>
-  return %0, %1 : tensor<?x?xf32>, tensor<?x?xf32>
-}
-// CHECK-LABEL: func @cse_single_block_with_commutative_ops
-//       CHECK:   %[[OP:.+]] = test.cse_of_single_block_op
-//   CHECK-NOT:   test.cse_of_single_block_op
-//       CHECK:   return %[[OP]], %[[OP]]
-
 func.func @failing_issue_59135(%arg0: tensor<2x2xi1>, %arg1: f32, %arg2 : tensor<2xi1>) -> (tensor<2xi1>, tensor<2xi1>) {
-  %false_2 = arith.constant false 
-  %true_5 = arith.constant true 
+  %false_2 = arith.constant false
+  %true_5 = arith.constant true
   %9 = test.cse_of_single_block_op inputs(%arg2) {
   ^bb0(%out: i1):
     %true_144 = arith.constant true
     test.region_yield %true_144 : i1
   } : tensor<2xi1> -> tensor<2xi1>
   %15 = test.cse_of_single_block_op inputs(%arg2) {
   ^bb0(%out: i1):
     %true_144 = arith.constant true