[flang][runtime] Enable more APIs in the offload build. (#75996)

This patch enables more numeric (mod, sum, matmul, etc.) APIs,
and some others.

I added new macros to disable warnings about using C++ STD methods
like operators of std::complex, which do not have __device__ attribute.
This may probably result in unresolved references, if the header files
implementation relies on libstdc++. I will need to follow up on this.

Author: vzakhari

flang/include/flang/ISO_Fortran_binding.h

@@ -189,8 +189,8 @@ RT_API_ATTRS void *CFI_address(
 RT_API_ATTRS int CFI_allocate(CFI_cdesc_t *, const CFI_index_t lower_bounds[],
     const CFI_index_t upper_bounds[], size_t elem_len);
 RT_API_ATTRS int CFI_deallocate(CFI_cdesc_t *);
-int CFI_establish(CFI_cdesc_t *, void *base_addr, CFI_attribute_t, CFI_type_t,
-    size_t elem_len, CFI_rank_t, const CFI_index_t extents[]);
+RT_API_ATTRS int CFI_establish(CFI_cdesc_t *, void *base_addr, CFI_attribute_t,
+    CFI_type_t, size_t elem_len, CFI_rank_t, const CFI_index_t extents[]);
 RT_API_ATTRS int CFI_is_contiguous(const CFI_cdesc_t *);
 RT_API_ATTRS int CFI_section(CFI_cdesc_t *, const CFI_cdesc_t *source,
     const CFI_index_t lower_bounds[], const CFI_index_t upper_bounds[],

flang/include/flang/Runtime/allocatable.h

@@ -26,22 +26,22 @@ extern "C" {
 // A descriptor must be initialized before being used for any purpose,
 // but needs reinitialization in a deallocated state only when there is
 // a change of type, rank, or corank.
-void RTNAME(AllocatableInitIntrinsic)(
+void RTDECL(AllocatableInitIntrinsic)(
     Descriptor &, TypeCategory, int kind, int rank = 0, int corank = 0);
-void RTNAME(AllocatableInitCharacter)(Descriptor &, SubscriptValue length = 0,
+void RTDECL(AllocatableInitCharacter)(Descriptor &, SubscriptValue length = 0,
     int kind = 1, int rank = 0, int corank = 0);
-void RTNAME(AllocatableInitDerived)(
+void RTDECL(AllocatableInitDerived)(
     Descriptor &, const typeInfo::DerivedType &, int rank = 0, int corank = 0);
 
 // Initializes the descriptor for an allocatable of intrinsic or derived type.
 // These functions are meant to be used in the allocate statement lowering. If
 // the descriptor is allocated, the initialization is skiped so the error
 // handling can be done by AllocatableAllocate.
-void RTNAME(AllocatableInitIntrinsicForAllocate)(
+void RTDECL(AllocatableInitIntrinsicForAllocate)(
     Descriptor &, TypeCategory, int kind, int rank = 0, int corank = 0);
-void RTNAME(AllocatableInitCharacterForAllocate)(Descriptor &,
+void RTDECL(AllocatableInitCharacterForAllocate)(Descriptor &,
     SubscriptValue length = 0, int kind = 1, int rank = 0, int corank = 0);
-void RTNAME(AllocatableInitDerivedForAllocate)(
+void RTDECL(AllocatableInitDerivedForAllocate)(
     Descriptor &, const typeInfo::DerivedType &, int rank = 0, int corank = 0);
 
 // Checks that an allocatable is not already allocated in statements
@@ -50,37 +50,37 @@ void RTNAME(AllocatableInitDerivedForAllocate)(
 // (If there's no STAT=, the error will be caught later anyway, but
 // this API allows the error to be caught before descriptor is modified.)
 // Return 0 on success (deallocated state), else the STAT= value.
-int RTNAME(AllocatableCheckAllocated)(Descriptor &,
+int RTDECL(AllocatableCheckAllocated)(Descriptor &,
     const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
     int sourceLine = 0);
 
 // For MOLD= allocation; sets bounds, cobounds, and length type
 // parameters from another descriptor. The destination descriptor must
 // be initialized and deallocated.
-void RTNAME(AllocatableApplyMold)(
+void RTDECL(AllocatableApplyMold)(
     Descriptor &, const Descriptor &mold, int rank = 0);
 
 // Explicitly sets the bounds and length type parameters of an initialized
 // deallocated allocatable.
-void RTNAME(AllocatableSetBounds)(
+void RTDECL(AllocatableSetBounds)(
     Descriptor &, int zeroBasedDim, SubscriptValue lower, SubscriptValue upper);
 
 // The upper cobound is ignored for the last codimension.
-void RTNAME(AllocatableSetCoBounds)(Descriptor &, int zeroBasedCoDim,
+void RTDECL(AllocatableSetCoBounds)(Descriptor &, int zeroBasedCoDim,
     SubscriptValue lower, SubscriptValue upper = 0);
 
 // Length type parameters are indexed in declaration order; i.e., 0 is the
 // first length type parameter in the deepest base type.  (Not for use
 // with CHARACTER; see above.)
-void RTNAME(AllocatableSetDerivedLength)(
+void RTDECL(AllocatableSetDerivedLength)(
     Descriptor &, int which, SubscriptValue);
 
 // When an explicit type-spec appears in an ALLOCATE statement for an
 // allocatable with an explicit (non-deferred) length type paramater for
 // a derived type or CHARACTER value, the explicit value has to match
 // the length type parameter's value.  This API checks that requirement.
 // Returns 0 for success, or the STAT= value on failure with hasStat==true.
-int RTNAME(AllocatableCheckLengthParameter)(Descriptor &,
+int RTDECL(AllocatableCheckLengthParameter)(Descriptor &,
     int which /* 0 for CHARACTER length */, SubscriptValue other,
     bool hasStat = false, const Descriptor *errMsg = nullptr,
     const char *sourceFile = nullptr, int sourceLine = 0);
@@ -94,38 +94,38 @@ int RTNAME(AllocatableCheckLengthParameter)(Descriptor &,
 // Successfully allocated memory is initialized if the allocatable has a
 // derived type, and is always initialized by AllocatableAllocateSource().
 // Performs all necessary coarray synchronization and validation actions.
-int RTNAME(AllocatableAllocate)(Descriptor &, bool hasStat = false,
+int RTDECL(AllocatableAllocate)(Descriptor &, bool hasStat = false,
     const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
     int sourceLine = 0);
-int RTNAME(AllocatableAllocateSource)(Descriptor &, const Descriptor &source,
+int RTDECL(AllocatableAllocateSource)(Descriptor &, const Descriptor &source,
     bool hasStat = false, const Descriptor *errMsg = nullptr,
     const char *sourceFile = nullptr, int sourceLine = 0);
 
 // Implements the intrinsic subroutine MOVE_ALLOC (16.9.137 in F'2018,
 // but note the order of first two arguments is reversed for consistency
 // with the other APIs for allocatables.)  The destination descriptor
 // must be initialized.
-std::int32_t RTNAME(MoveAlloc)(Descriptor &to, Descriptor &from,
+std::int32_t RTDECL(MoveAlloc)(Descriptor &to, Descriptor &from,
     const typeInfo::DerivedType *, bool hasStat = false,
     const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
     int sourceLine = 0);
 
 // Deallocates an allocatable.  Finalizes elements &/or components as needed.
 // The allocatable is left in an initialized state suitable for reallocation
 // with the same bounds, cobounds, and length type parameters.
-int RTNAME(AllocatableDeallocate)(Descriptor &, bool hasStat = false,
+int RTDECL(AllocatableDeallocate)(Descriptor &, bool hasStat = false,
     const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
     int sourceLine = 0);
 
 // Same as AllocatableDeallocate but also set the dynamic type as the declared
 // type as mentioned in 7.3.2.3 note 7.
-int RTNAME(AllocatableDeallocatePolymorphic)(Descriptor &,
+int RTDECL(AllocatableDeallocatePolymorphic)(Descriptor &,
     const typeInfo::DerivedType *, bool hasStat = false,
     const Descriptor *errMsg = nullptr, const char *sourceFile = nullptr,
     int sourceLine = 0);
 
 // Variant of above that does not finalize; for intermediate results
-void RTNAME(AllocatableDeallocateNoFinal)(
+void RTDECL(AllocatableDeallocateNoFinal)(
     Descriptor &, const char *sourceFile = nullptr, int sourceLine = 0);
 } // extern "C"
 } // namespace Fortran::runtime

flang/include/flang/Runtime/api-attrs.h

@@ -121,4 +121,15 @@
 #undef RT_DEVICE_COMPILATION
 #endif
 
+#if defined(__CUDACC__)
+#define RT_DIAG_PUSH _Pragma("nv_diagnostic push")
+#define RT_DIAG_POP _Pragma("nv_diagnostic pop")
+#define RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN \
+  _Pragma("nv_diag_suppress 20011") _Pragma("nv_diag_suppress 20014")
+#else /* !defined(__CUDACC__) */
+#define RT_DIAG_PUSH
+#define RT_DIAG_POP
+#define RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN
+#endif /* !defined(__CUDACC__) */
+
 #endif /* !FORTRAN_RUNTIME_API_ATTRS_H_ */

flang/include/flang/Runtime/derived-api.h

@@ -29,37 +29,37 @@ extern "C" {
 // Initializes and allocates an object's components, if it has a derived type
 // with any default component initialization or automatic components.
 // The descriptor must be initialized and non-null.
-void RTNAME(Initialize)(
+void RTDECL(Initialize)(
     const Descriptor &, const char *sourceFile = nullptr, int sourceLine = 0);
 
 // Finalizes an object and its components.  Deallocates any
 // allocatable/automatic components.  Does not deallocate the descriptor's
 // storage.
-void RTNAME(Destroy)(const Descriptor &);
+void RTDECL(Destroy)(const Descriptor &);
 
 // Finalizes the object and its components.
-void RTNAME(Finalize)(
+void RTDECL(Finalize)(
     const Descriptor &, const char *sourceFile = nullptr, int sourceLine = 0);
 
 /// Deallocates any allocatable/automatic components.
 /// Does not deallocate the descriptor's storage.
 /// Does not perform any finalization.
-void RTNAME(DestroyWithoutFinalization)(const Descriptor &);
+void RTDECL(DestroyWithoutFinalization)(const Descriptor &);
 
 // Intrinsic or defined assignment, with scalar expansion but not type
 // conversion.
-void RTNAME(Assign)(const Descriptor &, const Descriptor &,
+void RTDECL(Assign)(const Descriptor &, const Descriptor &,
     const char *sourceFile = nullptr, int sourceLine = 0);
 
 // Perform the test of the CLASS IS type guard statement of the SELECT TYPE
 // construct.
-bool RTNAME(ClassIs)(const Descriptor &, const typeInfo::DerivedType &);
+bool RTDECL(ClassIs)(const Descriptor &, const typeInfo::DerivedType &);
 
 // Perform the test of the SAME_TYPE_AS intrinsic.
-bool RTNAME(SameTypeAs)(const Descriptor &, const Descriptor &);
+bool RTDECL(SameTypeAs)(const Descriptor &, const Descriptor &);
 
 // Perform the test of the EXTENDS_TYPE_OF intrinsic.
-bool RTNAME(ExtendsTypeOf)(const Descriptor &, const Descriptor &);
+bool RTDECL(ExtendsTypeOf)(const Descriptor &, const Descriptor &);
 
 } // extern "C"
 } // namespace Fortran::runtime

flang/include/flang/Runtime/matmul-transpose.h

@@ -18,12 +18,12 @@ extern "C" {
 // The most general MATMUL(TRANSPOSE()).  All type and shape information is
 // taken from the arguments' descriptors, and the result is dynamically
 // allocated.
-void RTNAME(MatmulTranspose)(Descriptor &, const Descriptor &,
+void RTDECL(MatmulTranspose)(Descriptor &, const Descriptor &,
     const Descriptor &, const char *sourceFile = nullptr, int line = 0);
 
 // A non-allocating variant; the result's descriptor must be established
 // and have a valid base address.
-void RTNAME(MatmulTransposeDirect)(const Descriptor &, const Descriptor &,
+void RTDECL(MatmulTransposeDirect)(const Descriptor &, const Descriptor &,
     const Descriptor &, const char *sourceFile = nullptr, int line = 0);
 } // extern "C"
 } // namespace Fortran::runtime

flang/include/flang/Runtime/matmul.h

@@ -17,12 +17,12 @@ extern "C" {
 
 // The most general MATMUL.  All type and shape information is taken from the
 // arguments' descriptors, and the result is dynamically allocated.
-void RTNAME(Matmul)(Descriptor &, const Descriptor &, const Descriptor &,
+void RTDECL(Matmul)(Descriptor &, const Descriptor &, const Descriptor &,
     const char *sourceFile = nullptr, int line = 0);
 
 // A non-allocating variant; the result's descriptor must be established
 // and have a valid base address.
-void RTNAME(MatmulDirect)(const Descriptor &, const Descriptor &,
+void RTDECL(MatmulDirect)(const Descriptor &, const Descriptor &,
     const Descriptor &, const char *sourceFile = nullptr, int line = 0);
 } // extern "C"
 } // namespace Fortran::runtime