[flang-commits] [flang] 4d33826 - [flang] Fixed issues in nested LICM. (#192117)

[via flang-commits]

Author: Slava Zakharin
Date: 2026-04-14T13:54:19-07:00
New Revision: 4d33826d60acf6cc95a52b380954556a84c9cbb7

URL: https://github.com/llvm/llvm-project/commit/4d33826d60acf6cc95a52b380954556a84c9cbb7
DIFF: https://github.com/llvm/llvm-project/commit/4d33826d60acf6cc95a52b380954556a84c9cbb7.diff

LOG: [flang] Fixed issues in nested LICM. (#192117)

First change is to check the hoisting safety for all nested
operations of the candidate. This prevents hoistings of
region operations as in the added test.

When hoisting operations from nested regions we have to
check every parent region for `canMoveOutOf`, otherwise,
illegal hoisting may happen. This second change is NFC,
because all operations that support `OperationMoveOpInterface`
currently also support `LoopLikeOpInterface` and their regions
are not considered for nested hoisting. Anyway, it is worth
fixing it.

Added: 
    

Modified: 
    flang/lib/Optimizer/Transforms/LoopInvariantCodeMotion.cpp
    flang/test/Transforms/CUF/cuf-kernel-licm.fir

Removed: 
    


################################################################################
diff  --git a/flang/lib/Optimizer/Transforms/LoopInvariantCodeMotion.cpp b/flang/lib/Optimizer/Transforms/LoopInvariantCodeMotion.cpp
index d1c4046f38b19..b3b9fa2e01ee9 100644
--- a/flang/lib/Optimizer/Transforms/LoopInvariantCodeMotion.cpp
+++ b/flang/lib/Optimizer/Transforms/LoopInvariantCodeMotion.cpp
@@ -335,8 +335,23 @@ void LoopInvariantCodeMotion::runOnOperation() {
     auto isDefinedOutsideRegion = [&](Value value, Region *) {
       return loopLike.isDefinedOutsideOfLoop(value);
     };
+    // Check canMoveOutOf for the candidate and all its nested operations.
+    // Moving an operation with regions also moves its contents, so
+    // restrictions like cuf.kernel blocking !fir.ref operands must be
+    // checked transitively.
+    auto canMoveOutOfOp = [&](Operation *regionOwner, Operation *candidate) {
+      if (!fir::canMoveOutOf(regionOwner, candidate))
+        return false;
+      bool blocked = false;
+      candidate->walk([&](Operation *nested) {
+        if (nested != candidate && !fir::canMoveOutOf(regionOwner, nested))
+          blocked = true;
+        return blocked ? WalkResult::interrupt() : WalkResult::advance();
+      });
+      return !blocked;
+    };
     auto canMoveOutOfLoop = [&](Operation *op) {
-      if (!fir::canMoveOutOf(loopLike, op)) {
+      if (!canMoveOutOfOp(loopLike, op)) {
         LDBG() << "Cannot hoist " << *op << " out of the loop";
         return false;
       }
@@ -383,6 +398,26 @@ void LoopInvariantCodeMotion::runOnOperation() {
       return;
 
     auto shouldMoveFromNestedRegion = [&](Operation *op, Region *) {
+      // Check that all intermediate operations between op and the loop
+      // allow the candidate to be moved out.  For example, cuf.kernel
+      // restricts hoisting of operations with !fir.ref operands.
+      for (Operation *ancestor = op->getParentOp();
+           ancestor != loopLike.getOperation();
+           ancestor = ancestor->getParentOp()) {
+        if (!ancestor) {
+          // The operation is no longer nested inside the loop (a parent
+          // operation was already hoisted out). Nothing to do.
+          return false;
+        }
+        if (!canMoveOutOfOp(ancestor, op)) {
+          LDBG() << "Cannot hoist " << *op
+                 << " out of intermediate operation: ";
+          LDBG_OS([&](llvm::raw_ostream &os) {
+            ancestor->print(os, OpPrintingFlags().skipRegions());
+          });
+          return false;
+        }
+      }
       return canMoveOutOfLoop(op) &&
              shouldMoveOutOfLoop(op, loopLike,
                                  /*maybeConditionallyExecuted=*/true);

diff  --git a/flang/test/Transforms/CUF/cuf-kernel-licm.fir b/flang/test/Transforms/CUF/cuf-kernel-licm.fir
index 88e8dc5ee8ba3..7075d19903e31 100644
--- a/flang/test/Transforms/CUF/cuf-kernel-licm.fir
+++ b/flang/test/Transforms/CUF/cuf-kernel-licm.fir
@@ -83,6 +83,40 @@ func.func @_QPtest(%arg0: !fir.ref<!fir.array<10xf32>> {cuf.data_attr = #cuf.cud
 
 // -----
 
+// Verify that scf.if containing fir.convert with !fir.ref operand is NOT
+// hoisted out of cuf.kernel as a whole.  The canMoveOutOf restriction must
+// be checked transitively for all nested operations of the candidate,
+// not just the candidate itself.  fir.convert can convert !fir.ref to i64,
+// so scf.if may yield an integer result produced from a !fir.ref.
+// CHECK-LABEL: func.func @transitive_canMoveOutOf(
+// CHECK-SAME:    %[[ARG0:.*]]: !fir.ref<!fir.array<10xf32>>
+// CHECK-SAME:    %[[COND:.*]]: i1)
+// CHECK:         cuf.kernel
+// CHECK:           %[[ADDR:.*]] = scf.if %[[COND]] -> (i64)
+// CHECK:             fir.convert %[[ARG0]]
+// CHECK:           %[[REF:.*]] = fir.convert %[[ADDR]] : (i64) -> !fir.ref<f32>
+// CHECK:           fir.store %{{.*}} to %[[REF]]
+func.func @transitive_canMoveOutOf(%arg0: !fir.ref<!fir.array<10xf32>> {cuf.data_attr = #cuf.cuda<device>}, %cond: i1) {
+  %c0_i64 = arith.constant 0 : i64
+  %c1 = arith.constant 1 : index
+  %c10 = arith.constant 10 : index
+  %cst = arith.constant 1.000000e+00 : f32
+  cuf.kernel<<<*, *>>> (%arg2 : index) = (%c1 : index) to (%c10 : index)  step (%c1 : index) {
+    %addr = scf.if %cond -> (i64) {
+      %ptr = fir.convert %arg0 : (!fir.ref<!fir.array<10xf32>>) -> i64
+      scf.yield %ptr : i64
+    } else {
+      scf.yield %c0_i64 : i64
+    }
+    %ref = fir.convert %addr : (i64) -> !fir.ref<f32>
+    fir.store %cst to %ref : !fir.ref<f32>
+    "fir.end"() : () -> ()
+  }
+  return
+}
+
+// -----
+
 // Verify that Pure fir.address_of operation is not hoisted by LICM
 // out of cuf.kernel.
 // CHECK-LABEL: func.func @_QQmain() {