Skip to content

Commit d925006

Browse files
authored
[flang] add fir.proc_attrs attributes to func.func (#110002)
BIND(C) ABI need care in the TargetRewrite pass. currently, we are not able to accurately identify fun.func that are BIND(C) in FIR (the fir.bindc_name is used in other contexts, like for program names). This patch adds the fir.proc_attrs to func.func just like it was done for calls recently. This replace the previous named attribute for PURE/ELEMENTAL/RECURSIVE (note that RECURSIVE is changed to NON_RECURSIVE, which brings more data since RECURSIVE is the default for procedures that do not have explicit RECURSIVE/NON_RECUSRIVE attributes).
1 parent 69ef3b1 commit d925006

17 files changed

+84
-64
lines changed

flang/include/flang/Lower/CallInterface.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -248,6 +248,7 @@ class CallInterface {
248248
CallInterface(Fortran::lower::AbstractConverter &c) : converter{c} {}
249249
/// CRTP handle.
250250
T &side() { return *static_cast<T *>(this); }
251+
const T &side() const { return *static_cast<const T *>(this); }
251252
/// Entry point to be called by child ctor to analyze the signature and
252253
/// create/find the mlir::func::FuncOp. Child needs to be initialized first.
253254
void declare();

flang/include/flang/Optimizer/Dialect/FIROpsSupport.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -160,6 +160,10 @@ static constexpr llvm::StringRef getFuncRecursiveAttrName() {
160160
return "fir.func_recursive";
161161
}
162162

163+
static constexpr llvm::StringRef getFortranProcedureFlagsAttrName() {
164+
return "fir.proc_attrs";
165+
}
166+
163167
// Attribute for an alloca that is a trivial adaptor for converting a value to
164168
// pass-by-ref semantics for a VALUE parameter. The optimizer may be able to
165169
// eliminate these.

flang/lib/Lower/CallInterface.cpp

Lines changed: 27 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -582,6 +582,7 @@ mlir::Value Fortran::lower::CalleeInterface::getHostAssociatedTuple() const {
582582

583583
static void addSymbolAttribute(mlir::func::FuncOp func,
584584
const Fortran::semantics::Symbol &sym,
585+
fir::FortranProcedureFlagsEnumAttr procAttrs,
585586
mlir::MLIRContext &mlirContext) {
586587
const Fortran::semantics::Symbol &ultimate = sym.GetUltimate();
587588
// The link between an internal procedure and its host procedure is lost
@@ -611,16 +612,8 @@ static void addSymbolAttribute(mlir::func::FuncOp func,
611612
}
612613
}
613614

614-
// Set procedure attributes to the func op.
615-
if (IsPureProcedure(sym))
616-
func->setAttr(fir::getFuncPureAttrName(),
617-
mlir::UnitAttr::get(&mlirContext));
618-
if (IsElementalProcedure(sym))
619-
func->setAttr(fir::getFuncElementalAttrName(),
620-
mlir::UnitAttr::get(&mlirContext));
621-
if (sym.attrs().test(Fortran::semantics::Attr::RECURSIVE))
622-
func->setAttr(fir::getFuncRecursiveAttrName(),
623-
mlir::UnitAttr::get(&mlirContext));
615+
if (procAttrs)
616+
func->setAttr(fir::getFortranProcedureFlagsAttrName(), procAttrs);
624617

625618
// Only add this on bind(C) functions for which the symbol is not reflected in
626619
// the current context.
@@ -703,6 +696,7 @@ void Fortran::lower::CallInterface<T>::declare() {
703696
func = fir::FirOpBuilder::getNamedFunction(module, symbolTable, name);
704697
if (!func) {
705698
mlir::Location loc = side().getCalleeLocation();
699+
mlir::MLIRContext &mlirContext = converter.getMLIRContext();
706700
mlir::FunctionType ty = genFunctionType();
707701
func =
708702
fir::FirOpBuilder::createFunction(loc, module, name, ty, symbolTable);
@@ -712,7 +706,8 @@ void Fortran::lower::CallInterface<T>::declare() {
712706
mlir::StringAttr::get(&converter.getMLIRContext(),
713707
sym->name().ToString()));
714708
} else {
715-
addSymbolAttribute(func, *sym, converter.getMLIRContext());
709+
addSymbolAttribute(func, *sym, getProcedureAttrs(&mlirContext),
710+
mlirContext);
716711
}
717712
}
718713
for (const auto &placeHolder : llvm::enumerate(inputs))
@@ -1550,8 +1545,8 @@ template <typename T>
15501545
fir::FortranProcedureFlagsEnumAttr
15511546
Fortran::lower::CallInterface<T>::getProcedureAttrs(
15521547
mlir::MLIRContext *mlirContext) const {
1548+
fir::FortranProcedureFlagsEnum flags = fir::FortranProcedureFlagsEnum::none;
15531549
if (characteristic) {
1554-
fir::FortranProcedureFlagsEnum flags = fir::FortranProcedureFlagsEnum::none;
15551550
if (characteristic->IsBindC())
15561551
flags = flags | fir::FortranProcedureFlagsEnum::bind_c;
15571552
if (characteristic->IsPure())
@@ -1560,12 +1555,27 @@ Fortran::lower::CallInterface<T>::getProcedureAttrs(
15601555
flags = flags | fir::FortranProcedureFlagsEnum::elemental;
15611556
// TODO:
15621557
// - SIMPLE: F2023, not yet handled by semantics.
1563-
// - NON_RECURSIVE: not part of the characteristics. Maybe this should
1564-
// simply not be part of FortranProcedureFlagsEnum since cannot accurately
1565-
// be known on the caller side.
1566-
if (flags != fir::FortranProcedureFlagsEnum::none)
1567-
return fir::FortranProcedureFlagsEnumAttr::get(mlirContext, flags);
15681558
}
1559+
1560+
if constexpr (std::is_same_v<Fortran::lower::CalleeInterface, T>) {
1561+
// Only gather and set NON_RECURSIVE for procedure definition. It is
1562+
// meaningless on calls since this is not part of Fortran characteristics
1563+
// (Fortran 2023 15.3.1) so there is no way to always know if the procedure
1564+
// called is recursive or not.
1565+
if (const Fortran::semantics::Symbol *sym = side().getProcedureSymbol()) {
1566+
// Note: By default procedures are RECURSIVE unless
1567+
// -fno-automatic/-save/-Msave is set. NON_RECURSIVE is is made explicit
1568+
// in that case in FIR.
1569+
if (sym->attrs().test(Fortran::semantics::Attr::NON_RECURSIVE) ||
1570+
(sym->owner().context().languageFeatures().IsEnabled(
1571+
Fortran::common::LanguageFeature::DefaultSave) &&
1572+
!sym->attrs().test(Fortran::semantics::Attr::RECURSIVE))) {
1573+
flags = flags | fir::FortranProcedureFlagsEnum::non_recursive;
1574+
}
1575+
}
1576+
}
1577+
if (flags != fir::FortranProcedureFlagsEnum::none)
1578+
return fir::FortranProcedureFlagsEnumAttr::get(mlirContext, flags);
15691579
return nullptr;
15701580
}
15711581

flang/test/Lower/CUDA/cuda-device-proc.cuf

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -26,11 +26,11 @@ end
2626
! CHECK: %{{.*}} = fir.call @__syncthreads_count(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (!fir.ref<i32>) -> i32
2727
! CHECK: %{{.*}} = fir.call @__syncthreads_or(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (!fir.ref<i32>) -> i32
2828

29-
! CHECK: func.func private @__syncthreads() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads"}
30-
! CHECK: func.func private @__syncwarp(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncwarp"}
31-
! CHECK: func.func private @__threadfence() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence"}
32-
! CHECK: func.func private @__threadfence_block() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_block"}
33-
! CHECK: func.func private @__threadfence_system() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_system"}
34-
! CHECK: func.func private @__syncthreads_and(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_and"}
35-
! CHECK: func.func private @__syncthreads_count(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_count"}
36-
! CHECK: func.func private @__syncthreads_or(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_or"}
29+
! CHECK: func.func private @__syncthreads() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads", fir.proc_attrs = #fir.proc_attrs<bind_c>}
30+
! CHECK: func.func private @__syncwarp(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncwarp", fir.proc_attrs = #fir.proc_attrs<bind_c>}
31+
! CHECK: func.func private @__threadfence() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence", fir.proc_attrs = #fir.proc_attrs<bind_c>}
32+
! CHECK: func.func private @__threadfence_block() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_block", fir.proc_attrs = #fir.proc_attrs<bind_c>}
33+
! CHECK: func.func private @__threadfence_system() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_system", fir.proc_attrs = #fir.proc_attrs<bind_c>}
34+
! CHECK: func.func private @__syncthreads_and(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_and", fir.proc_attrs = #fir.proc_attrs<bind_c>}
35+
! CHECK: func.func private @__syncthreads_count(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_count", fir.proc_attrs = #fir.proc_attrs<bind_c>}
36+
! CHECK: func.func private @__syncthreads_or(!fir.ref<i32> {cuf.data_attr = #cuf.cuda<device>}) -> i32 attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads_or", fir.proc_attrs = #fir.proc_attrs<bind_c>}

flang/test/Lower/HLFIR/bindc-value-derived.f90

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,7 @@ subroutine test(x) bind(c)
1414
call use_it(x%i)
1515
end subroutine
1616
! CHECK-LABEL: func.func @test(
17-
! CHECK-SAME: %[[VAL_0:.*]]: !fir.type<_QMbindc_byvalTt{i:i32}> {fir.bindc_name = "x"}) attributes {fir.bindc_name = "test"} {
17+
! CHECK-SAME: %[[VAL_0:.*]]: !fir.type<_QMbindc_byvalTt{i:i32}>
1818
! CHECK: %[[VAL_1:.*]] = fir.alloca !fir.type<_QMbindc_byvalTt{i:i32}>
1919
! CHECK: fir.store %[[VAL_0]] to %[[VAL_1]] : !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
2020
! CHECK: %[[VAL_2:.*]]:2 = hlfir.declare %[[VAL_1]] dummy_scope %{{[0-9]+}} {fortran_attrs = #fir.var_attrs<value>, uniq_name = "_QMbindc_byvalFtestEx"} : (!fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>, !fir.dscope) -> (!fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>, !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>)
@@ -28,7 +28,7 @@ subroutine call_it(x)
2828
call test(x)
2929
end subroutine
3030
! CHECK-LABEL: func.func @_QMbindc_byvalPcall_it(
31-
! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>> {fir.bindc_name = "x"}) {
31+
! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
3232
! CHECK: %[[VAL_1:.*]]:2 = hlfir.declare %[[VAL_0]] dummy_scope %{{[0-9]+}} {uniq_name = "_QMbindc_byvalFcall_itEx"} : (!fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>, !fir.dscope) -> (!fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>, !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>)
3333
! CHECK: %[[VAL_2:.*]] = fir.load %[[VAL_1]]#1 : !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
3434
! CHECK: fir.call @test(%[[VAL_2]]) proc_attrs<bind_c> fastmath<contract> : (!fir.type<_QMbindc_byvalTt{i:i32}>) -> ()

flang/test/Lower/HLFIR/block_bindc_pocs.f90

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@ end module m
1111
!CHECK-DAG: %[[S0:.*]] = llvm.intr.stacksave : !llvm.ptr
1212
!CHECK-DAG: fir.call @test_proc() proc_attrs<bind_c> fastmath<contract> : () -> ()
1313
!CHECK-DAG: llvm.intr.stackrestore %[[S0]] : !llvm.ptr
14-
!CHECK-DAG: func.func private @test_proc() attributes {fir.bindc_name = "test_proc"}
14+
!CHECK-DAG: func.func private @test_proc() attributes {fir.bindc_name = "test_proc", fir.proc_attrs = #fir.proc_attrs<bind_c>}
1515
subroutine test
1616
BLOCK
1717
use m

flang/test/Lower/Intrinsics/signal.f90

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,14 +4,14 @@
44
module m
55
contains
66
! CHECK-LABEL: func.func @handler(
7-
! CHECK-SAME: %[[VAL_0:.*]]: i32 {fir.bindc_name = "signum"}) attributes {fir.bindc_name = "handler"} {
7+
! CHECK-SAME: %[[VAL_0:.*]]: i32
88
subroutine handler(signum) bind(C)
99
use iso_c_binding, only: c_int
1010
integer(c_int), value :: signum
1111
end subroutine
1212

1313
! CHECK-LABEL: func.func @_QMmPsetup_signals(
14-
! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<i32> {fir.bindc_name = "optional_status", fir.optional}) {
14+
! CHECK-SAME: %[[VAL_0:.*]]: !fir.ref<i32>
1515
subroutine setup_signals(optional_status)
1616
! not portable accross systems
1717
integer, parameter :: SIGFPE = 8

flang/test/Lower/OpenMP/declare-target-func-and-subr.f90

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -154,7 +154,7 @@ SUBROUTINE SUBR_DEFAULT_EXTENDEDLIST()
154154
!! -----
155155

156156
! DEVICE-LABEL: func.func @_QPrecursive_declare_target
157-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}
157+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}
158158
RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
159159
!$omp declare target to(RECURSIVE_DECLARE_TARGET) device_type(nohost)
160160
INTEGER :: INCREMENT, K
@@ -166,7 +166,7 @@ RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
166166
END FUNCTION RECURSIVE_DECLARE_TARGET
167167

168168
! DEVICE-LABEL: func.func @_QPrecursive_declare_target_enter
169-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
169+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
170170
RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT) RESULT(K)
171171
!$omp declare target enter(RECURSIVE_DECLARE_TARGET_ENTER) device_type(nohost)
172172
INTEGER :: INCREMENT, K

flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap-enter.f90

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ end function target_function_test_host
105105
!! -----
106106

107107
! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
108-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
108+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
109109
recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
110110
!$omp declare target enter(implicitly_captured_with_dev_type_recursive) device_type(host)
111111
integer :: increment, k
@@ -174,7 +174,7 @@ recursive subroutine implicitly_captured_recursive(increment)
174174
end program
175175

176176
! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
177-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
177+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
178178
recursive subroutine implicitly_captured_recursive(increment)
179179
integer :: increment
180180
if (increment == 10) then

flang/test/Lower/OpenMP/declare-target-implicit-func-and-subr-cap.f90

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -131,7 +131,7 @@ end function target_function_test_host
131131
!! -----
132132

133133
! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
134-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}}
134+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}}
135135
recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
136136
!$omp declare target to(implicitly_captured_with_dev_type_recursive) device_type(host)
137137
integer :: increment, k
@@ -200,7 +200,7 @@ recursive subroutine implicitly_captured_recursive(increment)
200200
end program
201201

202202
! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
203-
! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}}
203+
! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}}
204204
recursive subroutine implicitly_captured_recursive(increment)
205205
integer :: increment
206206
if (increment == 10) then

0 commit comments

Comments
 (0)