[flang-commits] [flang] [flang][cuda] Handle special default.nonTbpDefinedIoTable in pass (PR #120686)

[Valentin Clement バレンタイン クレメン via flang-commits]

https://github.com/clementval created https://github.com/llvm/llvm-project/pull/120686

`default.nonTbpDefinedIoTable` is a special global defined for IO that doesn't follow the mangling scheme and is then not handle correctly in the `CompilerGeneratedNames` pass. 

Add this exception in the pass so the pass can replace symbol in it. 

Note: The prefix might be changed so it follow the mangling scheme and no special handling would be needed. 

>From a5e102f08b59f69829c810f10a1c8f91ef8a1fe3 Mon Sep 17 00:00:00 2001
From: Valentin Clement <clementval at gmail.com>
Date: Thu, 19 Dec 2024 22:29:00 -0800
Subject: [PATCH] [flang][cuda] Handle specifal default.nonTbpDefinedIoTable in
 pass

---
 flang/lib/Optimizer/CodeGen/CodeGen.cpp        |  4 +++-
 .../Transforms/CompilerGeneratedNames.cpp      |  5 +++--
 .../CUDA/cuda-compiler-generated-names.mlir    | 18 ++++++++++++++++++
 3 files changed, 24 insertions(+), 3 deletions(-)

diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index aaf97d46d83d4f..786425befb9496 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -2990,10 +2990,12 @@ struct GlobalOpConversion : public fir::FIROpConversion<fir::GlobalOp> {
       g.setAlignment(*global.getAlignment());
 
     auto module = global->getParentOfType<mlir::ModuleOp>();
+    auto gpuMod = global->getParentOfType<mlir::gpu::GPUModuleOp>();
     // Add comdat if necessary
     if (fir::getTargetTriple(module).supportsCOMDAT() &&
         (linkage == mlir::LLVM::Linkage::Linkonce ||
-         linkage == mlir::LLVM::Linkage::LinkonceODR)) {
+         linkage == mlir::LLVM::Linkage::LinkonceODR) &&
+        !gpuMod) {
       addComdat(g, rewriter, module);
     }
 
diff --git a/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp b/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
index f92c60908b1496..3ecf119e522678 100644
--- a/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
+++ b/flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp
@@ -48,8 +48,9 @@ void CompilerGeneratedNamesConversionPass::runOnOperation() {
     auto symName = op.getAttrOfType<mlir::StringAttr>(
         mlir::SymbolTable::getSymbolAttrName());
     auto deconstructedName = fir::NameUniquer::deconstruct(symName);
-    if (deconstructedName.first != fir::NameUniquer::NameKind::NOT_UNIQUED &&
-        !fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) {
+    if ((deconstructedName.first != fir::NameUniquer::NameKind::NOT_UNIQUED &&
+         !fir::NameUniquer::isExternalFacingUniquedName(deconstructedName)) ||
+        symName.getValue().starts_with("default.nonTbpDefinedIoTable")) {
       std::string newName =
           fir::NameUniquer::replaceSpecialSymbols(symName.getValue().str());
       if (newName != symName) {
diff --git a/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir b/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
index 4507e444d1b510..1a6c67227d9fe7 100644
--- a/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
+++ b/flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir
@@ -8,6 +8,22 @@ module @mod1 attributes {gpu.container} {
       %0 = fir.embox %arg0() : (!fir.ref<!fir.type<_QMtest_dinitTtseq{i:i32}>>) -> !fir.box<!fir.type<_QMtest_dinitTtseq{i:i32}>>
       return
     }
+
+    fir.global @default.nonTbpDefinedIoTable constant : tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1> {
+      %true = arith.constant true
+      %c0_i64 = arith.constant 0 : i64
+      %0 = fir.undefined tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>
+      %1 = fir.insert_value %0, %c0_i64, [0 : index] : (tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>, i64) -> tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>
+      %2 = fir.zero_bits !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>
+      %3 = fir.insert_value %1, %2, [1 : index] : (tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>) -> tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>
+      %4 = fir.insert_value %3, %true, [2 : index] : (tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>, i1) -> tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>
+      fir.has_value %4 : tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>
+    }
+
+    func.func @special() {
+      %0 = fir.address_of(@default.nonTbpDefinedIoTable) : !fir.ref<tuple<i64, !fir.ref<!fir.array<0xtuple<!fir.ref<none>, !fir.ref<none>, i32, i1>>>, i1>>
+      return
+    }
   }
 }
 
@@ -15,3 +31,5 @@ module @mod1 attributes {gpu.container} {
 // CHECK: llvm.mlir.global linkonce constant @_QMtest_dinitEXdtXtseq
 // CHECK: llvm.mlir.addressof @_QMtest_dinitEXdtXtseq : !llvm.ptr
 
+// CHECK: llvm.mlir.global external constant @defaultXnonTbpDefinedIoTable()
+// CHECK: llvm.mlir.addressof @defaultXnonTbpDefinedIoTable