[flang-commits] [flang] [flang] only instantiate referenced use associated variables (PR #222318)
via flang-commits
flang-commits at lists.llvm.org
Wed Sep 9 07:04:05 PDT 2026
https://github.com/jeanPerier updated https://github.com/llvm/llvm-project/pull/222318
>From 519decd8d17562ddbbe15325dfb724a06d30c544 Mon Sep 17 00:00:00 2001
From: Jean Perier <jperier at nvidia.com>
Date: Tue, 8 Sep 2026 07:05:33 -0700
Subject: [PATCH 1/2] [flang] only instantiate referenced use associated
variables
Procedure lowering instantiated every variable brought in by USE, so an empty 'use m' generated IR for all module variables.
This created a lot of extra IR (including globals that are not pruned until LLVM) for simple include like "use openmp"/"use openacc"... which made it harder to read the IR for debugging purposes.
getScopeVariableList(FunctionLikeUnit) now collects the ultimate symbols referenced in the unit (including internal procedures) via visitAllSymbols, and SymbolDependenceAnalysis skips use-associated symbols that are not referenced.
Dependencies found through specification/initialization expressions are still analyzed recursively, so module variables used for extents, character lengths or initial targets are instantiated as needed. Equivalence sets, namelists and procedure pointers hidden behind generics are preserved.
Tests relying on unreferenced USE instantiation were updated and a regression test added.
Assited-by: AI
---
flang/include/flang/Lower/PFTBuilder.h | 6 +
flang/lib/Lower/Bridge.cpp | 5 +-
flang/lib/Lower/PFTBuilder.cpp | 50 +++++-
.../Integration/MIF/coarray_allocation5.f90 | 1 +
flang/test/Lower/CUDA/cuda-cooperative.cuf | 5 +-
flang/test/Lower/CUDA/cuda-device-proc.cuf | 10 +-
flang/test/Lower/CUDA/cuda-devptr.cuf | 2 +-
flang/test/Lower/CUDA/cuda-module-use.cuf | 2 +-
.../OpenMP/threadprivate-use-association.f90 | 2 -
flang/test/Lower/declare-with-storage.f90 | 2 +
.../use_associated_variable_instantiation.f90 | 158 ++++++++++++++++++
11 files changed, 229 insertions(+), 14 deletions(-)
create mode 100644 flang/test/Lower/use_associated_variable_instantiation.f90
diff --git a/flang/include/flang/Lower/PFTBuilder.h b/flang/include/flang/Lower/PFTBuilder.h
index f495046c66a32..d3cf3b08ec625 100644
--- a/flang/include/flang/Lower/PFTBuilder.h
+++ b/flang/include/flang/Lower/PFTBuilder.h
@@ -625,6 +625,12 @@ struct FunctionLikeUnit;
/// only instantiate referenced host module variables rather than all of them.
VariableList getHostModuleVariableList(const FunctionLikeUnit &funit);
+/// Create an ordered list of the equivalence sets and variables that appear in
+/// the scope of \p funit. Use associated variables that are not referenced in
+/// \p funit are left out to avoid instantiating unused module variables. The
+/// result is not cached.
+VariableList getScopeVariableList(const FunctionLikeUnit &funit);
+
void dump(VariableList &, std::string s = {}); // `s` is an optional dump label
/// Function-like units may contain evaluations (executable statements),
diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp
index 8e33239be0f8b..e501f65b43231 100644
--- a/flang/lib/Lower/Bridge.cpp
+++ b/flang/lib/Lower/Bridge.cpp
@@ -6211,10 +6211,11 @@ class FirConverter : public Fortran::lower::AbstractConverter {
Fortran::lower::pft::getHostModuleVariableList(funit))
instantiateVar(var, storeMap);
- // Map function equivalences and variables.
+ // Map function equivalences and variables. Use associated variables that
+ // are not referenced in the function-like unit are not instantiated.
mlir::Value primaryFuncResultStorage;
for (const Fortran::lower::pft::Variable &var :
- Fortran::lower::pft::getScopeVariableList(scope)) {
+ Fortran::lower::pft::getScopeVariableList(funit)) {
// Always instantiate aggregate storage blocks.
if (var.isAggregateStore()) {
instantiateVar(var, storeMap);
diff --git a/flang/lib/Lower/PFTBuilder.cpp b/flang/lib/Lower/PFTBuilder.cpp
index fad47e752d958..2fa5fe77cc5dc 100644
--- a/flang/lib/Lower/PFTBuilder.cpp
+++ b/flang/lib/Lower/PFTBuilder.cpp
@@ -1684,6 +1684,18 @@ struct SymbolDependenceAnalysis {
analyze(iter.second.get());
finalize();
}
+ /// Analyze the symbols of a subprogram \p scope, skipping the use associated
+ /// symbols whose ultimate symbols are not in \p referencedSymbols.
+ explicit SymbolDependenceAnalysis(
+ const semantics::Scope &scope,
+ const llvm::SetVector<const semantics::Symbol *> &referencedSymbols)
+ : referencedSymbols{&referencedSymbols} {
+ analyzeEquivalenceSets(scope);
+ for (const auto &iter : scope)
+ if (!isUnreferencedUseAssociated(iter.second.get()))
+ analyze(iter.second.get());
+ finalize();
+ }
explicit SymbolDependenceAnalysis(const semantics::Symbol &symbol) {
analyzeEquivalenceSets(symbol.owner());
analyze(symbol);
@@ -1727,7 +1739,8 @@ struct SymbolDependenceAnalysis {
// Analyze local, USEd, and host procedure scope equivalences.
for (const auto &iter : scope) {
const semantics::Symbol &ultimate = iter.second.get().GetUltimate();
- if (!skipSymbol(ultimate))
+ if (!skipSymbol(ultimate) &&
+ !isUnreferencedUseAssociated(iter.second.get()))
analyzeLocalEquivalenceSets(ultimate.owner());
}
// Add all aggregate stores to the front of the variable list.
@@ -1890,6 +1903,21 @@ struct SymbolDependenceAnalysis {
return depth;
}
+ /// Is \p sym a use associated symbol that is not referenced in the
+ /// subprogram being analyzed? Such symbols do not need to be instantiated.
+ bool isUnreferencedUseAssociated(const semantics::Symbol &sym) const {
+ if (!referencedSymbols || !sym.has<semantics::UseDetails>())
+ return false;
+ const semantics::Symbol &ultimate = sym.GetUltimate();
+ if (referencedSymbols->contains(&ultimate))
+ return false;
+ // Procedure pointers may be hidden behind a generic with the same name.
+ if (const auto *generic = ultimate.detailsIf<semantics::GenericDetails>())
+ if (const semantics::Symbol *specific = generic->specific())
+ return !referencedSymbols->contains(&specific->GetUltimate());
+ return true;
+ }
+
/// Skip symbol in alias analysis.
bool skipSymbol(const semantics::Symbol &sym) {
// Common block equivalences are largely managed by the front end.
@@ -1958,6 +1986,9 @@ struct SymbolDependenceAnalysis {
/// Set of scopes that have been analyzed for aliases.
llvm::SmallPtrSet<const semantics::Scope *, 4> analyzedScopes;
std::vector<Fortran::lower::pft::Variable::AggregateStore> stores;
+ /// Ultimate symbols referenced in the analyzed subprogram, if the analysis
+ /// is restricted to referenced use associated symbols.
+ const llvm::SetVector<const semantics::Symbol *> *referencedSymbols = nullptr;
};
} // namespace
@@ -2202,6 +2233,23 @@ lower::pft::getScopeVariableList(const semantics::Scope &scope) {
return sda.getVariableList();
}
+/// Create an ordered list of equivalences and variables in the scope of \p
+/// funit. Use associated variables that are not referenced in \p funit are
+/// left out. The result is not cached.
+lower::pft::VariableList
+lower::pft::getScopeVariableList(const FunctionLikeUnit &funit) {
+ const semantics::Scope &scope = funit.getScope();
+ LLVM_DEBUG(llvm::dbgs() << "\ngetScopeVariableList of [sub]program scope <"
+ << &scope << "> " << scope.GetName() << "\n");
+ llvm::SetVector<const semantics::Symbol *> referencedSymbols;
+ Fortran::lower::pft::visitAllSymbols(
+ funit, [&](const Fortran::semantics::Symbol &sym) {
+ referencedSymbols.insert(&sym.GetUltimate());
+ });
+ SymbolDependenceAnalysis sda(scope, referencedSymbols);
+ return sda.getVariableList();
+}
+
/// Create an ordered list of equivalences and variables that \p symbol
/// depends on (no caching). Include \p symbol at the end of the list.
lower::pft::VariableList
diff --git a/flang/test/Integration/MIF/coarray_allocation5.f90 b/flang/test/Integration/MIF/coarray_allocation5.f90
index 03b37768759d4..8f341cd33b252 100644
--- a/flang/test/Integration/MIF/coarray_allocation5.f90
+++ b/flang/test/Integration/MIF/coarray_allocation5.f90
@@ -16,6 +16,7 @@ end module m_coarray_test
program test
use m_coarray_test
+ module_coarray = 1.0
end program
! LLVM: @llvm.global_ctors = appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 0, ptr @__mif_save_coarrays_allocate, ptr null }]
diff --git a/flang/test/Lower/CUDA/cuda-cooperative.cuf b/flang/test/Lower/CUDA/cuda-cooperative.cuf
index 657a87c0b5a05..e13b8bdf40920 100644
--- a/flang/test/Lower/CUDA/cuda-cooperative.cuf
+++ b/flang/test/Lower/CUDA/cuda-cooperative.cuf
@@ -57,8 +57,9 @@ attributes(grid_global) subroutine w1()
gg = this_warp()
end subroutine
-! CHECK: %[[WARPSIZE:.*]] = fir.alloca i32 {bindc_name = "__builtin_warpsize", uniq_name = "_QM__fortran_builtinsEC__builtin_warpsize"}
-! CHECK: %[[WARPSIZE_DECL:.*]]:2 = hlfir.declare %[[WARPSIZE]] {uniq_name = "_QM__fortran_builtinsEC__builtin_warpsize"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK-LABEL: func.func @_QPw1
+! CHECK-NOT: __builtin_warpsize
+! CHECK: hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFw1Egg"}
! CHECK: %[[COALESCED_GROUP:.*]] = fir.alloca !fir.type<_QMcooperative_groupsTcoalesced_group{_QMcooperative_groupsTcoalesced_group.handle:!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>,size:i32,rank:i32}>
! CHECK: %[[C32:.*]] = arith.constant 32 : i32
! CHECK: %[[SIZE_COORD:.*]] = fir.coordinate_of %[[COALESCED_GROUP]], size : (!fir.ref<!fir.type<_QMcooperative_groupsTcoalesced_group{_QMcooperative_groupsTcoalesced_group.handle:!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>,size:i32,rank:i32}>>) -> !fir.ref<i32>
diff --git a/flang/test/Lower/CUDA/cuda-device-proc.cuf b/flang/test/Lower/CUDA/cuda-device-proc.cuf
index 97fb1ceefd8d2..62e076c8c8d43 100644
--- a/flang/test/Lower/CUDA/cuda-device-proc.cuf
+++ b/flang/test/Lower/CUDA/cuda-device-proc.cuf
@@ -467,11 +467,11 @@ attributes(global) subroutine test_bulk_g2s(a)
end subroutine
! CHECK-LABEL: func.func @_QPtest_bulk_g2s
-! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %4 {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)
-! CHECK: %[[DST:.*]]:2 = hlfir.declare %16(%17) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>)
-! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %19 {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)
+! CHECK: %[[DST:.*]]:2 = hlfir.declare %{{.*}}(%{{.*}}) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>)
+! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
! CHECK: %[[SRC:.*]] = hlfir.designate %{{.*}} (%{{.*}}) : (!fir.box<!fir.array<?xf64>>, i64) -> !fir.ref<f64>
-! CHECK: %[[COUNT_LOAD:.*]] = fir.load %20#0 : !fir.ref<i32>
+! CHECK: %[[COUNT_LOAD:.*]] = fir.load %[[COUNT]]#0 : !fir.ref<i32>
! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr
! CHECK: %[[BARRIER_3:.*]] = llvm.addrspacecast %[[BARRIER_PTR]] : !llvm.ptr to !llvm.ptr<3>
! CHECK: %[[DST_PTR:.*]] = fir.convert %[[DST]]#0 : (!fir.ref<!fir.array<1024xf64>>) -> !llvm.ptr
@@ -499,7 +499,7 @@ attributes(device) subroutine testAtomicCasLoop(aa, n)
end subroutine
! CHECK-LABEL: func.func @_QPtestatomiccasloop
-! CHECK: %[[CMP_XCHG:.*]] = llvm.cmpxchg %15, %c0_i32, %c1_i32 acq_rel monotonic : !llvm.ptr, i32
+! CHECK: %[[CMP_XCHG:.*]] = llvm.cmpxchg %{{.*}}, %c0_i32, %c1_i32 acq_rel monotonic : !llvm.ptr, i32
! CHECK: %[[CMP_XCHG_EV:.*]] = llvm.extractvalue %[[CMP_XCHG]][0] : !llvm.struct<(i32, i1)>
! CHECK: %{{.*}} = arith.constant 1 : i32
! CHECK: %{{.*}} = arith.cmpi eq, %[[CMP_XCHG_EV]], %{{.*}} : i32
diff --git a/flang/test/Lower/CUDA/cuda-devptr.cuf b/flang/test/Lower/CUDA/cuda-devptr.cuf
index 924bfa55c7e5b..f7eb429a60ee5 100644
--- a/flang/test/Lower/CUDA/cuda-devptr.cuf
+++ b/flang/test/Lower/CUDA/cuda-devptr.cuf
@@ -65,7 +65,7 @@ end subroutine
! CHECK-LABEL: func.func @_QPassign_c_devptr
! CHECK: %[[P:.*]] = fir.declare %arg0 dummy_scope %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFassign_c_devptrEp"}
-! CHECK: %[[C_DEVLOC_RES:.*]] = fir.declare %15 {uniq_name = ".tmp.intrinsic_result"} : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>) -> !fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>
+! CHECK: %[[C_DEVLOC_RES:.*]] = fir.declare %{{.*}} {uniq_name = ".tmp.intrinsic_result"} : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>) -> !fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>
! CHECK: %[[RES_CPTR_COORD:.*]] = fir.coordinate_of %[[C_DEVLOC_RES]], cptr : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>) -> !fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>>
! CHECK: %[[P_CPTR_COORD:.*]] = fir.coordinate_of %[[P]], cptr : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>>) -> !fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>>
! CHECK: %[[RES_ADDR_COORD:.*]] = fir.coordinate_of %[[RES_CPTR_COORD]], __address : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>>) -> !fir.ref<i64>
diff --git a/flang/test/Lower/CUDA/cuda-module-use.cuf b/flang/test/Lower/CUDA/cuda-module-use.cuf
index 130fefab24d90..8d61244f11256 100644
--- a/flang/test/Lower/CUDA/cuda-module-use.cuf
+++ b/flang/test/Lower/CUDA/cuda-module-use.cuf
@@ -5,7 +5,7 @@
subroutine sub1()
use cuf_mod
-! md = 1.0 ! currently a TODO
+ md = 1.0
end
! CHECK-LABEL: func.func @_QPsub1()
diff --git a/flang/test/Lower/OpenMP/threadprivate-use-association.f90 b/flang/test/Lower/OpenMP/threadprivate-use-association.f90
index 5daf87783acf9..dc472d1e9691b 100644
--- a/flang/test/Lower/OpenMP/threadprivate-use-association.f90
+++ b/flang/test/Lower/OpenMP/threadprivate-use-association.f90
@@ -50,8 +50,6 @@ program main
call sub()
! CHECK-LABEL: @_QQmain()
-!CHECK-DAG: [[ADDR0:%.*]] = fir.address_of(@blk_) : !fir.ref<!fir.array<24xi8>>
-!CHECK-DAG: [[NEWADDR0:%.*]] = omp.threadprivate [[ADDR0]] : !fir.ref<!fir.array<24xi8>> -> !fir.ref<!fir.array<24xi8>>
!CHECK-DAG: [[ADDR1:%.*]] = fir.address_of(@blk_) : !fir.ref<!fir.array<24xi8>>
!CHECK-DAG: [[NEWADDR1:%.*]] = omp.threadprivate [[ADDR1]] : !fir.ref<!fir.array<24xi8>> -> !fir.ref<!fir.array<24xi8>>
!CHECK-DAG: [[ADDR2:%.*]] = fir.address_of(@_QMtestEy) : !fir.ref<f32>
diff --git a/flang/test/Lower/declare-with-storage.f90 b/flang/test/Lower/declare-with-storage.f90
index 9004d239f3bdb..5abad0242fc24 100644
--- a/flang/test/Lower/declare-with-storage.f90
+++ b/flang/test/Lower/declare-with-storage.f90
@@ -27,6 +27,7 @@ end module data3
subroutine test1
use data1
use data2
+ call sub1(m1, m2, m3, m4)
end subroutine test1
! ALL-LABEL: func.func @_QPtest1() {
! HLFIR: %[[VAL_1:.*]] = fir.address_of(@common1_) : !fir.ref<tuple<i32, !fir.array<31xi8>>>
@@ -96,6 +97,7 @@ end subroutine test2
! Test common2 with equivalence.
subroutine test3
use data3
+ call sub3(x, y)
end subroutine test3
! ALL-LABEL: func.func @_QPtest3() {
! HLFIR: %[[VAL_1:.*]] = fir.address_of(@common2_) : !fir.ref<!fir.array<52xi8>>
diff --git a/flang/test/Lower/use_associated_variable_instantiation.f90 b/flang/test/Lower/use_associated_variable_instantiation.f90
new file mode 100644
index 0000000000000..13e26ab0f4087
--- /dev/null
+++ b/flang/test/Lower/use_associated_variable_instantiation.f90
@@ -0,0 +1,158 @@
+! Test that use associated module variables are only instantiated
+! in a subprogram when they are referenced in it.
+
+! RUN: %flang_fc1 -emit-hlfir %s -o - | FileCheck %s
+
+module test_use_instantiation
+ integer :: var1, var2, var3
+ integer :: a, b, c
+ equivalence (a, b)
+ integer :: n = 10, l = 3
+ integer, target :: pointee_target
+ real, target :: tgt(10)
+ integer, pointer :: p => null()
+ procedure(real), pointer :: proc_ptr
+ namelist /nml/ var3
+end module
+
+subroutine foo()
+ use test_use_instantiation
+ call bar(var1)
+end subroutine
+
+subroutine foo_equiv()
+ use test_use_instantiation
+ b = 1
+end subroutine
+
+subroutine foo_spec_expr()
+ use test_use_instantiation
+ real :: x(n)
+ call bar2(x)
+end subroutine
+
+subroutine foo_char_len()
+ use test_use_instantiation
+ character(l) :: ch
+ call bar3(ch)
+end subroutine
+
+subroutine foo_equiv_spec_expr()
+ use test_use_instantiation
+ real :: x(a)
+ call bar2(x)
+end subroutine
+
+subroutine foo_init_target()
+ use test_use_instantiation
+ real, pointer :: ptr(:) => tgt
+ call bar4(ptr)
+end subroutine
+
+subroutine foo_init_target_scalar()
+ use test_use_instantiation
+ integer, pointer :: iptr => pointee_target
+ call bar5(iptr)
+end subroutine
+
+subroutine foo_namelist()
+ use test_use_instantiation
+ read(*, nml)
+end subroutine
+
+subroutine foo_internal()
+ use test_use_instantiation
+ call internal()
+contains
+ subroutine internal()
+ var2 = 42
+ end subroutine
+end subroutine
+
+subroutine foo_unused()
+ use test_use_instantiation
+end subroutine
+
+! CHECK-LABEL: func.func @_QPfoo(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR1:.*]] = fir.address_of(@_QMtest_use_instantiationEvar1) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR1]]
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_equiv(
+! CHECK-NOT: hlfir.declare
+! CHECK: %[[ADDR_A:.*]] = fir.address_of(@_QMtest_use_instantiationEa) : !fir.ref<!fir.array<4xi8>>
+! CHECK: hlfir.declare %{{.*}} storage(%[[ADDR_A]][0]) {uniq_name = "_QMtest_use_instantiationEb"}
+! CHECK-NOT: hlfir.declare
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_spec_expr(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR_N:.*]] = fir.address_of(@_QMtest_use_instantiationEn) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR_N]]
+! CHECK-NOT: fir.address_of
+! CHECK: hlfir.declare %{{.*}} {uniq_name = "_QFfoo_spec_exprEx"}
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_char_len(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR_L:.*]] = fir.address_of(@_QMtest_use_instantiationEl) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR_L]]
+! CHECK-NOT: fir.address_of
+! CHECK: hlfir.declare %{{.*}} typeparams %{{.*}} {uniq_name = "_QFfoo_char_lenEch"}
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_equiv_spec_expr(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR_A2:.*]] = fir.address_of(@_QMtest_use_instantiationEa) : !fir.ref<!fir.array<4xi8>>
+! CHECK: hlfir.declare %{{.*}} storage(%[[ADDR_A2]][0]) {uniq_name = "_QMtest_use_instantiationEa"}
+! CHECK-NOT: fir.address_of
+! CHECK: hlfir.declare %{{.*}} {uniq_name = "_QFfoo_equiv_spec_exprEx"}
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_init_target(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR_TGT:.*]] = fir.address_of(@_QMtest_use_instantiationEtgt) : !fir.ref<!fir.array<10xf32>>
+! CHECK: hlfir.declare %[[ADDR_TGT]]
+! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
+! CHECK: hlfir.declare %{{.*}} {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFfoo_init_targetEptr"}
+! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_init_target_scalar(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR_PT:.*]] = fir.address_of(@_QMtest_use_instantiationEpointee_target) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR_PT]]
+! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
+! CHECK: hlfir.declare %{{.*}} {fortran_attrs = #fir.var_attrs<pointer>, uniq_name = "_QFfoo_init_target_scalarEiptr"}
+! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_namelist(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR3:.*]] = fir.address_of(@_QMtest_use_instantiationEvar3) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR3]]
+! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_internal(
+! CHECK-NOT: fir.address_of
+! CHECK: fir.address_of(@_QMtest_use_instantiationEvar2) : !fir.ref<i32>
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func private @_QFfoo_internalPinternal(
+! CHECK-NOT: fir.address_of
+! CHECK: %[[ADDR2:.*]] = fir.address_of(@_QMtest_use_instantiationEvar2) : !fir.ref<i32>
+! CHECK: hlfir.declare %[[ADDR2]]
+! CHECK-NOT: fir.address_of
+! CHECK: return
+
+! CHECK-LABEL: func.func @_QPfoo_unused(
+! CHECK-NOT: fir.address_of
+! CHECK-NOT: hlfir.declare
+! CHECK: return
>From 13c99a828048cd352ea07f73d0cff28f7340bd56 Mon Sep 17 00:00:00 2001
From: Jean Perier <jperier at nvidia.com>
Date: Wed, 9 Sep 2026 07:03:35 -0700
Subject: [PATCH 2/2] move CHECK under each subroutines
---
.../use_associated_variable_instantiation.f90 | 98 +++++++++----------
1 file changed, 44 insertions(+), 54 deletions(-)
diff --git a/flang/test/Lower/use_associated_variable_instantiation.f90 b/flang/test/Lower/use_associated_variable_instantiation.f90
index 13e26ab0f4087..26c2db237322d 100644
--- a/flang/test/Lower/use_associated_variable_instantiation.f90
+++ b/flang/test/Lower/use_associated_variable_instantiation.f90
@@ -19,60 +19,6 @@ subroutine foo()
use test_use_instantiation
call bar(var1)
end subroutine
-
-subroutine foo_equiv()
- use test_use_instantiation
- b = 1
-end subroutine
-
-subroutine foo_spec_expr()
- use test_use_instantiation
- real :: x(n)
- call bar2(x)
-end subroutine
-
-subroutine foo_char_len()
- use test_use_instantiation
- character(l) :: ch
- call bar3(ch)
-end subroutine
-
-subroutine foo_equiv_spec_expr()
- use test_use_instantiation
- real :: x(a)
- call bar2(x)
-end subroutine
-
-subroutine foo_init_target()
- use test_use_instantiation
- real, pointer :: ptr(:) => tgt
- call bar4(ptr)
-end subroutine
-
-subroutine foo_init_target_scalar()
- use test_use_instantiation
- integer, pointer :: iptr => pointee_target
- call bar5(iptr)
-end subroutine
-
-subroutine foo_namelist()
- use test_use_instantiation
- read(*, nml)
-end subroutine
-
-subroutine foo_internal()
- use test_use_instantiation
- call internal()
-contains
- subroutine internal()
- var2 = 42
- end subroutine
-end subroutine
-
-subroutine foo_unused()
- use test_use_instantiation
-end subroutine
-
! CHECK-LABEL: func.func @_QPfoo(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR1:.*]] = fir.address_of(@_QMtest_use_instantiationEvar1) : !fir.ref<i32>
@@ -80,6 +26,10 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of
! CHECK: return
+subroutine foo_equiv()
+ use test_use_instantiation
+ b = 1
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_equiv(
! CHECK-NOT: hlfir.declare
! CHECK: %[[ADDR_A:.*]] = fir.address_of(@_QMtest_use_instantiationEa) : !fir.ref<!fir.array<4xi8>>
@@ -87,6 +37,11 @@ subroutine foo_unused()
! CHECK-NOT: hlfir.declare
! CHECK: return
+subroutine foo_spec_expr()
+ use test_use_instantiation
+ real :: x(n)
+ call bar2(x)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_spec_expr(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR_N:.*]] = fir.address_of(@_QMtest_use_instantiationEn) : !fir.ref<i32>
@@ -96,6 +51,11 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of
! CHECK: return
+subroutine foo_char_len()
+ use test_use_instantiation
+ character(l) :: ch
+ call bar3(ch)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_char_len(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR_L:.*]] = fir.address_of(@_QMtest_use_instantiationEl) : !fir.ref<i32>
@@ -105,6 +65,11 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of
! CHECK: return
+subroutine foo_equiv_spec_expr()
+ use test_use_instantiation
+ real :: x(a)
+ call bar2(x)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_equiv_spec_expr(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR_A2:.*]] = fir.address_of(@_QMtest_use_instantiationEa) : !fir.ref<!fir.array<4xi8>>
@@ -114,6 +79,11 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of
! CHECK: return
+subroutine foo_init_target()
+ use test_use_instantiation
+ real, pointer :: ptr(:) => tgt
+ call bar4(ptr)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_init_target(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR_TGT:.*]] = fir.address_of(@_QMtest_use_instantiationEtgt) : !fir.ref<!fir.array<10xf32>>
@@ -123,6 +93,11 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
! CHECK: return
+subroutine foo_init_target_scalar()
+ use test_use_instantiation
+ integer, pointer :: iptr => pointee_target
+ call bar5(iptr)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_init_target_scalar(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR_PT:.*]] = fir.address_of(@_QMtest_use_instantiationEpointee_target) : !fir.ref<i32>
@@ -132,6 +107,10 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
! CHECK: return
+subroutine foo_namelist()
+ use test_use_instantiation
+ read(*, nml)
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_namelist(
! CHECK-NOT: fir.address_of
! CHECK: %[[ADDR3:.*]] = fir.address_of(@_QMtest_use_instantiationEvar3) : !fir.ref<i32>
@@ -139,6 +118,14 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of(@_QMtest_use_instantiationE
! CHECK: return
+subroutine foo_internal()
+ use test_use_instantiation
+ call internal()
+contains
+ subroutine internal()
+ var2 = 42
+ end subroutine
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_internal(
! CHECK-NOT: fir.address_of
! CHECK: fir.address_of(@_QMtest_use_instantiationEvar2) : !fir.ref<i32>
@@ -152,6 +139,9 @@ subroutine foo_unused()
! CHECK-NOT: fir.address_of
! CHECK: return
+subroutine foo_unused()
+ use test_use_instantiation
+end subroutine
! CHECK-LABEL: func.func @_QPfoo_unused(
! CHECK-NOT: fir.address_of
! CHECK-NOT: hlfir.declare
More information about the flang-commits
mailing list