[flang][cuda] Make default.nonTbpDefinedIoTable compiler generated

[clementval]

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. Update how it is generated with doGenerated so it can be handle without special handling.

Also do not generate comdat in gpu module as the current code is not handling nested module correctly.

[llvmbot]

@llvm/pr-subscribers-flang-codegen

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

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

Changes

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.

---

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

3 Files Affected:

• (modified) flang/lib/Optimizer/CodeGen/CodeGen.cpp (+3-1)
• (modified) flang/lib/Optimizer/Transforms/CompilerGeneratedNames.cpp (+3-2)
• (modified) flang/test/Fir/CUDA/cuda-compiler-generated-names.mlir (+18)

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

[jeanPerier]

Maybe "default.nonTbpDefinedIoTable" name could also be generated as a compiler generated name using fir::NameUniquer::doGenerated here so that it later falls into the "Namekind::GENERATED" which is more accurate.
[edit: that way I believe not special handling would be needed here]

[clementval]

Maybe "default.nonTbpDefinedIoTable" name could also be generated as a compiler generated name using fir::NameUniquer::doGenerated here so that it later falls into the "Namekind::GENERATED" which is more accurate. [edit: that way I believe not special handling would be needed here]

Yeah that was my other proposition. Just updated the patch so it's cleaner.

[razvanlupusoru]

LGTM!

[vdonaldson]

Thanks Valentin