[flang][cuda] Break circular global init when copying them to the GPU module (#211100)

PTX does not support cycle in globals. 

```
error: Circular dependency found in global variable set
```

To avoid cycle in type descriptor, check if there is a cycle and break
it.

The algorithm uses Tarjan’s strongly connected components algorithm for
finding strongly connected components in a directed graph.
- Each fir.global is a graph node.
- Each fir.address_of in a global initializer is an edge to another
global.
- An SCC containing multiple globals—or a self-edge—represents a
circular initializer dependency.

flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp

@@ -16,6 +16,8 @@
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/DenseSet.h"
+#include "llvm/ADT/SmallVector.h"
+#include <functional>
 
 namespace fir {
 #define GEN_PASS_DEF_CUFDEVICEGLOBAL
@@ -114,6 +116,81 @@ processPotentialTypeDescriptor(mlir::Type candidateType,
     processTypeDescriptor(recTy, symbolTable, candidates);
 }
 
+/// NVPTX cannot emit global initializers that form a reference cycle (see
+/// VisitGlobalVariableForEmission). Fortran type-info globals often do
+/// (mutually recursive derived types). Make a subset of the GPU copies extern
+/// declarations so that the remaining initializer dependency graph is
+/// acyclic, while preserving as many complete initializers as possible.
+static void dropCyclicGlobalInitializers(mlir::gpu::GPUModuleOp gpuMod) {
+  llvm::DenseMap<llvm::StringRef, fir::GlobalOp> byName;
+  llvm::SmallVector<fir::GlobalOp, 16> globals;
+  for (auto global : gpuMod.getOps<fir::GlobalOp>()) {
+    byName[global.getSymName()] = global;
+    globals.push_back(global);
+  }
+  if (globals.empty())
+    return;
+
+  llvm::sort(globals, [](fir::GlobalOp lhs, fir::GlobalOp rhs) {
+    return lhs.getSymName() < rhs.getSymName();
+  });
+
+  // Adjacency: global -> targets referenced via fir.address_of in its body.
+  llvm::DenseMap<fir::GlobalOp, llvm::SmallVector<fir::GlobalOp, 4>> adj;
+  for (fir::GlobalOp global : globals) {
+    global.walk([&](fir::AddrOfOp addrOf) {
+      fir::GlobalOp target =
+          byName.lookup(addrOf.getSymbol().getRootReference().getValue());
+      if (!target)
+        return;
+      adj[global].push_back(target);
+    });
+  }
+
+  // Greedily construct a feedback vertex set. Dropping one initializer
+  // removes all outgoing dependency edges from that global. Restart the DFS
+  // after each cut until no back edge remains.
+  llvm::DenseSet<fir::GlobalOp> declarationOnly;
+  while (true) {
+    // 0: unvisited, 1: active, 2: complete.
+    llvm::DenseMap<fir::GlobalOp, unsigned> state;
+    fir::GlobalOp cut;
+    std::function<bool(fir::GlobalOp)> findCycle =
+        [&](fir::GlobalOp global) -> bool {
+      if (declarationOnly.contains(global))
+        return false;
+      state[global] = 1;
+      for (fir::GlobalOp target : adj.lookup(global)) {
+        if (declarationOnly.contains(target))
+          continue;
+        if (state.lookup(target) == 1) {
+          cut = global;
+          return true;
+        }
+        if (state.lookup(target) == 0 && findCycle(target))
+          return true;
+      }
+      state[global] = 2;
+      return false;
+    };
+
+    for (fir::GlobalOp global : globals)
+      if (state.lookup(global) == 0 && findCycle(global))
+        break;
+    if (!cut)
+      break;
+    declarationOnly.insert(cut);
+  }
+
+  for (fir::GlobalOp global : declarationOnly) {
+    global.getRegion().getBlocks().clear();
+    global.removeInitValAttr();
+    // No initializer: use default external linkage so NVPTX emits
+    // `.extern .global` with no initializer dependency edges.
+    global.removeLinkNameAttr();
+  }
+}
+
 class CUFDeviceGlobal : public fir::impl::CUFDeviceGlobalBase<CUFDeviceGlobal> {
 public:
   using CUFDeviceGlobalBase::CUFDeviceGlobalBase;
@@ -190,6 +267,9 @@ public:
         clonedGlobal.removeLinkNameAttr();
       gpuSymTable.insert(cloned);
     }
+    // Type-info globals for mutually recursive derived types form initializer
+    // cycles; NVPTX rejects those. Drop initializers from cyclic GPU copies.
+    dropCyclicGlobalInitializers(gpuMod);
   }
 };
 } // namespace

flang/test/Fir/CUDA/cuda-device-global-cycle.fir

@@ -1,7 +1,8 @@
 // RUN: fir-opt --cuf-device-global %s | FileCheck %s
 
 // Mutually referencing type-descriptor globals must not cause infinite
-// recursion in processAddrOfOp when recurseInGlobal is true.
+// recursion. One GPU copy must drop its initializer to break the cycle, while
+// the other keeps its complete initializer.
 module attributes {gpu.container_module} {
   fir.global linkonce @_QMmE.dt.t1 constant : i64 {
     %0 = fir.address_of(@_QMmE.dt.t2) : !fir.ref<i64>
@@ -20,6 +21,15 @@ module attributes {gpu.container_module} {
   }
 }
 
-// CHECK-LABEL: gpu.module @cuda_device_mod
-// CHECK-DAG: fir.global linkonce @_QMmE.dt.t1
-// CHECK-DAG: fir.global linkonce @_QMmE.dt.t2
+// Host copies keep the cycle.
+// CHECK: fir.global linkonce @_QMmE.dt.t1
+// CHECK: fir.address_of(@_QMmE.dt.t2)
+// CHECK: fir.global linkonce @_QMmE.dt.t2
+// CHECK: fir.address_of(@_QMmE.dt.t1)
+
+// Cut the lexicographically last global in the cycle; the other keeps its init.
+// CHECK-LABEL: gpu.module @cuda_device_mod {
+// CHECK-DAG: fir.global @_QMmE.dt.t2 constant : i64{{$}}
+// CHECK-DAG: fir.global linkonce @_QMmE.dt.t1 constant : i64 {
+// CHECK-DAG: fir.address_of(@_QMmE.dt.t2)
+// CHECK-DAG: fir.has_value