[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

file*

25

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

131

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

131

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/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/device.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); }
+  HOST_DEVICE_ATTRS T &operator[](int index) { return *(this + index); }
+  const HOST_DEVICE_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 *);
+HOST_DEVICE_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,
+HOST_DEVICE_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/descriptor.h

 // to implement nontrivial dummy arguments, pointers, allocatables,
 // function results, and the special behaviors of instances of derived types.
 // This header file includes and extends the published language
 // interoperability header that is required by the Fortran 2018 standard
 // as a subset of definitions suitable for exposure to user C/C++ code.
 // User C code is welcome to depend on that ISO_Fortran_binding.h file,
 // but should never reference this internal header.
 
+#include "flang/Runtime/device.h"
+
 #include "flang/ISO_Fortran_binding.h"
 #include "flang/Runtime/memory.h"
 #include "flang/Runtime/type-code.h"
 #include <algorithm>
 #include <cassert>
 #include <cinttypes>
 #include <cstddef>
 #include <cstdio>

 
 static constexpr int maxRank{CFI_MAX_RANK};
 
 // 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; }
+  HOST_DEVICE_ATTRS SubscriptValue LowerBound() const {
+    return raw_.lower_bound;
+  }
+  HOST_DEVICE_ATTRS SubscriptValue Extent() const { return raw_.extent; }
+  HOST_DEVICE_ATTRS SubscriptValue UpperBound() const {
+    return LowerBound() + Extent() - 1;
+  }
+  HOST_DEVICE_ATTRS SubscriptValue ByteStride() const { return raw_.sm; }
 
