[flang] add fir.proc_attrs attributes to func.func (llvm#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).

Author: jeanPerier

flang/include/flang/Lower/CallInterface.h

@@ -248,6 +248,7 @@ class CallInterface {
   CallInterface(Fortran::lower::AbstractConverter &c) : converter{c} {}
   /// CRTP handle.
   T &side() { return *static_cast<T *>(this); }
+  const T &side() const { return *static_cast<const T *>(this); }
   /// Entry point to be called by child ctor to analyze the signature and
   /// create/find the mlir::func::FuncOp. Child needs to be initialized first.
   void declare();

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

@@ -160,6 +160,10 @@ static constexpr llvm::StringRef getFuncRecursiveAttrName() {
   return "fir.func_recursive";
 }
 
+static constexpr llvm::StringRef getFortranProcedureFlagsAttrName() {
+  return "fir.proc_attrs";
+}
+
 // Attribute for an alloca that is a trivial adaptor for converting a value to
 // pass-by-ref semantics for a VALUE parameter. The optimizer may be able to
 // eliminate these.

flang/lib/Lower/CallInterface.cpp

@@ -582,6 +582,7 @@ mlir::Value Fortran::lower::CalleeInterface::getHostAssociatedTuple() const {
 
 static void addSymbolAttribute(mlir::func::FuncOp func,
                                const Fortran::semantics::Symbol &sym,
+                               fir::FortranProcedureFlagsEnumAttr procAttrs,
                                mlir::MLIRContext &mlirContext) {
   const Fortran::semantics::Symbol &ultimate = sym.GetUltimate();
   // The link between an internal procedure and its host procedure is lost
@@ -611,16 +612,8 @@ static void addSymbolAttribute(mlir::func::FuncOp func,
     }
   }
 
-  // Set procedure attributes to the func op.
-  if (IsPureProcedure(sym))
-    func->setAttr(fir::getFuncPureAttrName(),
-                  mlir::UnitAttr::get(&mlirContext));
-  if (IsElementalProcedure(sym))
-    func->setAttr(fir::getFuncElementalAttrName(),
-                  mlir::UnitAttr::get(&mlirContext));
-  if (sym.attrs().test(Fortran::semantics::Attr::RECURSIVE))
-    func->setAttr(fir::getFuncRecursiveAttrName(),
-                  mlir::UnitAttr::get(&mlirContext));
+  if (procAttrs)
+    func->setAttr(fir::getFortranProcedureFlagsAttrName(), procAttrs);
 
   // Only add this on bind(C) functions for which the symbol is not reflected in
   // the current context.
@@ -703,6 +696,7 @@ void Fortran::lower::CallInterface<T>::declare() {
     func = fir::FirOpBuilder::getNamedFunction(module, symbolTable, name);
     if (!func) {
       mlir::Location loc = side().getCalleeLocation();
+      mlir::MLIRContext &mlirContext = converter.getMLIRContext();
       mlir::FunctionType ty = genFunctionType();
       func =
           fir::FirOpBuilder::createFunction(loc, module, name, ty, symbolTable);
@@ -712,7 +706,8 @@ void Fortran::lower::CallInterface<T>::declare() {
                         mlir::StringAttr::get(&converter.getMLIRContext(),
                                               sym->name().ToString()));
         } else {
-          addSymbolAttribute(func, *sym, converter.getMLIRContext());
+          addSymbolAttribute(func, *sym, getProcedureAttrs(&mlirContext),
+                             mlirContext);
         }
       }
       for (const auto &placeHolder : llvm::enumerate(inputs))
@@ -1550,8 +1545,8 @@ template <typename T>
 fir::FortranProcedureFlagsEnumAttr
 Fortran::lower::CallInterface<T>::getProcedureAttrs(
     mlir::MLIRContext *mlirContext) const {
+  fir::FortranProcedureFlagsEnum flags = fir::FortranProcedureFlagsEnum::none;
   if (characteristic) {
-    fir::FortranProcedureFlagsEnum flags = fir::FortranProcedureFlagsEnum::none;
     if (characteristic->IsBindC())
       flags = flags | fir::FortranProcedureFlagsEnum::bind_c;
     if (characteristic->IsPure())
@@ -1560,12 +1555,27 @@ Fortran::lower::CallInterface<T>::getProcedureAttrs(
       flags = flags | fir::FortranProcedureFlagsEnum::elemental;
     // TODO:
     // - SIMPLE: F2023, not yet handled by semantics.
-    // - NON_RECURSIVE: not part of the characteristics. Maybe this should
-    //   simply not be part of FortranProcedureFlagsEnum since cannot accurately
-    //   be known on the caller side.
-    if (flags != fir::FortranProcedureFlagsEnum::none)
-      return fir::FortranProcedureFlagsEnumAttr::get(mlirContext, flags);
   }
+
+  if constexpr (std::is_same_v<Fortran::lower::CalleeInterface, T>) {
+    // Only gather and set NON_RECURSIVE for procedure definition. It is
+    // meaningless on calls since this is not part of Fortran characteristics
+    // (Fortran 2023 15.3.1) so there is no way to always know if the procedure
+    // called is recursive or not.
+    if (const Fortran::semantics::Symbol *sym = side().getProcedureSymbol()) {
+      // Note: By default procedures are RECURSIVE unless
+      // -fno-automatic/-save/-Msave is set. NON_RECURSIVE is is made explicit
+      // in that case in FIR.
+      if (sym->attrs().test(Fortran::semantics::Attr::NON_RECURSIVE) ||
+          (sym->owner().context().languageFeatures().IsEnabled(
+               Fortran::common::LanguageFeature::DefaultSave) &&
+           !sym->attrs().test(Fortran::semantics::Attr::RECURSIVE))) {
+        flags = flags | fir::FortranProcedureFlagsEnum::non_recursive;
+      }
+    }
+  }
+  if (flags != fir::FortranProcedureFlagsEnum::none)
+    return fir::FortranProcedureFlagsEnumAttr::get(mlirContext, flags);
   return nullptr;
 }
 

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

@@ -26,11 +26,11 @@ end
 ! CHECK: %{{.*}} = fir.call @__syncthreads_count(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (!fir.ref<i32>) -> i32
 ! CHECK: %{{.*}} = fir.call @__syncthreads_or(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (!fir.ref<i32>) -> i32
 
-! CHECK: func.func private @__syncthreads() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__syncthreads"}
-! 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"}
-! CHECK: func.func private @__threadfence() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence"}
-! CHECK: func.func private @__threadfence_block() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_block"}
-! CHECK: func.func private @__threadfence_system() attributes {cuf.proc_attr = #cuf.cuda_proc<device>, fir.bindc_name = "__threadfence_system"}
-! 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"}
-! 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"}
-! 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"}
+! 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>}
+! 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>}
+! 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>}
+! 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>}
+! 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>}
+! 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>}
+! 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>}
+! 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

