[flang-commits] [flang] [flang][cuda] Keep cycle-broken device globals defined (PR #211897)

[via flang-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-flang-fir-hlfir

Author: Valentin Clement (バレンタイン クレメン) (clementval)

<details>
<summary>Changes</summary>

Breaking cyclic global initializers by removing an initializer produced an
external declaration, which could leave references unresolved during device
linking.

Replace selected cyclic initializers with zero initializers instead. This
breaks the NVPTX dependency cycle while preserving device symbol definitions.

---
Full diff: https://github.com/llvm/llvm-project/pull/211897.diff


2 Files Affected:

- (modified) flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp (+21-14) 
- (modified) flang/test/Fir/CUDA/cuda-device-global-cycle.fir (+8-6) 


``````````diff
diff --git a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp
index ba0a6aec9992a..5c5d76588673e 100644
--- a/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp
+++ b/flang/lib/Optimizer/Transforms/CUDA/CUFDeviceGlobal.cpp
@@ -12,6 +12,7 @@
 #include "flang/Optimizer/Support/InternalNames.h"
 #include "flang/Optimizer/Transforms/Passes.h"
 #include "mlir/Dialect/LLVMIR/NVVMDialect.h"
+#include "mlir/IR/Builders.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
@@ -118,9 +119,10 @@ processPotentialTypeDescriptor(mlir::Type candidateType,
 
 /// 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.
+/// (mutually recursive derived types). Replace the initializer of a feedback
+/// vertex set of GPU copies with a zero value so the remaining dependency
+/// graph is acyclic, while keeping those symbols as real definitions (not
+/// `.extern`) so nvlink can resolve references from other device objects.
 static void dropCyclicGlobalInitializers(mlir::gpu::GPUModuleOp gpuMod) {
   llvm::DenseMap<llvm::StringRef, fir::GlobalOp> byName;
   llvm::SmallVector<fir::GlobalOp, 16> globals;
@@ -147,21 +149,21 @@ static void dropCyclicGlobalInitializers(mlir::gpu::GPUModuleOp gpuMod) {
     });
   }
 
-  // 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;
+  // Greedily construct a feedback vertex set. Replacing one initializer with
+  // a zero value removes all outgoing dependency edges from that global.
+  // Restart the DFS after each cut until no back edge remains.
+  llvm::DenseSet<fir::GlobalOp> zeroInitOnly;
   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))
+      if (zeroInitOnly.contains(global))
         return false;
       state[global] = 1;
       for (fir::GlobalOp target : adj.lookup(global)) {
-        if (declarationOnly.contains(target))
+        if (zeroInitOnly.contains(target))
           continue;
         if (state.lookup(target) == 1) {
           cut = global;
@@ -179,15 +181,20 @@ static void dropCyclicGlobalInitializers(mlir::gpu::GPUModuleOp gpuMod) {
         break;
     if (!cut)
       break;
-    declarationOnly.insert(cut);
+    zeroInitOnly.insert(cut);
   }
 
-  for (fir::GlobalOp global : declarationOnly) {
+  for (fir::GlobalOp global : zeroInitOnly) {
     global.getRegion().getBlocks().clear();
     global.removeInitValAttr();
-    // No initializer: use default external linkage so NVPTX emits
-    // `.extern .global` with no initializer dependency edges.
-    global.removeLinkNameAttr();
+    // Keep linkage (typically linkonce). Emit a zero initializer so the
+    // symbol remains a definition with no fir.address_of dependency edges.
+    mlir::OpBuilder builder(global.getContext());
+    mlir::Block *block = builder.createBlock(&global.getRegion());
+    builder.setInsertionPointToStart(block);
+    mlir::Value zero =
+        fir::ZeroOp::create(builder, global.getLoc(), global.getType());
+    fir::HasValueOp::create(builder, global.getLoc(), zero);
   }
 }
 
diff --git a/flang/test/Fir/CUDA/cuda-device-global-cycle.fir b/flang/test/Fir/CUDA/cuda-device-global-cycle.fir
index 227ea252b1322..f929316545b8b 100644
--- a/flang/test/Fir/CUDA/cuda-device-global-cycle.fir
+++ b/flang/test/Fir/CUDA/cuda-device-global-cycle.fir
@@ -1,8 +1,8 @@
 // RUN: fir-opt --cuf-device-global %s | FileCheck %s
 
 // Mutually referencing type-descriptor globals must not cause infinite
-// recursion. One GPU copy must drop its initializer to break the cycle, while
-// the other keeps its complete initializer.
+// recursion. One GPU copy must get a zero initializer to break the cycle,
+// while the other keeps its complete initializer. Both remain definitions.
 module attributes {gpu.container_module} {
   fir.global linkonce @_QMmE.dt.t1 constant : i64 {
     %0 = fir.address_of(@_QMmE.dt.t2) : !fir.ref<i64>
@@ -27,9 +27,11 @@ module attributes {gpu.container_module} {
 // 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 {
+// Cut the source of the back edge; replace its init with zero_bits.
+// CHECK-LABEL: gpu.module @cuda_device_mod
+// CHECK-DAG: fir.global linkonce @_QMmE.dt.t2 constant : i64
+// CHECK-DAG: fir.zero_bits
+// CHECK-DAG: fir.has_value
+// CHECK-DAG: fir.global linkonce @_QMmE.dt.t1 constant : i64
 // CHECK-DAG: fir.address_of(@_QMmE.dt.t2)
 // CHECK-DAG: fir.has_value

``````````

</details>


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