-  Dimension &SetBounds(SubscriptValue lower, SubscriptValue upper) {
+  HOST_DEVICE_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) {
+  HOST_DEVICE_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 HOST_DEVICE_ATTRS DescriptorAddendum(
+      const typeInfo::DerivedType *dt = nullptr)
       : derivedType_{dt} {}
-  DescriptorAddendum &operator=(const DescriptorAddendum &);
+  HOST_DEVICE_ATTRS DescriptorAddendum &operator=(const DescriptorAddendum &);
 
-  const typeInfo::DerivedType *derivedType() const { return derivedType_; }
-  DescriptorAddendum &set_derivedType(const typeInfo::DerivedType *dt) {
+  const HOST_DEVICE_ATTRS typeInfo::DerivedType *derivedType() const {
+    return derivedType_;
+  }
+  HOST_DEVICE_ATTRS DescriptorAddendum &set_derivedType(
+      const typeInfo::DerivedType *dt) {
     derivedType_ = dt;
     return *this;
   }
 
-  std::size_t LenParameters() const;
+  HOST_DEVICE_ATTRS std::size_t LenParameters() const;
 
-  typeInfo::TypeParameterValue LenParameterValue(int which) const {
+  HOST_DEVICE_ATTRS typeInfo::TypeParameterValue LenParameterValue(
+      int which) const {
     return len_[which];
   }
-  static constexpr std::size_t SizeInBytes(int lenParameters) {
+  static constexpr HOST_DEVICE_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;
+  HOST_DEVICE_ATTRS std::size_t SizeInBytes() const;
 
-  void SetLenParameterValue(int which, typeInfo::TypeParameterValue x) {
+  HOST_DEVICE_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 &);
+  HOST_DEVICE_ATTRS Descriptor(const Descriptor &);
+  HOST_DEVICE_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 HOST_DEVICE_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,
+  HOST_DEVICE_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,
+  HOST_DEVICE_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,
+  HOST_DEVICE_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);
+  HOST_DEVICE_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}; }
+  HOST_DEVICE_ATTRS ISO::CFI_cdesc_t &raw() { return raw_; }
+  const HOST_DEVICE_ATTRS ISO::CFI_cdesc_t &raw() const { return raw_; }
+  HOST_DEVICE_ATTRS std::size_t ElementBytes() const { return raw_.elem_len; }
+  HOST_DEVICE_ATTRS int rank() const { return raw_.rank; }
+  HOST_DEVICE_ATTRS TypeCode type() const { return TypeCode{raw_.type}; }
 
-  Descriptor &set_base_addr(void *p) {
+  HOST_DEVICE_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 {
+  HOST_DEVICE_ATTRS bool IsPointer() const {
+    return raw_.attribute == CFI_attribute_pointer;
+  }
+  HOST_DEVICE_ATTRS bool IsAllocatable() const {
     return raw_.attribute == CFI_attribute_allocatable;
   }
-  bool IsAllocated() const { return raw_.base_addr != nullptr; }
+  HOST_DEVICE_ATTRS bool IsAllocated() const {
+    return raw_.base_addr != nullptr;
+  }
 
-  Dimension &GetDimension(int dim) {
+  HOST_DEVICE_ATTRS Dimension &GetDimension(int dim) {
     return *reinterpret_cast<Dimension *>(&raw_.dim[dim]);
   }
-  const Dimension &GetDimension(int dim) const {
+  const HOST_DEVICE_ATTRS Dimension &GetDimension(int dim) const {
     return *reinterpret_cast<const Dimension *>(&raw_.dim[dim]);
   }
 
-  std::size_t SubscriptByteOffset(
+  HOST_DEVICE_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 {
+  HOST_DEVICE_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>
+  HOST_DEVICE_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>
+  HOST_DEVICE_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>
+  HOST_DEVICE_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 {
+  HOST_DEVICE_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 {
+  HOST_DEVICE_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(
+  HOST_DEVICE_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(
+  HOST_DEVICE_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 {
+  HOST_DEVICE_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(
+  HOST_DEVICE_ATTRS std::size_t ZeroBasedElementNumber(
       const SubscriptValue *, const int *permutation = nullptr) const;
 
-  DescriptorAddendum *Addendum() {
+  HOST_DEVICE_ATTRS DescriptorAddendum *Addendum() {
     if (raw_.f18Addendum != 0) {
       return reinterpret_cast<DescriptorAddendum *>(&GetDimension(rank()));
     } else {
       return nullptr;
     }
   }
-  const DescriptorAddendum *Addendum() const {
+  const HOST_DEVICE_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 HOST_DEVICE_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;
+  HOST_DEVICE_ATTRS std::size_t SizeInBytes() const;
 
-  std::size_t Elements() const;
+  HOST_DEVICE_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();
+  HOST_DEVICE_ATTRS int Allocate();
 
   // Deallocates storage; does not call FINAL subroutines or
   // deallocate allocatable/automatic components.
-  int Deallocate();
+  HOST_DEVICE_ATTRS int Deallocate();
 
   // Deallocates storage, including allocatable and automatic
   // components.  Optionally invokes FINAL subroutines.
-  int Destroy(bool finalize = false, bool destroyPointers = false);
+  HOST_DEVICE_ATTRS int Destroy(
+      bool finalize = false, bool destroyPointers = false);
 
-  bool IsContiguous(int leadingDimensions = maxRank) const {
+  HOST_DEVICE_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,
+  HOST_DEVICE_ATTRS bool EstablishPointerSection(const Descriptor &source,
       const SubscriptValue *lower = nullptr,
       const SubscriptValue *upper = nullptr,
       const SubscriptValue *stride = nullptr);
 
-  void Check() const;
+  HOST_DEVICE_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 {
+  HOST_DEVICE_ATTRS Descriptor &descriptor() {
+    return *reinterpret_cast<Descriptor *>(storage_);
+  }
+  const HOST_DEVICE_ATTRS Descriptor &descriptor() const {
     return *reinterpret_cast<const Descriptor *>(storage_);
   }
 
-  void Check() {
+  HOST_DEVICE_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/device.h

+/*===-- include/flang/Runtime/device.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_DEVICE_H
+#define FORTRAN_RUNTIME_DEVICE_H
+
+#ifdef __CUDA__
+#define HOST_DEVICE_ATTRS __host__ __device__
+
+#ifdef __CUDA_ARCH__
+#define CUDA_DEVICE 1
+#endif
+
+#else
+#define HOST_DEVICE_ATTRS
+#endif
+
+#endif /* !FORTRAN_RUNTIME_DEVICE_H */

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/device.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) HOST_DEVICE_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 HOST_DEVICE_ATTRS TypeCode(ISO::CFI_type_t t) : raw_{t} {}
+  HOST_DEVICE_ATTRS TypeCode(TypeCategory, int kind);
 
-  int raw() const { return raw_; }
+  HOST_DEVICE_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;
+  HOST_DEVICE_ATTRS std::optional<std::pair<TypeCategory, int>>
+  GetCategoryAndKind() const;
 
-  bool operator==(const TypeCode &that) const { return raw_ == that.raw_; }
+  HOST_DEVICE_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

if (CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
   if ("${CMAKE_VERSION}" VERSION_LESS "3.20.0")
     message(WARNING
       "Your CMake version is ${CMAKE_VERSION}. Starting with LLVM 17.0.0, the "
       "minimum version of CMake required to build LLVM will become 3.20.0, and "
       "using an older CMake will become an error. Please upgrade your CMake to "
       "at least 3.20.0 now to avoid issues in the future!")
   endif()
 
-  project(FlangRuntime C CXX)
+  option(FLANG_EXPERIMENTAL_CUDA_RUNTIME
+    "Compile Fortran runtime as CUDA sources (experimental)" OFF
+    )
+
+  if (FLANG_EXPERIMENTAL_CUDA_RUNTIME)
+    set(enable_cuda_lang CUDA)
+  endif()
+
+  project(FlangRuntime C CXX ${enable_cuda_lang})
 
   set(CMAKE_CXX_STANDARD 17)
   set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
   set(CMAKE_CXX_EXTENSIONS OFF)
 
   set(FLANG_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}/..")
 
   set(LLVM_COMMON_CMAKE_UTILS "${FLANG_SOURCE_DIR}/../cmake")
else ()
     add_compile_definitions(FLANG_LITTLE_ENDIAN=1)
   endif ()
   include_directories(BEFORE
     ${FLANG_SOURCE_DIR}/include)
 endif()
 
 include(CheckCXXSymbolExists)
 include(CheckCXXSourceCompiles)
+# TODO: in CUDA build mode these checks should be done
+#       for the device side as well, and probably instead
+#       of failing during CMake configuration the corresponding
+#       functionality has to be disabled/customized for the device
+#       compilation.
 check_cxx_symbol_exists(strerror string.h HAVE_STRERROR)
 check_cxx_symbol_exists(strerror_r string.h HAVE_STRERROR_R)
 # Can't use symbol exists here as the function is overloaded in C++
 check_cxx_source_compiles(
   "#include <string.h>
    int main() {
      char buf[4096];
      return strerror_s(buf, 4096, 0);

 
 # 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
+)
 
+if (FLANG_EXPERIMENTAL_CUDA_RUNTIME)
+  # List of files that are buildable.
+  set(supported_files
+    transformational.cpp
+    )
+
+  # 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})
+
+  # Treat all sources as CUDA files.
+  set(sources ${supported_files})
+  set_source_files_properties(${sources} PROPERTIES LANGUAGE CUDA)
+  if ("${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
+    # Allow varargs.
+    add_compile_options(-Xclang -fcuda-allow-variadic-functions)
+  endif()
+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[],
-    const Descriptor &from, const SubscriptValue fromAt[], Terminator &);
+HOST_DEVICE_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 &);
+HOST_DEVICE_ATTRS void CopyArray(
+    const Descriptor &to, const Descriptor &from, Terminator &);
 
 } // namespace Fortran::runtime
 #endif // FORTRAN_RUNTIME_COPY_H_

flang/runtime/device-tools.h

+//===-- runtime/device-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_DEVICE_TOOLS_H
+#define FORTRAN_RUNTIME_DEVICE_TOOLS_H
+
+#include "flang/Runtime/device.h"
+#include <algorithm>
+
+// The files defines a set of utilities/classes that might be not
+// supported on the offload devices via standard libraries (e.g. STD C++).
+// In some cases the implementations are temporary until the device
+// toolchains implement support on their side.
+
+namespace Fortran::runtime {
+
+// Provides alternative implementation for std::fill_n(), if
+// it is not supported.
+template <typename A>
+static inline HOST_DEVICE_ATTRS void FillN(
+    A *start, std::size_t count, const A &value) {
+#if defined(CUDA_DEVICE) && defined(__clang__)
+  // Clang does not support std::fill_n as a device function.
+  for (std::size_t j{0}; j < count; ++j)
+    start[j] = value;
+#else
+  std::fill_n(start, count, value);
+#endif
+}
+
+} // namespace Fortran::runtime
+#endif // FORTRAN_RUNTIME_DEVICE_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/device.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 HOST_DEVICE_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]] HOST_DEVICE_ATTRS void Crash(const char *message, ...) const;
+  [[noreturn]] HOST_DEVICE_ATTRS void CrashArgs(
+      const char *message, va_list &) const;
+  [[noreturn]] HOST_DEVICE_ATTRS void CheckFailed(
       const char *predicate, const char *file, int line) const;
-  [[noreturn]] void CheckFailed(const char *predicate) const;
+  [[noreturn]] HOST_DEVICE_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 "device-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 HOST_DEVICE_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,
-    Terminator &, const char *funcName, const char *toName,
+HOST_DEVICE_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);
+HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_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)...);
     case 4:
       return FUNC<TypeCategory::Integer, 4>{}(std::forward<A>(x)...);
     case 8:
       return FUNC<TypeCategory::Integer, 8>{}(std::forward<A>(x)...);
-#ifdef __SIZEOF_INT128__
     case 16:
       return FUNC<TypeCategory::Integer, 16>{}(std::forward<A>(x)...);
-#endif
     default:
       terminator.Crash("not yet implemented: INTEGER(KIND=%d)", kind);
     }
   case TypeCategory::Real:
     switch (kind) {
 #if 0 // TODO: REAL(2 & 3)
     case 2:
       return FUNC<TypeCategory::Real, 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 HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_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)
+  HOST_DEVICE_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) {
+  HOST_DEVICE_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 {
+  HOST_DEVICE_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 HOST_DEVICE_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, ' ');
+      FillN(result.OffsetElement<char>(), bytes, ' ');
       break;
     case 2:
-      std::fill_n(result.OffsetElement<char16_t>(), bytes / 2,
+      FillN(result.OffsetElement<char16_t>(), bytes / 2,
           static_cast<char16_t>(' '));
       break;
     case 4:
-      std::fill_n(result.OffsetElement<char32_t>(), bytes / 4,
+      FillN(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 HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_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 HOST_DEVICE_ATTRS void DoBesselJnX0(Descriptor &result,
   *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 HOST_DEVICE_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 HOST_DEVICE_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;
*result.Element<CppTypeFor<CAT, KIND>>(at) =
         -std::numeric_limits<CppTypeFor<CAT, KIND>>::infinity();
   }
 }
 
 extern "C" {
 
 // BESSEL_JN
 // TODO: REAL(2 & 3)
-void RTNAME(BesselJn_4)(Descriptor &result, int32_t n1, int32_t n2,
+void RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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 RTDECL(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);