[RFC][flang] Experimental device build of Flang runtime.

[vzakhari]

Summary

These are initial changes to experiment with building the Fortran runtime
as a CUDA or OpenMP target offload library.

The initial patch defines a set of macros that have to be used consistently
in Flang runtime source code so that it can be built for different
offload devices using different programming models (CUDA, HIP, OpenMP target
offload). Currently supported modes are:

• CUDA: Flang runtime may be built as a fatlib for the host and a set of CUDA architectures specified during the build. The packaging of the device code is done by the CUDA toolchain and may differ from toolchan to toolchain.
• OpenMP offload:
  • host_device mode: Flang runtime may be built as a fatlib for the host and a set of OpenMP offload architectures. The packaging of the device code is done by the OpenMP offload compiler and may differ from compiler to compiler.

OpenMP offload 'nohost' mode is a TODO to match the build setup
of libomptarget/DeviceRTL. Flang runtime will be built as LLVM Bitcode
library using Clang/LLVM toolchain. The host part of the library
will be "empty", so there will be two distributable object: the host
Flang runtime and dummy host library with device Flang runtime pieces
packaged using clang-offload-packager and clang.

In all supported modes, enabling parts of Flang runtime for the device
compilation can be done iteratively to make the patches observable.
Note that at any point in time the resulting library may have unresolved
references to not yet enabled parts of Flang runtime.

Example cmake/make commands for building with Clang for NVPTX target:
cmake \
-DFLANG_EXPERIMENTAL_CUDA_RUNTIME=ON \
-DCMAKE_CUDA_ARCHITECTURES=80 \
-DCMAKE_C_COMPILER=/clang_nvptx/bin/clang \
-DCMAKE_CXX_COMPILER=/clang_nvptx/bin/clang++ \
-DCMAKE_CUDA_COMPILER=/clang_nvptx/bin/clang \
/llvm-project/flang/runtime/
make -j FortranRuntime

Example cmake/make commands for building with Clang OpenMP offload:
cmake \
-DFLANG_EXPERIMENTAL_OMP_OFFLOAD_BUILD="host_device" \
-DCMAKE_C_COMPILER=clang \
-DCMAKE_CXX_COMPILER=clang++ \
-DFLANG_OMP_DEVICE_ARCHITECTURES="sm_80" \
../flang/runtime/
make -j FortranRuntime

[tschuett]

@jhuber6

[ronlieb]

@JonChesterfield to review

[JonChesterfield]

Tagging a few more potential reviewers. Thanks for posting this.

Is the current plan for amdgpu to compile this as hip?

Long term are there plans to move.off cuda as the source language, e.g. to freestanding C++ or the device-only variant of openmp which essentially serves as a C++ dialect with GPU extensions for conjuring specific IR?

[jhuber6]

As mentioned in the RFC, I think the easiest way to accomplish this in would be to use the "new driver" that's currently opt-in for CUDA. I would also recommend using OpenMP instead if we want this library to be generic, but we may also be able to have a separate build using HIP. There's some documentation on the compilation pipeline for offloading at https://clang.llvm.org/docs/OffloadingDesign.html and a talk at https://www.youtube.com/watch?v=4NnzymmQe7k. Right now to support multiple GPUs we simply provide redundant builds. It's inefficient space-wise but it's the simplest solution to cross-architecture compatibility without relying on mandatory LLVM passes or config libraries. Creating a library that can be linked with the proposed Fortran pipeline would work as follows:

clang++ -x cu cuda.cpp --offload-arch=sm_70,sm_80 -foffload-lto -fvisibility=hidden --offload-new-driver -fgpu-rdc -c
llvm-ar rcs libcuda.a cuda.o

This will create a fatbinary containing code for sm_70 and sm_80 similar to CUDA's support. We use -foffload-lto to improve performance, -fvisibility=hidden allows LTO to optimize out unused symbols. We should then be able to link this libcuda.a with the Fortran application the user is trying to compile. The only change needed here is a special option to make the kind metadata in the fat binary to be empty so we don't try to build code for the CUDA runtime when we link it.

[klausler]

flang/include/flang/Runtime/entry-names.h

30

As I mentioned, "decl" seems weird to me when applied to both declarations and definitions.

I suggest "RTENTRY" or "RTAPI", but please feel free to retain RTDECL if you think it is best.

flang/runtime/device-tools.h

15 ↗