@@ -14,7 +14,7 @@ subroutine test(x) bind(c)
     call use_it(x%i)
   end subroutine
 ! CHECK-LABEL:   func.func @test(
-! CHECK-SAME:                    %[[VAL_0:.*]]: !fir.type<_QMbindc_byvalTt{i:i32}> {fir.bindc_name = "x"}) attributes {fir.bindc_name = "test"} {
+! CHECK-SAME:                    %[[VAL_0:.*]]: !fir.type<_QMbindc_byvalTt{i:i32}>
 ! CHECK:           %[[VAL_1:.*]] = fir.alloca !fir.type<_QMbindc_byvalTt{i:i32}>
 ! CHECK:           fir.store %[[VAL_0]] to %[[VAL_1]] : !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
 ! 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)
     call test(x)
   end subroutine
 ! CHECK-LABEL:   func.func @_QMbindc_byvalPcall_it(
-! CHECK-SAME:                                      %[[VAL_0:.*]]: !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>> {fir.bindc_name = "x"}) {
+! CHECK-SAME:                                      %[[VAL_0:.*]]: !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
 ! 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}>>)
 ! CHECK:           %[[VAL_2:.*]] = fir.load %[[VAL_1]]#1 : !fir.ref<!fir.type<_QMbindc_byvalTt{i:i32}>>
 ! 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

