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

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

Differential Revision: https://reviews.llvm.org/D151173

Author: vzakhari

flang/docs/GettingStarted.md

@@ -180,6 +180,76 @@ directory:
 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.

flang/include/flang/ISO_Fortran_binding.h

@@ -18,6 +18,8 @@
  * implementation.
  */
 
+#include "Runtime/api-attrs.h"
+
 #ifdef __cplusplus
 namespace Fortran {
 namespace ISO {
@@ -121,8 +123,8 @@ namespace cfi_internal {
 // 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; }
 };
@@ -174,11 +176,11 @@ extern "C" {
 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,

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

@@ -0,0 +1,91 @@
+/*===-- 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_ */