(On Diff #524538)

file*

25 ↗

(On Diff #524538)

This could be called "fill_n" in the Fortran::runtime namespace, and defined via "using std::fill_n;" by default.

[vzakhari]

Thank you all for the reviews and the comments on discourse. I will upload the updated changes shortly. Please let me know if the new approach could work for all of us.

flang/include/flang/Runtime/entry-names.h

30

Thank you for the review, Peter!

I decided to go with two macros: RTDECL and RTDEF. Having the two does not make much sense in the context of this patch, but it may help with other applications. For example, I wanted to experiment with Clang cpu dispatch feature for the Flang runtime exported entry points. With the two macros I could expand RTDECL into __attribute__((cpu_dispatch(cpu1,cpu2,generic))) and RTDEF into __attribute__((cpu_specific(cpu1,cpu2,generic))).

[tschuett]

There will be instructions on the Flang webpage?

[vzakhari]

In D151173#4391608, @tschuett wrote:

There will be instructions on the Flang webpage?

Yes, sure, I can add the build instructions for Flang runtime to https://github.com/llvm/llvm-project/blob/main/flang/docs/GettingStarted.md

[jdoerfert]

Do we really need to support "old" CUDA style offloading/linking? If so, we also need HIP, SYCL, ...
I would suggest only to support the "new" offload linker. That's where Clang (driver) is headed anyway, I think.

If we would do that, we could also drop the CUDA stuff, and benefit from fun things like "auto declare target" for all the internal functions we use. All we would really annotate as "__device__/declare target" is the interface. And we could properly internalize the rest as we do with the OpenMP DeviceRTL.

flang/include/flang/ISO_Fortran_binding.h

129

The fact that we now have some but not all members available on the device is not ideal, IMHO. I would have hoped the opposing approach would at least be tried; compile the entire thing for the device, opt-out where necessary.

flang/runtime/terminator.h

39

FWIW, we should just support them (by defining our own "ABI", one for all GPU archs). Basically, the "printf" special handling generalized.

[vzakhari]

In D151173#4391902, @jdoerfert wrote:

Do we really need to support "old" CUDA style offloading/linking?

Yes, I would like to be able to build Flang runtime for NVIDIA GPUs so that the resulting format is compatible with nvfortran driver, i.e. the Flang runtime library can be linked using nvlink.

If so, we also need HIP, SYCL, ...

I added the new macros such that they can be used for HIP and SYCL as well. E.g. for HIP setting RT_API_ATTRS to __host__ __device__ should work as well as for CUDA; for SYCL, setting SYCL_EXTERNAL in RTDEF should probably be enough. The CMake changes will probably differ depending on the used model, but at least the source code will be somewhat uniform.

I would suggest only to support the "new" offload linker. That's where Clang (driver) is headed anyway, I think.

If we would do that, we could also drop the CUDA stuff, and benefit from fun things like "auto declare target" for all the internal functions we use. All we would really annotate as "__device__/declare target" is the interface.

Actually, the "auto declare target" is already used in this patch. "declare target" is only in effect for the definitions (see RT_EXT_API_GROUP_BEGIN/END usage in transformational.cpp).

And we could properly internalize the rest as we do with the OpenMP DeviceRTL.

I am not sure what you mean by this, could you please clarify what you by "internalize"?

[vzakhari]

flang/include/flang/ISO_Fortran_binding.h

129

I was trying to minimize the changes for the initial patch. Going forward, I agree that enabling all class methods or even the whole single file makes more sense.

[vzakhari]

Updated build doc and rebased.

[vzakhari]

Reverted accidentally removed __SIZEOF_INT128__ guard. This should fix the Windows build.

[vzakhari]

Rebase before merging.

flang/docs/GettingStarted.md

 ```
 
 To run the flang tests on this build, execute the command in the "flang/build"
 directory:
 ```bash
 ninja check-flang
 ```
 
+### Building flang runtime for accelerators
+Flang runtime can be built for accelerators in experimental mode, i.e.
+complete enabling is WIP.  CUDA and OpenMP target offload builds
+are currently supported.
+
+#### Building out-of-tree
+
+##### CUDA build
+Clang with NVPTX backend and NVCC compilers are supported.
+
+```bash
+cd llvm-project/flang
+mkdir -rf build_flang_runtime
+mkdir build_flang_runtime
+cd build_flang_runtime
+
+cmake \
+  -DFLANG_EXPERIMENTAL_CUDA_RUNTIME=ON \
+  -DCMAKE_CUDA_ARCHITECTURES=80 \
+  -DCMAKE_C_COMPILER=clang \
+  -DCMAKE_CXX_COMPILER=clang++ \
+  -DCMAKE_CUDA_COMPILER=clang \
+  ../runtime/
+make -j FortranRuntime
+```
+
+```bash
+cd llvm-project/flang
+mkdir -rf build_flang_runtime
+mkdir build_flang_runtime
+cd build_flang_runtime
+
+cmake \
+  -DFLANG_EXPERIMENTAL_CUDA_RUNTIME=ON \
+  -DCMAKE_CUDA_ARCHITECTURES=80 \
+  -DCMAKE_C_COMPILER=clang \
+  -DCMAKE_CXX_COMPILER=clang++ \
+  -DCMAKE_CUDA_COMPILER=nvcc \
+  ../runtime/
+make -j FortranRuntime
+```
+
+The result of the build is a "fat" library with the host and device
+code.  Note that the packaging of the libraries is different
+between [Clang](https://clang.llvm.org/docs/OffloadingDesign.html#linking-target-device-code) and NVCC, so the library must be linked using
+compatible compiler drivers.
+
+##### OpenMP target offload build
+Only Clang compiler is currently supported.
+
+```
+cd llvm-project/flang
+mkdir -rf build_flang_runtime
+mkdir build_flang_runtime
+cd build_flang_runtime
+
+cmake \
+  -DFLANG_EXPERIMENTAL_OMP_OFFLOAD_BUILD="host_device" \
+  -DCMAKE_C_COMPILER=clang \
+  -DCMAKE_CXX_COMPILER=clang++ \
+  -DFLANG_OMP_DEVICE_ARCHITECTURES="all" \
+  ../runtime/
+make -j FortranRuntime
+```
+
+The result of the build is a "device-only" library, i.e. the host
+part of the library is just a container for the device code.
+The resulting library may be linked to user programs using
+Clang-like device linking pipeline.
+
 ## Supported C++ compilers
 
 Flang is written in C++17.
 
 The code has been compiled and tested with GCC versions from 7.2.0 to 9.3.0.
 
 The code has been compiled and tested with clang version 7.0, 8.0, 9.0 and 10.0
 using either GNU's libstdc++ or LLVM's libc++.

flang/include/flang/ISO_Fortran_binding.h

 #include <stddef.h>
 
 /* Standard interface to Fortran from C and C++.
  * These interfaces are named in subclause 18.5 of the Fortran 2018
  * standard, with most of the actual details being left to the
  * implementation.
  */
 
+#include "Runtime/api-attrs.h"
+
 #ifdef __cplusplus
 namespace Fortran {
 namespace ISO {
 inline namespace Fortran_2018 {
 #endif
 
 /* 18.5.4 */
 #define CFI_VERSION 20180515

 
 #ifdef __cplusplus
 namespace cfi_internal {
 // C++ does not support flexible array.
 // The below structure emulates a flexible array. This structure does not take
 // care of getting the memory storage. Note that it already contains one element
 // because a struct cannot be empty.
 template <typename T> struct FlexibleArray : T {
-  T &operator[](int index) { return *(this + index); }
-  const T &operator[](int index) const { return *(this + index); }
+  RT_API_ATTRS T &operator[](int index) { return *(this + index); }
+  const RT_API_ATTRS T &operator[](int index) const { return *(this + index); }
   operator T *() { return this; }
   operator const T *() const { return this; }
 };
 } // namespace cfi_internal
 #endif
 
 /* 18.5.3 generic data descriptor */
 typedef struct CFI_cdesc_t {
   /* These three members must appear first, in exactly this order. */
   void *base_addr;

 
 /* 18.5.5 procedural interfaces*/
 #ifdef __cplusplus
 extern "C" {
 #endif
 void *CFI_address(const CFI_cdesc_t *, const CFI_index_t subscripts[]);
 int CFI_allocate(CFI_cdesc_t *, const CFI_index_t lower_bounds[],
     const CFI_index_t upper_bounds[], size_t elem_len);
-int CFI_deallocate(CFI_cdesc_t *);
+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[]);
 int CFI_is_contiguous(const CFI_cdesc_t *);
-int CFI_section(CFI_cdesc_t *, const CFI_cdesc_t *source,
+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[],
     const CFI_index_t strides[]);
 int CFI_select_part(CFI_cdesc_t *, const CFI_cdesc_t *source,
     size_t displacement, size_t elem_len);
 int CFI_setpointer(
     CFI_cdesc_t *, const CFI_cdesc_t *source, const CFI_index_t lower_bounds[]);
 #ifdef __cplusplus
 } // extern "C"
 } // inline namespace Fortran_2018
 }
 }
 #endif
 
 #endif /* CFI_ISO_FORTRAN_BINDING_H_ */

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

+/*===-- include/flang/Runtime/api-attrs.h ---------------------------*- C -*-=//
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===------------------------------------------------------------------------===
+ */
+
+/*
+ * The file defines a set macros that can be used to apply
+ * different attributes/pragmas to functions/variables
+ * declared/defined/used in Flang runtime library.
+ */
+
+#ifndef FORTRAN_RUNTIME_API_ATTRS_H_
+#define FORTRAN_RUNTIME_API_ATTRS_H_
+
+/*
+ * RT_EXT_API_GROUP_BEGIN/END pair is placed around definitions
+ * of functions exported by Flang runtime library. They are the entry
+ * points that are referenced in the Flang generated code.
+ * The macros may be expanded into any construct that is valid to appear
+ * at C++ module scope.
+ */
+#ifndef RT_EXT_API_GROUP_BEGIN
+#if defined(OMP_NOHOST_BUILD)
+#define RT_EXT_API_GROUP_BEGIN \
+  _Pragma("omp begin declare target device_type(nohost)")
+#elif defined(OMP_OFFLOAD_BUILD)
+#define RT_EXT_API_GROUP_BEGIN _Pragma("omp declare target")
+#else
+#define RT_EXT_API_GROUP_BEGIN
+#endif
+#endif /* !defined(RT_EXT_API_GROUP_BEGIN) */
+
+#ifndef RT_EXT_API_GROUP_END
+#if defined(OMP_NOHOST_BUILD) || defined(OMP_OFFLOAD_BUILD)
+#define RT_EXT_API_GROUP_END _Pragma("omp end declare target")
+#else
+#define RT_EXT_API_GROUP_END
+#endif
+#endif /* !defined(RT_EXT_API_GROUP_END) */
+
+/*
+ * RT_VAR_GROUP_BEGIN/END pair is placed around definitions
+ * of module scope variables referenced by Flang runtime (directly
+ * or indirectly).
+ * The macros may be expanded into any construct that is valid to appear
+ * at C++ module scope.
+ */
+#ifndef RT_VAR_GROUP_BEGIN
+#define RT_VAR_GROUP_BEGIN RT_EXT_API_GROUP_BEGIN
+#endif /* !defined(RT_VAR_GROUP_BEGIN) */
+
+#ifndef RT_VAR_GROUP_END
+#define RT_VAR_GROUP_END RT_EXT_API_GROUP_END
+#endif /* !defined(RT_VAR_GROUP_END) */
+
+/*
+ * Each non-exported function used by Flang runtime (e.g. via
+ * calling it or taking its address, etc.) is marked with
+ * RT_API_ATTRS. The macros is placed at both declaration and
+ * definition of such a function.
+ * The macros may be expanded into a construct that is valid
+ * to appear as part of a C++ decl-specifier.
+ */
+#ifndef RT_API_ATTRS
+#if defined(__CUDACC__) || defined(__CUDA__)
+#define RT_API_ATTRS __host__ __device__
+#else
+#define RT_API_ATTRS
+#endif
+#endif /* !defined(RT_API_ATTRS) */
+
+/*
+ * Each const/constexpr module scope variable referenced by Flang runtime
+ * (directly or indirectly) is marked with RT_CONST_VAR_ATTRS.
+ * The macros is placed at both declaration and definition of such a variable.
+ * The macros may be expanded into a construct that is valid
+ * to appear as part of a C++ decl-specifier.
+ */
+#ifndef RT_CONST_VAR_ATTRS
+#if defined(__CUDACC__) || defined(__CUDA__)
+#define RT_CONST_VAR_ATTRS __constant__
+#else
+#define RT_CONST_VAR_ATTRS
+#endif
+#endif /* !defined(RT_CONST_VAR_ATTRS) */
+
+#endif /* !FORTRAN_RUNTIME_API_ATTRS_H_ */

flang/include/flang/Runtime/descriptor.h

 using TypeParameterValue = std::int64_t;
 class DerivedType;
 } // namespace Fortran::runtime::typeInfo
 
 namespace Fortran::runtime {
 
 using SubscriptValue = ISO::CFI_index_t;
 
-static constexpr int maxRank{CFI_MAX_RANK};
+RT_VAR_GROUP_BEGIN
+static constexpr RT_CONST_VAR_ATTRS int maxRank{CFI_MAX_RANK};
+RT_VAR_GROUP_END
 
 // A C++ view of the sole interoperable standard descriptor (ISO::CFI_cdesc_t)
 // and its type and per-dimension information.
 
 class Dimension {
 public:
-  SubscriptValue LowerBound() const { return raw_.lower_bound; }
-  SubscriptValue Extent() const { return raw_.extent; }
-  SubscriptValue UpperBound() const { return LowerBound() + Extent() - 1; }
-  SubscriptValue ByteStride() const { return raw_.sm; }
+  RT_API_ATTRS SubscriptValue LowerBound() const { return raw_.lower_bound; }
+  RT_API_ATTRS SubscriptValue Extent() const { return raw_.extent; }
+  RT_API_ATTRS SubscriptValue UpperBound() const {
+    return LowerBound() + Extent() - 1;
+  }
+  RT_API_ATTRS SubscriptValue ByteStride() const { return raw_.sm; }
 
-  Dimension &SetBounds(SubscriptValue lower, SubscriptValue upper) {
+  RT_API_ATTRS Dimension &SetBounds(
+      SubscriptValue lower, SubscriptValue upper) {
     if (upper >= lower) {
       raw_.lower_bound = lower;
       raw_.extent = upper - lower + 1;
     } else {
       raw_.lower_bound = 1;
       raw_.extent = 0;
     }
     return *this;
   }
   // Do not use this API to cause the LB of an empty dimension
   // to be anything other than 1.  Use SetBounds() instead if you can.
   Dimension &SetLowerBound(SubscriptValue lower) {
     raw_.lower_bound = lower;
     return *this;
   }
   Dimension &SetUpperBound(SubscriptValue upper) {
     auto lower{raw_.lower_bound};
     raw_.extent = upper >= lower ? upper - lower + 1 : 0;
     return *this;
   }
   Dimension &SetExtent(SubscriptValue extent) {
     raw_.extent = extent;
     return *this;
   }
-  Dimension &SetByteStride(SubscriptValue bytes) {
+  RT_API_ATTRS Dimension &SetByteStride(SubscriptValue bytes) {
     raw_.sm = bytes;
     return *this;
   }
 
 private:
   ISO::CFI_dim_t raw_;
 };
 
 // The storage for this object follows the last used dim[] entry in a
 // Descriptor (CFI_cdesc_t) generic descriptor.  Space matters here, since
 // descriptors serve as POINTER and ALLOCATABLE components of derived type
 // instances.  The presence of this structure is implied by the flag
 // CFI_cdesc_t.f18Addendum, and the number of elements in the len_[]
 // array is determined by derivedType_->LenParameters().
 class DescriptorAddendum {
 public:
-  explicit DescriptorAddendum(const typeInfo::DerivedType *dt = nullptr)
+  explicit RT_API_ATTRS DescriptorAddendum(
+      const typeInfo::DerivedType *dt = nullptr)
       : derivedType_{dt} {}
-  DescriptorAddendum &operator=(const DescriptorAddendum &);
+  RT_API_ATTRS DescriptorAddendum &operator=(const DescriptorAddendum &);
 
-  const typeInfo::DerivedType *derivedType() const { return derivedType_; }
-  DescriptorAddendum &set_derivedType(const typeInfo::DerivedType *dt) {
+  const RT_API_ATTRS typeInfo::DerivedType *derivedType() const {
+    return derivedType_;
+  }
+  RT_API_ATTRS DescriptorAddendum &set_derivedType(
+      const typeInfo::DerivedType *dt) {
     derivedType_ = dt;
     return *this;
   }
 
-  std::size_t LenParameters() const;
+  RT_API_ATTRS std::size_t LenParameters() const;
 
-  typeInfo::TypeParameterValue LenParameterValue(int which) const {
+  RT_API_ATTRS typeInfo::TypeParameterValue LenParameterValue(int which) const {
     return len_[which];
   }
-  static constexpr std::size_t SizeInBytes(int lenParameters) {
+  static constexpr RT_API_ATTRS std::size_t SizeInBytes(int lenParameters) {
     // TODO: Don't waste that last word if lenParameters == 0
     return sizeof(DescriptorAddendum) +
         std::max(lenParameters - 1, 0) * sizeof(typeInfo::TypeParameterValue);
   }
-  std::size_t SizeInBytes() const;
+  RT_API_ATTRS std::size_t SizeInBytes() const;
 
-  void SetLenParameterValue(int which, typeInfo::TypeParameterValue x) {
+  RT_API_ATTRS void SetLenParameterValue(
+      int which, typeInfo::TypeParameterValue x) {
     len_[which] = x;
   }
 
   void Dump(FILE * = stdout) const;
 
 private:
   const typeInfo::DerivedType *derivedType_;
   typeInfo::TypeParameterValue len_[1]; // must be the last component
public:
   // the size is going to be correct only by accident, since the true size of
   // a descriptor depends on the number of its dimensions and the presence and
   // size of an addendum, which depends on the type of the data.
   // Use the class template StaticDescriptor (below) to declare a descriptor
   // whose type and rank are fixed and known at compilation time.  Use the
   // Create() static member functions otherwise to dynamically allocate a
   // descriptor.
 
-  Descriptor(const Descriptor &);
-  Descriptor &operator=(const Descriptor &);
+  RT_API_ATTRS Descriptor(const Descriptor &);
+  RT_API_ATTRS Descriptor &operator=(const Descriptor &);
 
   // Returns the number of bytes occupied by an element of the given
   // category and kind including any alignment padding required
   // between adjacent elements.
-  static std::size_t BytesFor(TypeCategory category, int kind);
+  static RT_API_ATTRS std::size_t BytesFor(TypeCategory category, int kind);
 
-  void Establish(TypeCode t, std::size_t elementBytes, void *p = nullptr,
-      int rank = maxRank, const SubscriptValue *extent = nullptr,
-      ISO::CFI_attribute_t attribute = CFI_attribute_other,
-      bool addendum = false);
-  void Establish(TypeCategory, int kind, void *p = nullptr, int rank = maxRank,
+  RT_API_ATTRS void Establish(TypeCode t, std::size_t elementBytes,
+      void *p = nullptr, int rank = maxRank,
       const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other,
       bool addendum = false);
-  void Establish(int characterKind, std::size_t characters, void *p = nullptr,
+  RT_API_ATTRS void Establish(TypeCategory, int kind, void *p = nullptr,
       int rank = maxRank, const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other,
       bool addendum = false);
-  void Establish(const typeInfo::DerivedType &dt, void *p = nullptr,
-      int rank = maxRank, const SubscriptValue *extent = nullptr,
+  RT_API_ATTRS void Establish(int characterKind, std::size_t characters,
+      void *p = nullptr, int rank = maxRank,
+      const SubscriptValue *extent = nullptr,
+      ISO::CFI_attribute_t attribute = CFI_attribute_other,
+      bool addendum = false);
+  RT_API_ATTRS void Establish(const typeInfo::DerivedType &dt,
+      void *p = nullptr, int rank = maxRank,
+      const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other);
 
+  // CUDA_TODO: Clang does not support unique_ptr on device.
   static OwningPtr<Descriptor> Create(TypeCode t, std::size_t elementBytes,
       void *p = nullptr, int rank = maxRank,
       const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other,
       int derivedTypeLenParameters = 0);
   static OwningPtr<Descriptor> Create(TypeCategory, int kind, void *p = nullptr,
       int rank = maxRank, const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other);
   static OwningPtr<Descriptor> Create(int characterKind,
       SubscriptValue characters, void *p = nullptr, int rank = maxRank,
       const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other);
   static OwningPtr<Descriptor> Create(const typeInfo::DerivedType &dt,
       void *p = nullptr, int rank = maxRank,
       const SubscriptValue *extent = nullptr,
       ISO::CFI_attribute_t attribute = CFI_attribute_other);
 
-  ISO::CFI_cdesc_t &raw() { return raw_; }
-  const ISO::CFI_cdesc_t &raw() const { return raw_; }
-  std::size_t ElementBytes() const { return raw_.elem_len; }
-  int rank() const { return raw_.rank; }
-  TypeCode type() const { return TypeCode{raw_.type}; }
+  RT_API_ATTRS ISO::CFI_cdesc_t &raw() { return raw_; }
+  const RT_API_ATTRS ISO::CFI_cdesc_t &raw() const { return raw_; }
+  RT_API_ATTRS std::size_t ElementBytes() const { return raw_.elem_len; }
+  RT_API_ATTRS int rank() const { return raw_.rank; }
+  RT_API_ATTRS TypeCode type() const { return TypeCode{raw_.type}; }
 
-  Descriptor &set_base_addr(void *p) {
+  RT_API_ATTRS Descriptor &set_base_addr(void *p) {
     raw_.base_addr = p;
     return *this;
   }
 
-  bool IsPointer() const { return raw_.attribute == CFI_attribute_pointer; }
-  bool IsAllocatable() const {
+  RT_API_ATTRS bool IsPointer() const {
+    return raw_.attribute == CFI_attribute_pointer;
+  }
+  RT_API_ATTRS bool IsAllocatable() const {
     return raw_.attribute == CFI_attribute_allocatable;
   }
-  bool IsAllocated() const { return raw_.base_addr != nullptr; }
+  RT_API_ATTRS bool IsAllocated() const { return raw_.base_addr != nullptr; }
 
-  Dimension &GetDimension(int dim) {
+  RT_API_ATTRS Dimension &GetDimension(int dim) {
     return *reinterpret_cast<Dimension *>(&raw_.dim[dim]);
   }
-  const Dimension &GetDimension(int dim) const {
+  const RT_API_ATTRS Dimension &GetDimension(int dim) const {
     return *reinterpret_cast<const Dimension *>(&raw_.dim[dim]);
   }
 
-  std::size_t SubscriptByteOffset(
+  RT_API_ATTRS std::size_t SubscriptByteOffset(
       int dim, SubscriptValue subscriptValue) const {
     const Dimension &dimension{GetDimension(dim)};
     return (subscriptValue - dimension.LowerBound()) * dimension.ByteStride();
   }
 
-  std::size_t SubscriptsToByteOffset(const SubscriptValue subscript[]) const {
+  RT_API_ATTRS std::size_t SubscriptsToByteOffset(
+      const SubscriptValue subscript[]) const {
     std::size_t offset{0};
     for (int j{0}; j < raw_.rank; ++j) {
       offset += SubscriptByteOffset(j, subscript[j]);
     }
     return offset;
   }
 
-  template <typename A = char> A *OffsetElement(std::size_t offset = 0) const {
+  template <typename A = char>
+  RT_API_ATTRS A *OffsetElement(std::size_t offset = 0) const {
     return reinterpret_cast<A *>(
         reinterpret_cast<char *>(raw_.base_addr) + offset);
   }
 
-  template <typename A> A *Element(const SubscriptValue subscript[]) const {
+  template <typename A>
+  RT_API_ATTRS A *Element(const SubscriptValue subscript[]) const {
     return OffsetElement<A>(SubscriptsToByteOffset(subscript));
   }
 
-  template <typename A> A *ZeroBasedIndexedElement(std::size_t n) const {
+  template <typename A>
+  RT_API_ATTRS A *ZeroBasedIndexedElement(std::size_t n) const {
     SubscriptValue at[maxRank];
     if (SubscriptsForZeroBasedElementNumber(at, n)) {
       return Element<A>(at);
     }
     return nullptr;
   }
 
-  int GetLowerBounds(SubscriptValue subscript[]) const {
+  RT_API_ATTRS int GetLowerBounds(SubscriptValue subscript[]) const {
     for (int j{0}; j < raw_.rank; ++j) {
       subscript[j] = GetDimension(j).LowerBound();
     }
     return raw_.rank;
   }
 
-  int GetShape(SubscriptValue subscript[]) const {
+  RT_API_ATTRS int GetShape(SubscriptValue subscript[]) const {
     for (int j{0}; j < raw_.rank; ++j) {
       subscript[j] = GetDimension(j).Extent();
     }
     return raw_.rank;
   }
 
   // When the passed subscript vector contains the last (or first)
   // subscripts of the array, these wrap the subscripts around to
   // their first (or last) values and return false.
-  bool IncrementSubscripts(
+  RT_API_ATTRS bool IncrementSubscripts(
       SubscriptValue subscript[], const int *permutation = nullptr) const {
     for (int j{0}; j < raw_.rank; ++j) {
       int k{permutation ? permutation[j] : j};
       const Dimension &dim{GetDimension(k)};
       if (subscript[k]++ < dim.UpperBound()) {
         return true;
       }
       subscript[k] = dim.LowerBound();
     }
     return false;
   }
 
-  bool DecrementSubscripts(
+  RT_API_ATTRS bool DecrementSubscripts(
       SubscriptValue[], const int *permutation = nullptr) const;
 
   // False when out of range.
-  bool SubscriptsForZeroBasedElementNumber(SubscriptValue subscript[],
-      std::size_t elementNumber, const int *permutation = nullptr) const {
+  RT_API_ATTRS bool SubscriptsForZeroBasedElementNumber(
+      SubscriptValue subscript[], std::size_t elementNumber,
+      const int *permutation = nullptr) const {
     if (raw_.rank == 0) {
       return elementNumber == 0;
     }
     std::size_t dimCoefficient[maxRank];
     int k0{permutation ? permutation[0] : 0};
     dimCoefficient[0] = 1;
     auto coefficient{static_cast<std::size_t>(GetDimension(k0).Extent())};
     for (int j{1}; j < raw_.rank; ++j) {
for (int j{raw_.rank - 1}; j > 0; --j) {
       std::size_t quotient{elementNumber / dimCoefficient[j]};
       subscript[k] = quotient + dim.LowerBound();
       elementNumber -= quotient * dimCoefficient[j];
     }
     subscript[k0] = elementNumber + GetDimension(k0).LowerBound();
     return true;
   }
 
-  std::size_t ZeroBasedElementNumber(
+  RT_API_ATTRS std::size_t ZeroBasedElementNumber(
       const SubscriptValue *, const int *permutation = nullptr) const;
 
-  DescriptorAddendum *Addendum() {
+  RT_API_ATTRS DescriptorAddendum *Addendum() {
     if (raw_.f18Addendum != 0) {
       return reinterpret_cast<DescriptorAddendum *>(&GetDimension(rank()));
     } else {
       return nullptr;
     }
   }
-  const DescriptorAddendum *Addendum() const {
+  const RT_API_ATTRS DescriptorAddendum *Addendum() const {
     if (raw_.f18Addendum != 0) {
       return reinterpret_cast<const DescriptorAddendum *>(
           &GetDimension(rank()));
     } else {
       return nullptr;
     }
   }
 
   // Returns size in bytes of the descriptor (not the data)
-  static constexpr std::size_t SizeInBytes(
+  static constexpr RT_API_ATTRS std::size_t SizeInBytes(
       int rank, bool addendum = false, int lengthTypeParameters = 0) {
     std::size_t bytes{sizeof(Descriptor) - sizeof(Dimension)};
     bytes += rank * sizeof(Dimension);
     if (addendum || lengthTypeParameters > 0) {
       bytes += DescriptorAddendum::SizeInBytes(lengthTypeParameters);
     }
     return bytes;
   }
 
-  std::size_t SizeInBytes() const;
+  RT_API_ATTRS std::size_t SizeInBytes() const;
 
-  std::size_t Elements() const;
+  RT_API_ATTRS std::size_t Elements() const;
 
   // Allocate() assumes Elements() and ElementBytes() work;
   // define the extents of the dimensions and the element length
   // before calling.  It (re)computes the byte strides after
   // allocation.  Does not allocate automatic components or
   // perform default component initialization.
-  int Allocate();
+  RT_API_ATTRS int Allocate();
 
   // Deallocates storage; does not call FINAL subroutines or
   // deallocate allocatable/automatic components.
-  int Deallocate();
+  RT_API_ATTRS int Deallocate();
 
   // Deallocates storage, including allocatable and automatic
   // components.  Optionally invokes FINAL subroutines.
-  int Destroy(bool finalize = false, bool destroyPointers = false);
+  RT_API_ATTRS int Destroy(bool finalize = false, bool destroyPointers = false);
 
-  bool IsContiguous(int leadingDimensions = maxRank) const {
+  RT_API_ATTRS bool IsContiguous(int leadingDimensions = maxRank) const {
     auto bytes{static_cast<SubscriptValue>(ElementBytes())};
     if (leadingDimensions > raw_.rank) {
       leadingDimensions = raw_.rank;
     }
     for (int j{0}; j < leadingDimensions; ++j) {
       const Dimension &dim{GetDimension(j)};
       if (bytes != dim.ByteStride()) {
         return false;
       }
       bytes *= dim.Extent();
     }
     return true;
   }
 
   // Establishes a pointer to a section or element.
-  bool EstablishPointerSection(const Descriptor &source,
+  RT_API_ATTRS bool EstablishPointerSection(const Descriptor &source,
       const SubscriptValue *lower = nullptr,
       const SubscriptValue *upper = nullptr,
       const SubscriptValue *stride = nullptr);
 
-  void Check() const;
+  RT_API_ATTRS void Check() const;
 
   void Dump(FILE * = stdout) const;
 
 private:
   ISO::CFI_cdesc_t raw_;
 };
 static_assert(sizeof(Descriptor) == sizeof(ISO::CFI_cdesc_t));
 

 class alignas(Descriptor) StaticDescriptor {
 public:
   static constexpr int maxRank{MAX_RANK};
   static constexpr int maxLengthTypeParameters{MAX_LEN_PARMS};
   static constexpr bool hasAddendum{ADDENDUM || MAX_LEN_PARMS > 0};
   static constexpr std::size_t byteSize{
       Descriptor::SizeInBytes(maxRank, hasAddendum, maxLengthTypeParameters)};
 
-  Descriptor &descriptor() { return *reinterpret_cast<Descriptor *>(storage_); }
-  const Descriptor &descriptor() const {
+  RT_API_ATTRS Descriptor &descriptor() {
+    return *reinterpret_cast<Descriptor *>(storage_);
+  }
+  const RT_API_ATTRS Descriptor &descriptor() const {
     return *reinterpret_cast<const Descriptor *>(storage_);
   }
 
-  void Check() {
+  RT_API_ATTRS void Check() {
     assert(descriptor().rank() <= maxRank);
     assert(descriptor().SizeInBytes() <= byteSize);
     if (DescriptorAddendum * addendum{descriptor().Addendum()}) {
       assert(hasAddendum);
       assert(addendum->LenParameters() <= maxLengthTypeParameters);
     } else {
       assert(!hasAddendum);
       assert(maxLengthTypeParameters == 0);

flang/include/flang/Runtime/entry-names.h

  * library function or object with extra characters so that it
  * (a) is not in the user's name space,
  * (b) doesn't conflict with other libraries, and
  * (c) prevents incompatible versions of the runtime library from linking
  *
  * The value of REVISION should not be changed until/unless the API to the
  * runtime library must change in some way that breaks backward compatibility.
  */
+#ifndef FORTRAN_RUNTIME_ENTRY_NAMES_H
+#define FORTRAN_RUNTIME_ENTRY_NAMES_H
+
+#include "flang/Runtime/api-attrs.h"
+
 #ifndef RTNAME
 #define NAME_WITH_PREFIX_AND_REVISION(prefix, revision, name) \
   prefix##revision##name
 #define RTNAME(name) NAME_WITH_PREFIX_AND_REVISION(_Fortran, A, name)
 #endif
 
+#ifndef RTDECL
+#define RTDECL(name) RT_API_ATTRS RTNAME(name)
+#endif
+
+#ifndef RTDEF
+#define RTDEF(name) RT_API_ATTRS RTNAME(name)
+#endif
+
 #ifndef RTNAME_STRING
 #define RTNAME_STRINGIFY_(x) #x
 #define RTNAME_STRINGIFY(x) RTNAME_STRINGIFY_(x)
 #define RTNAME_STRING(name) RTNAME_STRINGIFY(RTNAME(name))
 #endif
+
+#endif /* !FORTRAN_RUNTIME_ENTRY_NAMES_H */

flang/include/flang/Runtime/float128.h

  * In order to check for libc++'s _LIBCPP_VERSION macro
  * we need to include at least one libc++ header file.
  */
 #include <cstddef>
 #endif
 
 #undef HAS_FLOAT128
 #if (defined(__FLOAT128__) || defined(__SIZEOF_FLOAT128__)) && \
-    !defined(_LIBCPP_VERSION)
+    !defined(_LIBCPP_VERSION) && !defined(__CUDA_ARCH__)
 /*
  * It may still be worth checking for compiler versions,
  * since earlier versions may define the macros above, but
  * still do not support __float128 fully.
  */
 #if __x86_64__
 #if __GNUC__ >= 7 || __clang_major__ >= 7
 #define HAS_FLOAT128 1
 #endif
 #elif defined __PPC__ && __GNUC__ >= 8
 #define HAS_FLOAT128 1
 #endif
 #endif /* (defined(__FLOAT128__) || defined(__SIZEOF_FLOAT128__)) && \
-          !defined(_LIBCPP_VERSION) */
+          !defined(_LIBCPP_VERSION)  && !defined(__CUDA_ARCH__) */
 
 #endif /* FORTRAN_RUNTIME_FLOAT128_H_ */

flang/include/flang/Runtime/transformational.h

 #include <cinttypes>
 
 namespace Fortran::runtime {
 
 class Descriptor;
 
 extern "C" {
 
-void RTNAME(Reshape)(Descriptor &result, const Descriptor &source,
+void RTDECL(Reshape)(Descriptor &result, const Descriptor &source,
     const Descriptor &shape, const Descriptor *pad = nullptr,
     const Descriptor *order = nullptr, const char *sourceFile = nullptr,
     int line = 0);
 
-void RTNAME(BesselJn_2)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselJn_2)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn2, float bn2_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJn_3)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselJn_3)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn2, float bn2_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJn_4)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselJn_4)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn2, float bn2_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJn_8)(Descriptor &result, int32_t n1, int32_t n2, double x,
+void RTDECL(BesselJn_8)(Descriptor &result, int32_t n1, int32_t n2, double x,
     double bn2, double bn2_1, const char *sourceFile = nullptr, int line = 0);
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselJn_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJn_10)(Descriptor &result, int32_t n1, int32_t n2,
     long double x, long double bn2, long double bn2_1,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselJn_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJn_16)(Descriptor &result, int32_t n1, int32_t n2,
     CppFloat128Type x, CppFloat128Type bn2, CppFloat128Type bn2_1,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
-void RTNAME(BesselJnX0_2)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_2)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJnX0_3)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_3)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselJnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselJnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselJnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselJnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
-void RTNAME(BesselYn_2)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselYn_2)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn1, float bn1_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYn_3)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselYn_3)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn1, float bn1_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYn_4)(Descriptor &result, int32_t n1, int32_t n2, float x,
+void RTDECL(BesselYn_4)(Descriptor &result, int32_t n1, int32_t n2, float x,
     float bn1, float bn1_1, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYn_8)(Descriptor &result, int32_t n1, int32_t n2, double x,
+void RTDECL(BesselYn_8)(Descriptor &result, int32_t n1, int32_t n2, double x,
     double bn1, double bn1_1, const char *sourceFile = nullptr, int line = 0);
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselYn_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYn_10)(Descriptor &result, int32_t n1, int32_t n2,
     long double x, long double bn1, long double bn1_1,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselYn_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYn_16)(Descriptor &result, int32_t n1, int32_t n2,
     CppFloat128Type x, CppFloat128Type bn1, CppFloat128Type bn1_1,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
-void RTNAME(BesselYnX0_2)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_2)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYnX0_3)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_3)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(BesselYnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselYnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselYnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(BesselYnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile = nullptr, int line = 0);
 #endif
 
-void RTNAME(Cshift)(Descriptor &result, const Descriptor &source,
+void RTDECL(Cshift)(Descriptor &result, const Descriptor &source,
     const Descriptor &shift, int dim = 1, const char *sourceFile = nullptr,
     int line = 0);
-void RTNAME(CshiftVector)(Descriptor &result, const Descriptor &source,
+void RTDECL(CshiftVector)(Descriptor &result, const Descriptor &source,
     std::int64_t shift, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(Eoshift)(Descriptor &result, const Descriptor &source,
+void RTDECL(Eoshift)(Descriptor &result, const Descriptor &source,
     const Descriptor &shift, const Descriptor *boundary = nullptr, int dim = 1,
     const char *sourceFile = nullptr, int line = 0);
-void RTNAME(EoshiftVector)(Descriptor &result, const Descriptor &source,
+void RTDECL(EoshiftVector)(Descriptor &result, const Descriptor &source,
     std::int64_t shift, const Descriptor *boundary = nullptr,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(Pack)(Descriptor &result, const Descriptor &source,
+void RTDECL(Pack)(Descriptor &result, const Descriptor &source,
     const Descriptor &mask, const Descriptor *vector = nullptr,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(Spread)(Descriptor &result, const Descriptor &source, int dim,
+void RTDECL(Spread)(Descriptor &result, const Descriptor &source, int dim,
     std::int64_t ncopies, const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(Transpose)(Descriptor &result, const Descriptor &matrix,
+void RTDECL(Transpose)(Descriptor &result, const Descriptor &matrix,
     const char *sourceFile = nullptr, int line = 0);
 
-void RTNAME(Unpack)(Descriptor &result, const Descriptor &vector,
+void RTDECL(Unpack)(Descriptor &result, const Descriptor &vector,
     const Descriptor &mask, const Descriptor &field,
     const char *sourceFile = nullptr, int line = 0);
 
 } // extern "C"
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_TRANSFORMATIONAL_H_

flang/include/flang/Runtime/type-code.h

 
 namespace Fortran::runtime {
 
 using common::TypeCategory;
 
 class TypeCode {
 public:
   TypeCode() {}
-  explicit TypeCode(ISO::CFI_type_t t) : raw_{t} {}
-  TypeCode(TypeCategory, int kind);
+  explicit RT_API_ATTRS TypeCode(ISO::CFI_type_t t) : raw_{t} {}
+  RT_API_ATTRS TypeCode(TypeCategory, int kind);
 
-  int raw() const { return raw_; }
+  RT_API_ATTRS int raw() const { return raw_; }
 
   constexpr bool IsValid() const {
     return raw_ >= CFI_type_signed_char && raw_ <= CFI_TYPE_LAST;
   }
   constexpr bool IsInteger() const {
     return raw_ >= CFI_type_signed_char && raw_ <= CFI_type_ptrdiff_t;
   }
   constexpr bool IsReal() const {
public:
   }
   constexpr bool IsLogical() const {
     return raw_ == CFI_type_Bool ||
         (raw_ >= CFI_type_int_least8_t && raw_ <= CFI_type_int_least64_t);
   }
   constexpr bool IsDerived() const { return raw_ == CFI_type_struct; }
   constexpr bool IsIntrinsic() const { return IsValid() && !IsDerived(); }
 
-  std::optional<std::pair<TypeCategory, int>> GetCategoryAndKind() const;
+  RT_API_ATTRS std::optional<std::pair<TypeCategory, int>>
+  GetCategoryAndKind() const;
 
-  bool operator==(const TypeCode &that) const { return raw_ == that.raw_; }
+  RT_API_ATTRS bool operator==(const TypeCode &that) const {
+    return raw_ == that.raw_;
+  }
   bool operator!=(const TypeCode &that) const { return raw_ != that.raw_; }
 
 private:
   ISO::CFI_type_t raw_{CFI_type_other};
 };
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_TYPE_CODE_H_

flang/runtime/CMakeLists.txt

 
 # Disable libstdc++/libc++ assertions, even in an LLVM_ENABLE_ASSERTIONS build,
 # to avoid an unwanted dependency on libstdc++/libc++.so.
 add_definitions(-U_GLIBCXX_ASSERTIONS)
 add_definitions(-U_LIBCPP_ENABLE_ASSERTIONS)
 
 add_subdirectory(FortranMain)
 
-add_flang_library(FortranRuntime
+set(sources
   ISO_Fortran_binding.cpp
   allocatable.cpp
   array-constructor.cpp
   assign.cpp
   buffer.cpp
   command.cpp
   complex-powi.cpp
   complex-reduction.c
set(sources
   time-intrinsic.cpp
   tools.cpp
   transformational.cpp
   type-code.cpp
   type-info.cpp
   unit.cpp
   unit-map.cpp
   utf.cpp
+)
+
+option(FLANG_EXPERIMENTAL_CUDA_RUNTIME
+  "Compile Fortran runtime as CUDA sources (experimental)" OFF
+  )
+
+# List of files that are buildable for all devices.
+set(supported_files
+  transformational.cpp
+  )
+
+if (FLANG_EXPERIMENTAL_CUDA_RUNTIME)
+  enable_language(CUDA)
+
+  # Add the unsupported files to LLVM_OPTIONAL_SOURCES.
+  set(todo_files ${sources})
+  list(REMOVE_ITEM todo_files ${supported_files})
+  list(APPEND LLVM_OPTIONAL_SOURCES ${todo_files})
+
+  # TODO: figure out how to make target property CUDA_SEPARABLE_COMPILATION
+  # work, and avoid setting CMAKE_CUDA_SEPARABLE_COMPILATION.
+  set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
+
+  # Treat all sources as CUDA files.
+  set(sources ${supported_files})
+  set_source_files_properties(${sources} PROPERTIES LANGUAGE CUDA)
+  if ("${CMAKE_CUDA_COMPILER_ID}" MATCHES "Clang")
+    # Allow varargs.
+    add_compile_options(-Xclang -fcuda-allow-variadic-functions)
+  endif()
+endif()
+
+set(FLANG_EXPERIMENTAL_OMP_OFFLOAD_BUILD "off" CACHE STRING
+  "Compile Fortran runtime as OpenMP target offload sources (experimental). Valid options are 'off', 'host_device', 'nohost'")
+
+set(FLANG_OMP_DEVICE_ARCHITECTURES "all" CACHE STRING
+  "List of OpenMP device architectures to be used to compile the Fortran runtime (e.g. 'gfx1103;sm_90')")
+
+if (NOT FLANG_EXPERIMENTAL_OMP_OFFLOAD_BUILD STREQUAL "off")
+  # 'host_device' build only works with Clang compiler currently.
+  # The build is done with the CMAKE_C/CXX_COMPILER, i.e. it does not use
+  # the in-tree built Clang. We may have a mode that would use the in-tree
+  # built Clang.
+  #
+  # 'nohost' is supposed to produce an LLVM Bitcode library,
+  # and it has to be done with a C/C++ compiler producing LLVM Bitcode
+  # compatible with the LLVM toolchain version distributed with the Flang
+  # compiler.
+  # In general, the in-tree built Clang should be used for 'nohost' build.
+  # Note that 'nohost' build does not produce the host version of Flang
+  # runtime library, so there will be two separate distributable objects.
+  # 'nohost' build is a TODO.
 
+  if (NOT FLANG_EXPERIMENTAL_OMP_OFFLOAD_BUILD STREQUAL "host_device")
+    message(FATAL_ERROR "Unsupported OpenMP offload build of Flang runtime")
+  endif()
+
+  # Add the unsupported files to LLVM_OPTIONAL_SOURCES.
+  set(todo_files ${sources})
+  list(REMOVE_ITEM todo_files ${supported_files})
+  list(APPEND LLVM_OPTIONAL_SOURCES ${todo_files})
+  set(sources ${supported_files})
+
+  if ("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang" AND
+      "${CMAKE_C_COMPILER_ID}" MATCHES "Clang")
+
+    set(all_amdgpu_architectures
+      "gfx700;gfx701;gfx801;gfx803;gfx900;gfx902;gfx906"
+      "gfx908;gfx90a;gfx90c;gfx940;gfx1010;gfx1030"
+      "gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036"
+      "gfx1100;gfx1101;gfx1102;gfx1103"
+      )
+    set(all_nvptx_architectures
+      "sm_35;sm_37;sm_50;sm_52;sm_53;sm_60;sm_61;sm_62"
+      "sm_70;sm_72;sm_75;sm_80;sm_86;sm_89;sm_90"
+      )
+    set(all_gpu_architectures
+      "${all_amdgpu_architectures};${all_nvptx_architectures}"
+      )
+    # TODO: support auto detection on the build system.
+    if (FLANG_OMP_DEVICE_ARCHITECTURES STREQUAL "all")
+      set(FLANG_OMP_DEVICE_ARCHITECTURES ${all_gpu_architectures})
+    endif()
+    list(REMOVE_DUPLICATES FLANG_OMP_DEVICE_ARCHITECTURES)
+
+    string(REPLACE ";" "," compile_for_architectures
+      "${FLANG_OMP_DEVICE_ARCHITECTURES}"
+      )
+
+    add_compile_options(-fopenmp -fvisibility=hidden -fopenmp-cuda-mode)
+    add_compile_options(--offload-arch=${compile_for_architectures})
+    # Force LTO for the device part.
+    add_compile_options(-foffload-lto)
+  else()
+    message(FATAL_ERROR
+      "Flang runtime build is not supported for these compilers:\n"
+      "CMAKE_CXX_COMPILER_ID: ${CMAKE_CXX_COMPILER_ID}\n"
+      "CMAKE_C_COMPILER_ID: ${CMAKE_C_COMPILER_ID}")
+  endif()
+
+  # Enable "declare target" in the source code.
+  add_compile_definitions(OMP_OFFLOAD_BUILD)
+endif()
+
+add_flang_library(FortranRuntime
+  ${sources}
   LINK_LIBS
   FortranDecimal
 
   INSTALL_WITH_TOOLCHAIN
 )

flang/runtime/copy.h

 #define FORTRAN_RUNTIME_COPY_H_
 
 #include "flang/Runtime/descriptor.h"
 
 namespace Fortran::runtime {
 
 // Assigns to uninitialized storage.
 // Duplicates allocatable & automatic components.
-void CopyElement(const Descriptor &to, const SubscriptValue toAt[],
+RT_API_ATTRS void CopyElement(const Descriptor &to, const SubscriptValue toAt[],
     const Descriptor &from, const SubscriptValue fromAt[], Terminator &);
 
 // Copies data from one allocated descriptor's array to another.
-void CopyArray(const Descriptor &to, const Descriptor &from, Terminator &);
+RT_API_ATTRS void CopyArray(
+    const Descriptor &to, const Descriptor &from, Terminator &);
 
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_COPY_H_

flang/runtime/freestanding-tools.h

+//===-- runtime/freestanding-tools.h ----------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef FORTRAN_RUNTIME_FREESTANDING_TOOLS_H_
+#define FORTRAN_RUNTIME_FREESTANDING_TOOLS_H_
+
+#include "flang/Runtime/api-attrs.h"
+#include <algorithm>
+
+// The file defines a set of utilities/classes that might be
+// used to get reduce the dependency on external libraries (e.g. libstdc++).
+
+#if !defined(STD_FILL_N_UNSUPPORTED) && \
+    (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__)
+#define STD_FILL_N_UNSUPPORTED 1
+#endif
+
+namespace Fortran::runtime {
+
+#if STD_FILL_N_UNSUPPORTED
+// Provides alternative implementation for std::fill_n(), if
+// it is not supported.
+template <typename A>
+static inline RT_API_ATTRS void fill_n(
+    A *start, std::size_t count, const A &value) {
+#if STD_FILL_N_UNSUPPORTED
+  for (std::size_t j{0}; j < count; ++j)
+    start[j] = value;
+#else
+  std::fill_n(start, count, value);
+#endif
+}
+#else // !STD_FILL_N_UNSUPPORTED
+using std::fill_n;
+#endif // !STD_FILL_N_UNSUPPORTED
+
+} // namespace Fortran::runtime
+#endif // FORTRAN_RUNTIME_FREESTANDING_TOOLS_H_

flang/runtime/terminator.h

 //===-- runtime/terminator.h ------------------------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 // Termination of the image
 
 #ifndef FORTRAN_RUNTIME_TERMINATOR_H_
 #define FORTRAN_RUNTIME_TERMINATOR_H_
 
+#include "flang/Runtime/api-attrs.h"
 #include <cstdarg>
 
 namespace Fortran::runtime {
 
 // A mixin class for statement-specific image error termination
 // for errors detected in the runtime library
 class Terminator {
 public:
   Terminator() {}
   Terminator(const Terminator &) = default;
-  explicit Terminator(const char *sourceFileName, int sourceLine = 0)
+  explicit RT_API_ATTRS Terminator(
+      const char *sourceFileName, int sourceLine = 0)
       : sourceFileName_{sourceFileName}, sourceLine_{sourceLine} {}
 
   const char *sourceFileName() const { return sourceFileName_; }
   int sourceLine() const { return sourceLine_; }
 
   void SetLocation(const char *sourceFileName = nullptr, int sourceLine = 0) {
     sourceFileName_ = sourceFileName;
     sourceLine_ = sourceLine;
   }
-  [[noreturn]] void Crash(const char *message, ...) const;
-  [[noreturn]] void CrashArgs(const char *message, va_list &) const;
-  [[noreturn]] void CheckFailed(
+
+  // CUDA_TODO: Clang for CUDA does not support varargs, though
+  // it compiles it with -fcuda-allow-variadic-functions.
+  // We can try to replace varargs functions with variadic templates.
+  [[noreturn]] RT_API_ATTRS void Crash(const char *message, ...) const;
+  [[noreturn]] RT_API_ATTRS void CrashArgs(
+      const char *message, va_list &) const;
+  [[noreturn]] RT_API_ATTRS void CheckFailed(
       const char *predicate, const char *file, int line) const;
-  [[noreturn]] void CheckFailed(const char *predicate) const;
+  [[noreturn]] RT_API_ATTRS void CheckFailed(const char *predicate) const;
 
   // For test harnessing - overrides CrashArgs().
   static void RegisterCrashHandler(void (*)(const char *sourceFile,
       int sourceLine, const char *message, va_list &ap));
 
 private:
   const char *sourceFileName_{nullptr};
   int sourceLine_{0};

flang/runtime/tools.h

 //===-- runtime/tools.h -----------------------------------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef FORTRAN_RUNTIME_TOOLS_H_
 #define FORTRAN_RUNTIME_TOOLS_H_
 
+#include "freestanding-tools.h"
 #include "terminator.h"
 #include "flang/Runtime/cpp-type.h"
 #include "flang/Runtime/descriptor.h"
 #include "flang/Runtime/memory.h"
 #include <cstring>
 #include <functional>
 #include <map>
 #include <type_traits>

 int IdentifyValue(
     const char *value, std::size_t length, const char *possibilities[]);
 
 // Truncates or pads as necessary
 void ToFortranDefaultCharacter(
     char *to, std::size_t toLength, const char *from);
 
 // Utility for dealing with elemental LOGICAL arguments
-inline bool IsLogicalElementTrue(
+inline RT_API_ATTRS bool IsLogicalElementTrue(
     const Descriptor &logical, const SubscriptValue at[]) {
   // A LOGICAL value is false if and only if all of its bytes are zero.
   const char *p{logical.Element<char>(at)};
   for (std::size_t j{logical.ElementBytes()}; j-- > 0; ++p) {
     if (*p) {
       return true;
     }
   }
   return false;
 }
 
 // Check array conformability; a scalar 'x' conforms.  Crashes on error.
-void CheckConformability(const Descriptor &to, const Descriptor &x,
+RT_API_ATTRS void CheckConformability(const Descriptor &to, const Descriptor &x,
     Terminator &, const char *funcName, const char *toName,
     const char *fromName);
 
 // Helper to store integer value in result[at].
 template <int KIND> struct StoreIntegerAt {
   void operator()(const Fortran::runtime::Descriptor &result, std::size_t at,
       std::int64_t value) const {
     *result.ZeroBasedIndexedElement<Fortran::runtime::CppTypeFor<
         Fortran::common::TypeCategory::Integer, KIND>>(at) = value;
   }
 };
 
 // Validate a KIND= argument
-void CheckIntegerKind(Terminator &, int kind, const char *intrinsic);
+RT_API_ATTRS void CheckIntegerKind(
+    Terminator &, int kind, const char *intrinsic);
 
 template <typename TO, typename FROM>
 inline void PutContiguousConverted(TO *to, FROM *from, std::size_t count) {
   while (count-- > 0) {
     *to++ = *from++;
   }
 }
 
-static inline std::int64_t GetInt64(
+static inline RT_API_ATTRS std::int64_t GetInt64(
     const char *p, std::size_t bytes, Terminator &terminator) {
   switch (bytes) {
   case 1:
     return *reinterpret_cast<const CppTypeFor<TypeCategory::Integer, 1> *>(p);
   case 2:
     return *reinterpret_cast<const CppTypeFor<TypeCategory::Integer, 2> *>(p);
   case 4:
     return *reinterpret_cast<const CppTypeFor<TypeCategory::Integer, 4> *>(p);
inline bool SetInteger(INT &x, int kind, std::int64_t value) {
   }
 }
 
 // Maps intrinsic runtime type category and kind values to the appropriate
 // instantiation of a function object template and calls it with the supplied
 // arguments.
 template <template <TypeCategory, int> class FUNC, typename RESULT,
     typename... A>
-inline RESULT ApplyType(
+inline RT_API_ATTRS RESULT ApplyType(
     TypeCategory cat, int kind, Terminator &terminator, A &&...x) {
   switch (cat) {
   case TypeCategory::Integer:
     switch (kind) {
     case 1:
       return FUNC<TypeCategory::Integer, 1>{}(std::forward<A>(x)...);
     case 2:
       return FUNC<TypeCategory::Integer, 2>{}(std::forward<A>(x)...);
default:
     terminator.Crash(
         "not yet implemented: type category(%d)", static_cast<int>(cat));
   }
 }
 
 // Maps a runtime INTEGER kind value to the appropriate instantiation of
 // a function object template and calls it with the supplied arguments.
 template <template <int KIND> class FUNC, typename RESULT, typename... A>
-inline RESULT ApplyIntegerKind(int kind, Terminator &terminator, A &&...x) {
+inline RT_API_ATTRS RESULT ApplyIntegerKind(
+    int kind, Terminator &terminator, A &&...x) {
   switch (kind) {
   case 1:
     return FUNC<1>{}(std::forward<A>(x)...);
   case 2:
     return FUNC<2>{}(std::forward<A>(x)...);
   case 4:
     return FUNC<4>{}(std::forward<A>(x)...);
   case 8:
     return FUNC<8>{}(std::forward<A>(x)...);
 #ifdef __SIZEOF_INT128__
   case 16:
     return FUNC<16>{}(std::forward<A>(x)...);
 #endif
   default:
     terminator.Crash("not yet implemented: INTEGER(KIND=%d)", kind);
   }
 }
 
 template <template <int KIND> class FUNC, typename RESULT, typename... A>
-inline RESULT ApplyFloatingPointKind(
+inline RT_API_ATTRS RESULT ApplyFloatingPointKind(
     int kind, Terminator &terminator, A &&...x) {
   switch (kind) {
 #if 0 // TODO: REAL/COMPLEX (2 & 3)
   case 2:
     return FUNC<2>{}(std::forward<A>(x)...);
   case 3:
     return FUNC<3>{}(std::forward<A>(x)...);
 #endif
if constexpr (HasCppTypeFor<TypeCategory::Real, 16>) {
       return FUNC<16>{}(std::forward<A>(x)...);
     }
     break;
   }
   terminator.Crash("not yet implemented: REAL/COMPLEX(KIND=%d)", kind);
 }
 
 template <template <int KIND> class FUNC, typename RESULT, typename... A>
-inline RESULT ApplyCharacterKind(int kind, Terminator &terminator, A &&...x) {
+inline RT_API_ATTRS RESULT ApplyCharacterKind(
+    int kind, Terminator &terminator, A &&...x) {
   switch (kind) {
   case 1:
     return FUNC<1>{}(std::forward<A>(x)...);
   case 2:
     return FUNC<2>{}(std::forward<A>(x)...);
   case 4:
     return FUNC<4>{}(std::forward<A>(x)...);
   default:
     terminator.Crash("not yet implemented: CHARACTER(KIND=%d)", kind);
   }
 }
 
 template <template <int KIND> class FUNC, typename RESULT, typename... A>
-inline RESULT ApplyLogicalKind(int kind, Terminator &terminator, A &&...x) {
+inline RT_API_ATTRS RESULT ApplyLogicalKind(
+    int kind, Terminator &terminator, A &&...x) {
   switch (kind) {
   case 1:
     return FUNC<1>{}(std::forward<A>(x)...);
   case 2:
     return FUNC<2>{}(std::forward<A>(x)...);
   case 4:
     return FUNC<4>{}(std::forward<A>(x)...);
   case 8:

flang/runtime/transformational.cpp

 // work with arbitrary lower bounds.  This may be technically an extension
 // of the standard but it more likely to conform with its intent.
 
 #include "flang/Runtime/transformational.h"
 #include "copy.h"
 #include "terminator.h"
 #include "tools.h"
 #include "flang/Runtime/descriptor.h"
-#include <algorithm>
 
 namespace Fortran::runtime {
 
 // Utility for CSHIFT & EOSHIFT rank > 1 cases that determines the shift count
 // for each of the vector sections of the result.
 class ShiftControl {
 public:
-  ShiftControl(const Descriptor &s, Terminator &t, int dim)
+  RT_API_ATTRS ShiftControl(const Descriptor &s, Terminator &t, int dim)
       : shift_{s}, terminator_{t}, shiftRank_{s.rank()}, dim_{dim} {}
-  void Init(const Descriptor &source, const char *which) {
+  RT_API_ATTRS void Init(const Descriptor &source, const char *which) {
     int rank{source.rank()};
     RUNTIME_CHECK(terminator_, shiftRank_ == 0 || shiftRank_ == rank - 1);
     auto catAndKind{shift_.type().GetCategoryAndKind()};
     RUNTIME_CHECK(
         terminator_, catAndKind && catAndKind->first == TypeCategory::Integer);
     shiftElemLen_ = catAndKind->second;
     if (shiftRank_ > 0) {
       int k{0};
if (shiftRank_ > 0) {
           }
         }
       }
     } else {
       shiftCount_ =
           GetInt64(shift_.OffsetElement<char>(), shiftElemLen_, terminator_);
     }
   }
-  SubscriptValue GetShift(const SubscriptValue resultAt[]) const {
+  RT_API_ATTRS SubscriptValue GetShift(const SubscriptValue resultAt[]) const {
     if (shiftRank_ > 0) {
       SubscriptValue shiftAt[maxRank];
       int k{0};
       for (int j{0}; j < shiftRank_ + 1; ++j) {
         if (j + 1 != dim_) {
           shiftAt[k] = lb_[k] + resultAt[j] - 1;
           ++k;
         }
private:
   int shiftRank_;
   int dim_;
   SubscriptValue lb_[maxRank];
   std::size_t shiftElemLen_;
   SubscriptValue shiftCount_{};
 };
 
 // Fill an EOSHIFT result with default boundary values
-static void DefaultInitialize(
+static RT_API_ATTRS void DefaultInitialize(
     const Descriptor &result, Terminator &terminator) {
   auto catAndKind{result.type().GetCategoryAndKind()};
   RUNTIME_CHECK(
       terminator, catAndKind && catAndKind->first != TypeCategory::Derived);
   std::size_t elementLen{result.ElementBytes()};
   std::size_t bytes{result.Elements() * elementLen};
   if (catAndKind->first == TypeCategory::Character) {
     switch (int kind{catAndKind->second}) {
     case 1:
-      std::fill_n(result.OffsetElement<char>(), bytes, ' ');
+      Fortran::runtime::fill_n(result.OffsetElement<char>(), bytes, ' ');
       break;
     case 2:
-      std::fill_n(result.OffsetElement<char16_t>(), bytes / 2,
+      Fortran::runtime::fill_n(result.OffsetElement<char16_t>(), bytes / 2,
           static_cast<char16_t>(' '));
       break;
     case 4:
-      std::fill_n(result.OffsetElement<char32_t>(), bytes / 4,
+      Fortran::runtime::fill_n(result.OffsetElement<char32_t>(), bytes / 4,
           static_cast<char32_t>(' '));
       break;
     default:
       terminator.Crash("not yet implemented: EOSHIFT: CHARACTER kind %d", kind);
     }
   } else {
     std::memset(result.raw().base_addr, 0, bytes);
   }
 }
 
-static inline std::size_t AllocateResult(Descriptor &result,
+static inline RT_API_ATTRS std::size_t AllocateResult(Descriptor &result,
     const Descriptor &source, int rank, const SubscriptValue extent[],
     Terminator &terminator, const char *function) {
   std::size_t elementLen{source.ElementBytes()};
   const DescriptorAddendum *sourceAddendum{source.Addendum()};
   result.Establish(source.type(), elementLen, nullptr, rank, extent,
       CFI_attribute_allocatable, sourceAddendum != nullptr);
   if (sourceAddendum) {
     *result.Addendum() = *sourceAddendum;
   }
   for (int j{0}; j < rank; ++j) {
     result.GetDimension(j).SetBounds(1, extent[j]);
   }
   if (int stat{result.Allocate()}) {
     terminator.Crash(
         "%s: Could not allocate memory for result (stat=%d)", function, stat);
   }
   return elementLen;
 }
 
 template <TypeCategory CAT, int KIND>
-static inline std::size_t AllocateBesselResult(Descriptor &result, int32_t n1,
-    int32_t n2, Terminator &terminator, const char *function) {
+static inline RT_API_ATTRS std::size_t AllocateBesselResult(Descriptor &result,
+    int32_t n1, int32_t n2, Terminator &terminator, const char *function) {
   int rank{1};
   SubscriptValue extent[maxRank];
   for (int j{0}; j < maxRank; j++) {
     extent[j] = 0;
   }
   if (n1 <= n2) {
     extent[0] = n2 - n1 + 1;
   }
 
   std::size_t elementLen{Descriptor::BytesFor(CAT, KIND)};
   result.Establish(TypeCode{CAT, KIND}, elementLen, nullptr, rank, extent,
       CFI_attribute_allocatable, false);
   for (int j{0}; j < rank; ++j) {
     result.GetDimension(j).SetBounds(1, extent[j]);
   }
   if (int stat{result.Allocate()}) {
     terminator.Crash(
         "%s: Could not allocate memory for result (stat=%d)", function, stat);
   }
   return elementLen;
 }
 
 template <TypeCategory CAT, int KIND>
-static inline void DoBesselJn(Descriptor &result, int32_t n1, int32_t n2,
-    CppTypeFor<CAT, KIND> x, CppTypeFor<CAT, KIND> bn2,
+static inline RT_API_ATTRS void DoBesselJn(Descriptor &result, int32_t n1,
+    int32_t n2, CppTypeFor<CAT, KIND> x, CppTypeFor<CAT, KIND> bn2,
     CppTypeFor<CAT, KIND> bn2_1, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   AllocateBesselResult<CAT, KIND>(result, n1, n2, terminator, "BESSEL_JN");
 
   // The standard requires that n1 and n2 be non-negative. However, some other
   // compilers generate results even when n1 and/or n2 are negative. For now,
   // we also do not enforce the non-negativity constraint.
   if (n2 < n1) {
for (int n{n2 - 2}; n >= n1; --n) {
     *result.Element<CppTypeFor<CAT, KIND>>(at) = bn;
 
     bn_2 = bn_1;
     bn_1 = bn;
   }
 }
 
 template <TypeCategory CAT, int KIND>
-static inline void DoBesselJnX0(Descriptor &result, int32_t n1, int32_t n2,
-    const char *sourceFile, int line) {
+static inline RT_API_ATTRS void DoBesselJnX0(Descriptor &result, int32_t n1,
+    int32_t n2, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   AllocateBesselResult<CAT, KIND>(result, n1, n2, terminator, "BESSEL_JN");
 
   // The standard requires that n1 and n2 be non-negative. However, some other
   // compilers generate results even when n1 and/or n2 are negative. For now,
   // we also do not enforce the non-negativity constraint.
   if (n2 < n1) {
     return;
static inline RT_API_ATTRS void DoBesselJnX0(Descriptor &result, int32_t n1,
   *result.Element<CppTypeFor<CAT, KIND>>(at) = (n1 == 0) ? 1.0 : 0.0;
   for (int j{2}; j <= n2 - n1 + 1; ++j) {
     at[0] = j;
     *result.Element<CppTypeFor<CAT, KIND>>(at) = 0.0;
   }
 }
 
 template <TypeCategory CAT, int KIND>
-static inline void DoBesselYn(Descriptor &result, int32_t n1, int32_t n2,
-    CppTypeFor<CAT, KIND> x, CppTypeFor<CAT, KIND> bn1,
+static inline RT_API_ATTRS void DoBesselYn(Descriptor &result, int32_t n1,
+    int32_t n2, CppTypeFor<CAT, KIND> x, CppTypeFor<CAT, KIND> bn1,
     CppTypeFor<CAT, KIND> bn1_1, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   AllocateBesselResult<CAT, KIND>(result, n1, n2, terminator, "BESSEL_YN");
 
   // The standard requires that n1 and n2 be non-negative. However, some other
   // compilers generate results even when n1 and/or n2 are negative. For now,
   // we also do not enforce the non-negativity constraint.
   if (n2 < n1) {
for (int n{n1 + 2}; n <= n2; ++n) {
     *result.Element<CppTypeFor<CAT, KIND>>(at) = bn;
 
     bn_2 = bn_1;
     bn_1 = bn;
   }
 }
 
 template <TypeCategory CAT, int KIND>
-static inline void DoBesselYnX0(Descriptor &result, int32_t n1, int32_t n2,
-    const char *sourceFile, int line) {
+static inline RT_API_ATTRS void DoBesselYnX0(Descriptor &result, int32_t n1,
+    int32_t n2, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   AllocateBesselResult<CAT, KIND>(result, n1, n2, terminator, "BESSEL_YN");
 
   // The standard requires that n1 and n2 be non-negative. However, some other
   // compilers generate results even when n1 and/or n2 are negative. For now,
   // we also do not enforce the non-negativity constraint.
   if (n2 < n1) {
     return;
   }
 
   SubscriptValue at[maxRank];
   for (int j{0}; j < maxRank; ++j) {
     at[j] = 0;
   }
 
   // Y(n, 0.0) = -Inf, when n >= 0
   for (int j{1}; j <= n2 - n1 + 1; ++j) {
     at[0] = j;
     *result.Element<CppTypeFor<CAT, KIND>>(at) =
         -std::numeric_limits<CppTypeFor<CAT, KIND>>::infinity();
   }
 }
 
 extern "C" {
+RT_EXT_API_GROUP_BEGIN
 
 // BESSEL_JN
 // TODO: REAL(2 & 3)
-void RTNAME(BesselJn_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJn_4)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 4> x, CppTypeFor<TypeCategory::Real, 4> bn2,
     CppTypeFor<TypeCategory::Real, 4> bn2_1, const char *sourceFile, int line) {
   DoBesselJn<TypeCategory::Real, 4>(
       result, n1, n2, x, bn2, bn2_1, sourceFile, line);
 }
 
-void RTNAME(BesselJn_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJn_8)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 8> x, CppTypeFor<TypeCategory::Real, 8> bn2,
     CppTypeFor<TypeCategory::Real, 8> bn2_1, const char *sourceFile, int line) {
   DoBesselJn<TypeCategory::Real, 8>(
       result, n1, n2, x, bn2, bn2_1, sourceFile, line);
 }
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselJn_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJn_10)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 10> x,
     CppTypeFor<TypeCategory::Real, 10> bn2,
     CppTypeFor<TypeCategory::Real, 10> bn2_1, const char *sourceFile,
     int line) {
   DoBesselJn<TypeCategory::Real, 10>(
       result, n1, n2, x, bn2, bn2_1, sourceFile, line);
 }
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselJn_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJn_16)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 16> x,
     CppTypeFor<TypeCategory::Real, 16> bn2,
     CppTypeFor<TypeCategory::Real, 16> bn2_1, const char *sourceFile,
     int line) {
   DoBesselJn<TypeCategory::Real, 16>(
       result, n1, n2, x, bn2, bn2_1, sourceFile, line);
 }
 #endif
 
 // TODO: REAL(2 & 3)
-void RTNAME(BesselJnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselJnX0<TypeCategory::Real, 4>(result, n1, n2, sourceFile, line);
 }
 
-void RTNAME(BesselJnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselJnX0<TypeCategory::Real, 8>(result, n1, n2, sourceFile, line);
 }
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselJnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselJnX0<TypeCategory::Real, 10>(result, n1, n2, sourceFile, line);
 }
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselJnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselJnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselJnX0<TypeCategory::Real, 16>(result, n1, n2, sourceFile, line);
 }
 #endif
 
 // BESSEL_YN
 // TODO: REAL(2 & 3)
-void RTNAME(BesselYn_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYn_4)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 4> x, CppTypeFor<TypeCategory::Real, 4> bn1,
     CppTypeFor<TypeCategory::Real, 4> bn1_1, const char *sourceFile, int line) {
   DoBesselYn<TypeCategory::Real, 4>(
       result, n1, n2, x, bn1, bn1_1, sourceFile, line);
 }
 
-void RTNAME(BesselYn_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYn_8)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 8> x, CppTypeFor<TypeCategory::Real, 8> bn1,
     CppTypeFor<TypeCategory::Real, 8> bn1_1, const char *sourceFile, int line) {
   DoBesselYn<TypeCategory::Real, 8>(
       result, n1, n2, x, bn1, bn1_1, sourceFile, line);
 }
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselYn_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYn_10)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 10> x,
     CppTypeFor<TypeCategory::Real, 10> bn1,
     CppTypeFor<TypeCategory::Real, 10> bn1_1, const char *sourceFile,
     int line) {
   DoBesselYn<TypeCategory::Real, 10>(
       result, n1, n2, x, bn1, bn1_1, sourceFile, line);
 }
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselYn_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYn_16)(Descriptor &result, int32_t n1, int32_t n2,
     CppTypeFor<TypeCategory::Real, 16> x,
     CppTypeFor<TypeCategory::Real, 16> bn1,
     CppTypeFor<TypeCategory::Real, 16> bn1_1, const char *sourceFile,
     int line) {
   DoBesselYn<TypeCategory::Real, 16>(
       result, n1, n2, x, bn1, bn1_1, sourceFile, line);
 }
 #endif
 
 // TODO: REAL(2 & 3)
-void RTNAME(BesselYnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYnX0_4)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselYnX0<TypeCategory::Real, 4>(result, n1, n2, sourceFile, line);
 }
 
-void RTNAME(BesselYnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYnX0_8)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselYnX0<TypeCategory::Real, 8>(result, n1, n2, sourceFile, line);
 }
 
 #if LDBL_MANT_DIG == 64
-void RTNAME(BesselYnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYnX0_10)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselYnX0<TypeCategory::Real, 10>(result, n1, n2, sourceFile, line);
 }
 #endif
 
 #if LDBL_MANT_DIG == 113 || HAS_FLOAT128
-void RTNAME(BesselYnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDEF(BesselYnX0_16)(Descriptor &result, int32_t n1, int32_t n2,
     const char *sourceFile, int line) {
   DoBesselYnX0<TypeCategory::Real, 16>(result, n1, n2, sourceFile, line);
 }
 #endif
 
 // CSHIFT where rank of ARRAY argument > 1
-void RTNAME(Cshift)(Descriptor &result, const Descriptor &source,
+void RTDEF(Cshift)(Descriptor &result, const Descriptor &source,
     const Descriptor &shift, int dim, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   int rank{source.rank()};
   RUNTIME_CHECK(terminator, rank > 1);
   if (dim < 1 || dim > rank) {
     terminator.Crash(
         "CSHIFT: DIM=%d must be >= 1 and <= SOURCE= rank %d", dim, rank);
   }
for (resDim = 1; resDim <= dimExtent; ++resDim) {
         sourceDim = dimLB;
       }
     }
     result.IncrementSubscripts(resultAt);
   }
 }
 
 // CSHIFT where rank of ARRAY argument == 1
-void RTNAME(CshiftVector)(Descriptor &result, const Descriptor &source,
+void RTDEF(CshiftVector)(Descriptor &result, const Descriptor &source,
     std::int64_t shift, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   RUNTIME_CHECK(terminator, source.rank() == 1);
   const Dimension &sourceDim{source.GetDimension(0)};
   SubscriptValue extent{sourceDim.Extent()};
   AllocateResult(result, source, 1, &extent, terminator, "CSHIFT");
   SubscriptValue lb{sourceDim.LowerBound()};
   for (SubscriptValue j{0}; j < extent; ++j) {
     SubscriptValue resultAt{1 + j};
     SubscriptValue sourceAt{lb + (j + shift) % extent};
     if (sourceAt < lb) {
       sourceAt += extent;
     }
     CopyElement(result, &resultAt, source, &sourceAt, terminator);
   }
 }
 
 // EOSHIFT of rank > 1
-void RTNAME(Eoshift)(Descriptor &result, const Descriptor &source,
+void RTDEF(Eoshift)(Descriptor &result, const Descriptor &source,
     const Descriptor &shift, const Descriptor *boundary, int dim,
     const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   SubscriptValue extent[maxRank];
   int rank{source.GetShape(extent)};
   RUNTIME_CHECK(terminator, rank > 1);
   if (dim < 1 || dim > rank) {
     terminator.Crash(
for (std::size_t n{result.Elements()}; n > 0; n -= dimExtent) {
     result.IncrementSubscripts(resultAt);
     if (boundaryRank > 0) {
       boundary->IncrementSubscripts(boundaryAt);
     }
   }
 }
 
 // EOSHIFT of vector
-void RTNAME(EoshiftVector)(Descriptor &result, const Descriptor &source,
+void RTDEF(EoshiftVector)(Descriptor &result, const Descriptor &source,
     std::int64_t shift, const Descriptor *boundary, const char *sourceFile,
     int line) {
   Terminator terminator{sourceFile, line};
   RUNTIME_CHECK(terminator, source.rank() == 1);
   SubscriptValue extent{source.GetDimension(0).Extent()};
   std::size_t elementLen{
       AllocateResult(result, source, 1, &extent, terminator, "EOSHIFT")};
   if (boundary) {
if (sourceAt >= lb && sourceAt < lb + extent) {
       CopyElement(result, &j, source, &sourceAt, terminator);
     } else if (boundary) {
       CopyElement(result, &j, *boundary, 0, terminator);
     }
   }
 }
 
 // PACK
-void RTNAME(Pack)(Descriptor &result, const Descriptor &source,
+void RTDEF(Pack)(Descriptor &result, const Descriptor &source,
     const Descriptor &mask, const Descriptor *vector, const char *sourceFile,
     int line) {
   Terminator terminator{sourceFile, line};
   CheckConformability(source, mask, terminator, "PACK", "ARRAY=", "MASK=");
   auto maskType{mask.type().GetCategoryAndKind()};
   RUNTIME_CHECK(
       terminator, maskType && maskType->first == TypeCategory::Logical);
   SubscriptValue trues{0};
if (vector) {
     for (; resultAt <= extent; ++resultAt, ++vectorAt) {
       CopyElement(result, &resultAt, *vector, &vectorAt, terminator);
     }
   }
 }
 
 // RESHAPE
 // F2018 16.9.163
-void RTNAME(Reshape)(Descriptor &result, const Descriptor &source,
+void RTDEF(Reshape)(Descriptor &result, const Descriptor &source,
     const Descriptor &shape, const Descriptor *pad, const Descriptor *order,
     const char *sourceFile, int line) {
   // Compute and check the rank of the result.
   Terminator terminator{sourceFile, line};
   RUNTIME_CHECK(terminator, shape.rank() == 1);
   RUNTIME_CHECK(terminator, shape.type().IsInteger());
   SubscriptValue resultRank{shape.GetDimension(0).Extent()};
   if (resultRank < 0 || resultRank > static_cast<SubscriptValue>(maxRank)) {
for (; resultElement < resultElements; ++resultElement) {
       CopyElement(result, resultSubscript, *pad, padSubscript, terminator);
       pad->IncrementSubscripts(padSubscript);
       result.IncrementSubscripts(resultSubscript, dimOrder);
     }
   }
 }
 
 // SPREAD
-void RTNAME(Spread)(Descriptor &result, const Descriptor &source, int dim,
+void RTDEF(Spread)(Descriptor &result, const Descriptor &source, int dim,
     std::int64_t ncopies, const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   int rank{source.rank() + 1};
   RUNTIME_CHECK(terminator, rank <= maxRank);
   if (dim < 1 || dim > rank) {
     terminator.Crash("SPREAD: DIM=%d argument for rank-%d source array "
                      "must be greater than 1 and less than or equal to %d",
         dim, rank - 1, rank);
for (resultDim = 1; resultDim <= ncopies; ++resultDim) {
       CopyElement(result, resultAt, source, sourceAt, terminator);
     }
     result.IncrementSubscripts(resultAt);
     source.IncrementSubscripts(sourceAt);
   }
 }
 
 // TRANSPOSE
-void RTNAME(Transpose)(Descriptor &result, const Descriptor &matrix,
+void RTDEF(Transpose)(Descriptor &result, const Descriptor &matrix,
     const char *sourceFile, int line) {
   Terminator terminator{sourceFile, line};
   RUNTIME_CHECK(terminator, matrix.rank() == 2);
   SubscriptValue extent[2]{
       matrix.GetDimension(1).Extent(), matrix.GetDimension(0).Extent()};
   AllocateResult(result, matrix, 2, extent, terminator, "TRANSPOSE");
   SubscriptValue resultAt[2]{1, 1};
   SubscriptValue matrixLB[2];
   matrix.GetLowerBounds(matrixLB);
   for (std::size_t n{result.Elements()}; n-- > 0;
        result.IncrementSubscripts(resultAt)) {
     SubscriptValue matrixAt[2]{
         matrixLB[0] + resultAt[1] - 1, matrixLB[1] + resultAt[0] - 1};
     CopyElement(result, resultAt, matrix, matrixAt, terminator);
   }
 }
 
 // UNPACK
-void RTNAME(Unpack)(Descriptor &result, const Descriptor &vector,
+void RTDEF(Unpack)(Descriptor &result, const Descriptor &vector,
     const Descriptor &mask, const Descriptor &field, const char *sourceFile,
     int line) {
   Terminator terminator{sourceFile, line};
   RUNTIME_CHECK(terminator, vector.rank() == 1);
   int rank{mask.rank()};
   RUNTIME_CHECK(terminator, rank > 0);
   SubscriptValue extent[maxRank];
   mask.GetShape(extent);
if (IsLogicalElementTrue(mask, maskAt)) {
       CopyElement(result, resultAt, field, fieldAt, terminator);
     }
     result.IncrementSubscripts(resultAt);
     mask.IncrementSubscripts(maskAt);
     field.IncrementSubscripts(fieldAt);
   }
 }
 
+RT_EXT_API_GROUP_END
 } // extern "C"
 } // namespace Fortran::runtime