[flang-commits] [flang] [flang][CodeGen] Fix address space mismatch for CUF globals in AddrOfOpConversion (PR #190408)
via flang-commits
flang-commits at lists.llvm.org
Tue Apr 14 15:53:48 PDT 2026
https://github.com/khaki3 updated https://github.com/llvm/llvm-project/pull/190408
>From 175bbd538cc5df791c38e00961d5893a4cfb5a7a Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Fri, 3 Apr 2026 12:38:08 -0700
Subject: [PATCH 1/5] [flang][CodeGen] Fix address space mismatch for CUF
globals in AddrOfOpConversion
When lowering fir.address_of to llvm.mlir.addressof, the AddrOfOpConversion
only looked up llvm.mlir.global to determine the address space. During
conversion, the global may still be a fir.global (not yet converted), causing
the lookup to fail and fall back to address space 0. For CUF constant memory
globals (address space 4), this creates a mismatch between the addressof
pointer and the global's address space, triggering a verification error.
Fix by also looking up fir::GlobalOp when LLVM::GlobalOp is not found and
deriving the address space from its CUF data_attr.
Made-with: Cursor
---
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 55 ++++++++++++++++---------
1 file changed, 35 insertions(+), 20 deletions(-)
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index fc4dc85ff8748..77b31bbbb12c9 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -196,26 +196,37 @@ mlir::Value replaceWithAddrOfOrASCast(mlir::ConversionPatternRewriter &rewriter,
return mlir::LLVM::AddressOfOp::create(rewriter, loc, type, symName);
}
-static std::uint64_t getAddressSpace(fir::AddrOfOp addr,
- mlir::ConversionPatternRewriter &rewriter,
- std::uint64_t defaultAS) {
- auto global = addr->getParentOfType<mlir::ModuleOp>()
- .lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
- if (global)
- return global.getAddrSpace();
- auto firGlobal =
- addr->getParentOfType<mlir::ModuleOp>().lookupSymbol<fir::GlobalOp>(
- addr.getSymbol());
- if (firGlobal && firGlobal.getDataAttr() &&
- *firGlobal.getDataAttr() == cuf::DataAttribute::Constant)
- return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Constant);
- return defaultAS;
+/// Return the NVVM address space implied by a CUF data attribute on a
+/// fir::GlobalOp that has not yet been converted to llvm.mlir.global.
+/// Returns std::nullopt if no CUF-specific address space applies.
+static std::optional<unsigned> getCUFAddrSpace(fir::GlobalOp global) {
+ if (auto dataAttr = global.getDataAttr()) {
+ if (*dataAttr == cuf::DataAttribute::Constant)
+ return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Constant);
+ if (*dataAttr == cuf::DataAttribute::Shared)
+ return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared);
+ if (*dataAttr == cuf::DataAttribute::Managed)
+ return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Global);
+ }
+ return std::nullopt;
}
/// Lower `fir.address_of` operation to `llvm.address_of` operation.
struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
using FIROpConversion::FIROpConversion;
+ /// Look up the address space for a symbol in \p mod, handling both
+ /// already-converted llvm.mlir.global and not-yet-converted fir.global.
+ template <typename ModOp>
+ unsigned getAddrSpaceForGlobal(ModOp mod, mlir::SymbolRefAttr sym,
+ unsigned fallback) const {
+ if (auto g = mod.template lookupSymbol<mlir::LLVM::GlobalOp>(sym))
+ return g.getAddrSpace();
+ if (auto g = mod.template lookupSymbol<fir::GlobalOp>(sym))
+ return getCUFAddrSpace(g);
+ return fallback;
+ }
+
llvm::LogicalResult
matchAndRewrite(fir::AddrOfOp addr, OpAdaptor adaptor,
mlir::ConversionPatternRewriter &rewriter) const override {
@@ -224,7 +235,9 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
auto global = gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
replaceWithAddrOfOrASCast(
rewriter, addr->getLoc(),
- global ? global.getAddrSpace() : getGlobalAddressSpace(rewriter),
+ global ? global.getAddrSpace()
+ : getAddrSpaceForGlobal(gpuMod, addr.getSymbol(),
+ getGlobalAddressSpace(rewriter)),
getProgramAddressSpace(rewriter),
global ? global.getSymName()
: addr.getSymbol().getRootReference().getValue(),
@@ -232,12 +245,14 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
return mlir::success();
}
- std::uint64_t globalAS =
- getAddressSpace(addr, rewriter, getGlobalAddressSpace(rewriter));
- auto global = addr->getParentOfType<mlir::ModuleOp>()
- .lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
+ auto mod = addr->getParentOfType<mlir::ModuleOp>();
+ auto global = mod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
replaceWithAddrOfOrASCast(
- rewriter, addr->getLoc(), globalAS, getProgramAddressSpace(rewriter),
+ rewriter, addr->getLoc(),
+ global ? global.getAddrSpace()
+ : getAddrSpaceForGlobal(mod, addr.getSymbol(),
+ getGlobalAddressSpace(rewriter)),
+ getProgramAddressSpace(rewriter),
global ? global.getSymName()
: addr.getSymbol().getRootReference().getValue(),
convertType(addr.getType()), addr);
>From e5f0cbc823c63a0576289e6168f913084bc1d5eb Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Fri, 3 Apr 2026 13:02:32 -0700
Subject: [PATCH 2/5] [flang][CodeGen] Add test for host-side addressof with
CUF constant global
Add a test case verifying that a host-side fir.address_of referencing a
fir.global with cuf.cuda<constant> data_attr correctly produces an
llvm.mlir.addressof with addr_space 4 and an addrspacecast to ptr.
Made-with: Cursor
---
flang/test/Fir/CUDA/cuda-code-gen.mlir | 21 +++++++++++++++++++++
1 file changed, 21 insertions(+)
diff --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir
index 68fef43bef3bb..7ad13729afbac 100644
--- a/flang/test/Fir/CUDA/cuda-code-gen.mlir
+++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir
@@ -360,3 +360,24 @@ func.func private @_FortranACUFGetDeviceAddress(!fir.llvm_ptr<i8>, !fir.ref<i8>,
// CHECK-LABEL: llvm.func @sub16_
// CHECK: llvm.mlir.addressof @_QMdevice_dataEd16 : !llvm.ptr<4>
+
+// -----
+
+// Test that a host-side fir.address_of referencing a fir.global with CUF
+// constant data_attr produces an addrspacecast from ptr<4> to ptr.
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
+ fir.global @_QMmodEcval {data_attr = #cuf.cuda<constant>} : i32 {
+ %0 = fir.zero_bits i32
+ fir.has_value %0 : i32
+ }
+ func.func @_QQhost() {
+ %0 = fir.address_of(@_QMmodEcval) : !fir.ref<i32>
+ return
+ }
+}
+
+// CHECK: llvm.mlir.global external @_QMmodEcval() {addr_space = 4 : i32} : i32
+// CHECK-LABEL: llvm.func @_QQhost()
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @_QMmodEcval : !llvm.ptr<4>
+// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDR]] : !llvm.ptr<4> to !llvm.ptr
>From f4cfe50746d578f0907aba56bda581234ae240f2 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Mon, 6 Apr 2026 03:55:40 -0700
Subject: [PATCH 3/5] [flang][CodeGen] Fix getCUFAddrSpace to not override
target global address space
getCUFAddrSpace was returning 0 for globals without CUF data attributes,
which incorrectly overrode the target's global address space (e.g. 1 on
AMDGPU). Return std::nullopt instead so the fallback address space from
the target data layout is used for non-CUF globals.
Made-with: Cursor
---
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 77b31bbbb12c9..3a7aa707b3eb0 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -223,7 +223,8 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
if (auto g = mod.template lookupSymbol<mlir::LLVM::GlobalOp>(sym))
return g.getAddrSpace();
if (auto g = mod.template lookupSymbol<fir::GlobalOp>(sym))
- return getCUFAddrSpace(g);
+ if (auto as = getCUFAddrSpace(g))
+ return *as;
return fallback;
}
>From 990493abfd1c5fdaf555c391beb8e953681d19ac Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Tue, 14 Apr 2026 15:22:58 -0700
Subject: [PATCH 4/5] [flang][CodeGen] Add test for host-side addressof with
CUF shared global
Add a test verifying that a host-side fir.address_of referencing a
fir.global with cuf.cuda<shared> data_attr produces an addressof with
addr_space 3 and an addrspacecast. The upstream test already covers
the constant case; this tests the shared case added by this patch.
Made-with: Cursor
---
flang/test/Fir/CUDA/cuda-code-gen.mlir | 21 +++++++++++++++++++++
1 file changed, 21 insertions(+)
diff --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir
index 7ad13729afbac..7ccf16c2567b2 100644
--- a/flang/test/Fir/CUDA/cuda-code-gen.mlir
+++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir
@@ -381,3 +381,24 @@ module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> :
// CHECK-LABEL: llvm.func @_QQhost()
// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @_QMmodEcval : !llvm.ptr<4>
// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDR]] : !llvm.ptr<4> to !llvm.ptr
+
+// -----
+
+// Test that a host-side fir.address_of referencing a fir.global with CUF
+// shared data_attr produces an addrspacecast from ptr<3> to ptr.
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
+ fir.global @_QMmodEsval {data_attr = #cuf.cuda<shared>} : i32 {
+ %0 = fir.zero_bits i32
+ fir.has_value %0 : i32
+ }
+ func.func @_QQhost_shared() {
+ %0 = fir.address_of(@_QMmodEsval) : !fir.ref<i32>
+ return
+ }
+}
+
+// CHECK: llvm.mlir.global external @_QMmodEsval() {addr_space = 3 : i32} : i32
+// CHECK-LABEL: llvm.func @_QQhost_shared()
+// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @_QMmodEsval : !llvm.ptr<3>
+// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDR]] : !llvm.ptr<3> to !llvm.ptr
>From 07f091d7c2277a6cadef76dae48f8f917bedfad8 Mon Sep 17 00:00:00 2001
From: Kazuaki Matsumura <kmatsumura at nvidia.com>
Date: Tue, 14 Apr 2026 15:26:15 -0700
Subject: [PATCH 5/5] [flang][CodeGen] Update test to use shared instead of
constant, remove managed from getCUFAddrSpace
The upstream sub16_ test already covers the constant case. Replace the
duplicate constant test with a shared (addr_space 3) test.
Remove Managed from getCUFAddrSpace since managed globals only get
addr_space 1 inside gpu.module (guarded by gpuMod in GlobalOpConversion),
not on the host side.
Made-with: Cursor
---
flang/lib/Optimizer/CodeGen/CodeGen.cpp | 23 +++++----------
flang/test/Fir/CUDA/cuda-code-gen.mlir | 39 +++++--------------------
2 files changed, 15 insertions(+), 47 deletions(-)
diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
index 3a7aa707b3eb0..acf1ba4e52a14 100644
--- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp
+++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp
@@ -205,8 +205,6 @@ static std::optional<unsigned> getCUFAddrSpace(fir::GlobalOp global) {
return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Constant);
if (*dataAttr == cuf::DataAttribute::Shared)
return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Shared);
- if (*dataAttr == cuf::DataAttribute::Managed)
- return static_cast<unsigned>(mlir::NVVM::NVVMMemorySpace::Global);
}
return std::nullopt;
}
@@ -232,30 +230,23 @@ struct AddrOfOpConversion : public fir::FIROpConversion<fir::AddrOfOp> {
matchAndRewrite(fir::AddrOfOp addr, OpAdaptor adaptor,
mlir::ConversionPatternRewriter &rewriter) const override {
+ auto symName = addr.getSymbol().getRootReference().getValue();
if (auto gpuMod = addr->getParentOfType<mlir::gpu::GPUModuleOp>()) {
- auto global = gpuMod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
replaceWithAddrOfOrASCast(
rewriter, addr->getLoc(),
- global ? global.getAddrSpace()
- : getAddrSpaceForGlobal(gpuMod, addr.getSymbol(),
- getGlobalAddressSpace(rewriter)),
- getProgramAddressSpace(rewriter),
- global ? global.getSymName()
- : addr.getSymbol().getRootReference().getValue(),
+ getAddrSpaceForGlobal(gpuMod, addr.getSymbol(),
+ getGlobalAddressSpace(rewriter)),
+ getProgramAddressSpace(rewriter), symName,
convertType(addr.getType()), addr);
return mlir::success();
}
auto mod = addr->getParentOfType<mlir::ModuleOp>();
- auto global = mod.lookupSymbol<mlir::LLVM::GlobalOp>(addr.getSymbol());
replaceWithAddrOfOrASCast(
rewriter, addr->getLoc(),
- global ? global.getAddrSpace()
- : getAddrSpaceForGlobal(mod, addr.getSymbol(),
- getGlobalAddressSpace(rewriter)),
- getProgramAddressSpace(rewriter),
- global ? global.getSymName()
- : addr.getSymbol().getRootReference().getValue(),
+ getAddrSpaceForGlobal(mod, addr.getSymbol(),
+ getGlobalAddressSpace(rewriter)),
+ getProgramAddressSpace(rewriter), symName,
convertType(addr.getType()), addr);
return mlir::success();
}
diff --git a/flang/test/Fir/CUDA/cuda-code-gen.mlir b/flang/test/Fir/CUDA/cuda-code-gen.mlir
index 7ccf16c2567b2..0fdcf5442b06b 100644
--- a/flang/test/Fir/CUDA/cuda-code-gen.mlir
+++ b/flang/test/Fir/CUDA/cuda-code-gen.mlir
@@ -349,7 +349,7 @@ func.func @sub16_(%arg0: !fir.ref<f32> {fir.bindc_name = "h16"}) attributes {fir
return
}
fir.global linkonce @_QQclXc8657e47c19bb9e89730387c9d99c2da constant : !fir.char<1,38> {
- %0 = fir.string_lit "/local/home/vclement/lorado/dummy.cuf\00"(38) : !fir.char<1,38>
+ %0 = fir.string_lit "/path/to/source/test/dummy_module.cuf\00"(38) : !fir.char<1,38>
fir.has_value %0 : !fir.char<1,38>
}
fir.global @_QMdevice_dataEd16 {data_attr = #cuf.cuda<constant>} : f32 {
@@ -363,39 +363,16 @@ func.func private @_FortranACUFGetDeviceAddress(!fir.llvm_ptr<i8>, !fir.ref<i8>,
// -----
-// Test that a host-side fir.address_of referencing a fir.global with CUF
-// constant data_attr produces an addrspacecast from ptr<4> to ptr.
-
-module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
- fir.global @_QMmodEcval {data_attr = #cuf.cuda<constant>} : i32 {
- %0 = fir.zero_bits i32
- fir.has_value %0 : i32
- }
- func.func @_QQhost() {
- %0 = fir.address_of(@_QMmodEcval) : !fir.ref<i32>
- return
- }
-}
-
-// CHECK: llvm.mlir.global external @_QMmodEcval() {addr_space = 4 : i32} : i32
-// CHECK-LABEL: llvm.func @_QQhost()
-// CHECK: %[[ADDR:.*]] = llvm.mlir.addressof @_QMmodEcval : !llvm.ptr<4>
-// CHECK: %{{.*}} = llvm.addrspacecast %[[ADDR]] : !llvm.ptr<4> to !llvm.ptr
-
-// -----
-
// Test that a host-side fir.address_of referencing a fir.global with CUF
// shared data_attr produces an addrspacecast from ptr<3> to ptr.
-module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>} {
- fir.global @_QMmodEsval {data_attr = #cuf.cuda<shared>} : i32 {
- %0 = fir.zero_bits i32
- fir.has_value %0 : i32
- }
- func.func @_QQhost_shared() {
- %0 = fir.address_of(@_QMmodEsval) : !fir.ref<i32>
- return
- }
+fir.global @_QMmodEsval {data_attr = #cuf.cuda<shared>} : i32 {
+ %0 = fir.zero_bits i32
+ fir.has_value %0 : i32
+}
+func.func @_QQhost_shared() {
+ %0 = fir.address_of(@_QMmodEsval) : !fir.ref<i32>
+ return
}
// CHECK: llvm.mlir.global external @_QMmodEsval() {addr_space = 3 : i32} : i32
More information about the flang-commits
mailing list