@@ -11,7 +11,7 @@ end module m
 !CHECK-DAG: %[[S0:.*]] = llvm.intr.stacksave : !llvm.ptr
 !CHECK-DAG: fir.call @test_proc() proc_attrs<bind_c> fastmath<contract> : () -> ()
 !CHECK-DAG: llvm.intr.stackrestore %[[S0]] : !llvm.ptr
-!CHECK-DAG: func.func private @test_proc() attributes {fir.bindc_name = "test_proc"}
+!CHECK-DAG: func.func private @test_proc() attributes {fir.bindc_name = "test_proc", fir.proc_attrs = #fir.proc_attrs<bind_c>}
 subroutine test
     BLOCK
         use m

flang/test/Lower/Intrinsics/signal.f90

@@ -4,14 +4,14 @@
 module m
 contains
 ! CHECK-LABEL:   func.func @handler(
-! CHECK-SAME:                       %[[VAL_0:.*]]: i32 {fir.bindc_name = "signum"}) attributes {fir.bindc_name = "handler"} {
+! CHECK-SAME:                       %[[VAL_0:.*]]: i32
   subroutine handler(signum) bind(C)
     use iso_c_binding, only: c_int
     integer(c_int), value :: signum
   end subroutine
 
 ! CHECK-LABEL:   func.func @_QMmPsetup_signals(
-! CHECK-SAME:                                  %[[VAL_0:.*]]: !fir.ref<i32> {fir.bindc_name = "optional_status", fir.optional}) {
+! CHECK-SAME:                                  %[[VAL_0:.*]]: !fir.ref<i32>
   subroutine setup_signals(optional_status)
     ! not portable accross systems
     integer, parameter :: SIGFPE = 8

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

@@ -154,7 +154,7 @@ SUBROUTINE SUBR_DEFAULT_EXTENDEDLIST()
 !! -----
 
 ! DEVICE-LABEL: func.func @_QPrecursive_declare_target
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}
 RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
 !$omp declare target to(RECURSIVE_DECLARE_TARGET) device_type(nohost)
     INTEGER :: INCREMENT, K
@@ -166,7 +166,7 @@ RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET(INCREMENT) RESULT(K)
 END FUNCTION RECURSIVE_DECLARE_TARGET
 
 ! DEVICE-LABEL: func.func @_QPrecursive_declare_target_enter
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}
 RECURSIVE FUNCTION RECURSIVE_DECLARE_TARGET_ENTER(INCREMENT) RESULT(K)
 !$omp declare target enter(RECURSIVE_DECLARE_TARGET_ENTER) device_type(nohost)
     INTEGER :: INCREMENT, K

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

@@ -105,7 +105,7 @@ end function target_function_test_host
 !! -----
 
 ! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (enter)>{{.*}}}
 recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
 !$omp declare target enter(implicitly_captured_with_dev_type_recursive) device_type(host)
    integer :: increment, k
@@ -174,7 +174,7 @@ recursive subroutine implicitly_captured_recursive(increment)
 end program
 
 ! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (enter)>{{.*}}}
 recursive subroutine implicitly_captured_recursive(increment)
    integer :: increment
    if (increment == 10) then

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

@@ -131,7 +131,7 @@ end function target_function_test_host
 !! -----
 
 ! DEVICE-LABEL: func.func @_QPimplicitly_captured_with_dev_type_recursive
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (any), capture_clause = (to)>{{.*}}}
 recursive function implicitly_captured_with_dev_type_recursive(increment) result(k)
 !$omp declare target to(implicitly_captured_with_dev_type_recursive) device_type(host)
    integer :: increment, k
@@ -200,7 +200,7 @@ recursive subroutine implicitly_captured_recursive(increment)
 end program
 
 ! DEVICE-LABEL: func.func @_QPimplicitly_captured_recursive
-! DEVICE-SAME: {{.*}}attributes {fir.func_recursive, omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}}
+! DEVICE-SAME: {{.*}}attributes {omp.declare_target = #omp.declaretarget<device_type = (nohost), capture_clause = (to)>{{.*}}}
 recursive subroutine implicitly_captured_recursive(increment)
    integer :: increment
    if (increment == 10) then