1 /*===-- include/flang/Common/api-attrs.h ---------------------------*- C -*-=// 2 * 3 * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 * See https://llvm.org/LICENSE.txt for license information. 5 * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 * 7 *===------------------------------------------------------------------------=== 8 */ 9 10 /* 11 * The file defines a set macros that can be used to apply 12 * different attributes/pragmas to functions/variables 13 * declared/defined/used in Flang runtime library. 14 */ 15 16 #ifndef FORTRAN_RUNTIME_API_ATTRS_H_ 17 #define FORTRAN_RUNTIME_API_ATTRS_H_ 18 19 /* 20 * RT_EXT_API_GROUP_BEGIN/END pair is placed around definitions 21 * of functions exported by Flang runtime library. They are the entry 22 * points that are referenced in the Flang generated code. 23 * The macros may be expanded into any construct that is valid to appear 24 * at C++ module scope. 25 */ 26 #ifndef RT_EXT_API_GROUP_BEGIN 27 #if defined(OMP_NOHOST_BUILD) 28 #define RT_EXT_API_GROUP_BEGIN \ 29 _Pragma("omp begin declare target device_type(nohost)") 30 #elif defined(OMP_OFFLOAD_BUILD) 31 #define RT_EXT_API_GROUP_BEGIN _Pragma("omp declare target") 32 #else 33 #define RT_EXT_API_GROUP_BEGIN 34 #endif 35 #endif /* !defined(RT_EXT_API_GROUP_BEGIN) */ 36 37 #ifndef RT_EXT_API_GROUP_END 38 #if defined(OMP_NOHOST_BUILD) || defined(OMP_OFFLOAD_BUILD) 39 #define RT_EXT_API_GROUP_END _Pragma("omp end declare target") 40 #else 41 #define RT_EXT_API_GROUP_END 42 #endif 43 #endif /* !defined(RT_EXT_API_GROUP_END) */ 44 45 /* 46 * RT_OFFLOAD_API_GROUP_BEGIN/END pair is placed around definitions 47 * of functions that can be referenced in other modules of Flang 48 * runtime. For OpenMP offload, these functions are made "declare target" 49 * making sure they are compiled for the target even though direct 50 * references to them from other "declare target" functions may not 51 * be seen. Host-only functions should not be put in between these 52 * two macros. 53 */ 54 #define RT_OFFLOAD_API_GROUP_BEGIN RT_EXT_API_GROUP_BEGIN 55 #define RT_OFFLOAD_API_GROUP_END RT_EXT_API_GROUP_END 56 57 /* 58 * RT_OFFLOAD_VAR_GROUP_BEGIN/END pair is placed around definitions 59 * of variables (e.g. globals or static class members) that can be 60 * referenced in functions marked with RT_OFFLOAD_API_GROUP_BEGIN/END. 61 * For OpenMP offload, these variables are made "declare target". 62 */ 63 #define RT_OFFLOAD_VAR_GROUP_BEGIN RT_EXT_API_GROUP_BEGIN 64 #define RT_OFFLOAD_VAR_GROUP_END RT_EXT_API_GROUP_END 65 66 /* 67 * RT_VAR_GROUP_BEGIN/END pair is placed around definitions 68 * of module scope variables referenced by Flang runtime (directly 69 * or indirectly). 70 * The macros may be expanded into any construct that is valid to appear 71 * at C++ module scope. 72 */ 73 #ifndef RT_VAR_GROUP_BEGIN 74 #define RT_VAR_GROUP_BEGIN RT_EXT_API_GROUP_BEGIN 75 #endif /* !defined(RT_VAR_GROUP_BEGIN) */ 76 77 #ifndef RT_VAR_GROUP_END 78 #define RT_VAR_GROUP_END RT_EXT_API_GROUP_END 79 #endif /* !defined(RT_VAR_GROUP_END) */ 80 81 /* 82 * Each non-exported function used by Flang runtime (e.g. via 83 * calling it or taking its address, etc.) is marked with 84 * RT_API_ATTRS. The macros is placed at both declaration and 85 * definition of such a function. 86 * The macros may be expanded into a construct that is valid 87 * to appear as part of a C++ decl-specifier. 88 */ 89 #ifndef RT_API_ATTRS 90 #if defined(__CUDACC__) || defined(__CUDA__) 91 #define RT_API_ATTRS __host__ __device__ 92 #else 93 #define RT_API_ATTRS 94 #endif 95 #endif /* !defined(RT_API_ATTRS) */ 96 97 /* 98 * Each const/constexpr module scope variable referenced by Flang runtime 99 * (directly or indirectly) is marked with RT_CONST_VAR_ATTRS. 100 * The macros is placed at both declaration and definition of such a variable. 101 * The macros may be expanded into a construct that is valid 102 * to appear as part of a C++ decl-specifier. 103 */ 104 #ifndef RT_CONST_VAR_ATTRS 105 #if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) 106 #define RT_CONST_VAR_ATTRS __constant__ 107 #else 108 #define RT_CONST_VAR_ATTRS 109 #endif 110 #endif /* !defined(RT_CONST_VAR_ATTRS) */ 111 112 /* 113 * RT_VAR_ATTRS is marking non-const/constexpr module scope variables 114 * referenced by Flang runtime. 115 */ 116 #ifndef RT_VAR_ATTRS 117 #if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) 118 #define RT_VAR_ATTRS __device__ 119 #else 120 #define RT_VAR_ATTRS 121 #endif 122 #endif /* !defined(RT_VAR_ATTRS) */ 123 124 /* 125 * RT_DEVICE_COMPILATION is defined for any device compilation. 126 * Note that it can only be used reliably with compilers that perform 127 * separate host and device compilations. 128 */ 129 #if ((defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__)) || \ 130 (defined(_OPENMP) && (defined(__AMDGCN__) || defined(__NVPTX__))) 131 #define RT_DEVICE_COMPILATION 1 132 #else 133 #undef RT_DEVICE_COMPILATION 134 #endif 135 136 /* 137 * Recurrence in the call graph prevents computing minimal stack size 138 * required for a kernel execution. This macro can be used to disable 139 * some F18 runtime functionality that is implemented using recurrent 140 * function calls or to use alternative implementation. 141 */ 142 #if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) 143 #define RT_DEVICE_AVOID_RECURSION 1 144 #else 145 #undef RT_DEVICE_AVOID_RECURSION 146 #endif 147 148 #if defined(__CUDACC__) 149 #define RT_DIAG_PUSH _Pragma("nv_diagnostic push") 150 #define RT_DIAG_POP _Pragma("nv_diagnostic pop") 151 #define RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN \ 152 _Pragma("nv_diag_suppress 20011") _Pragma("nv_diag_suppress 20014") 153 #else /* !defined(__CUDACC__) */ 154 #define RT_DIAG_PUSH 155 #define RT_DIAG_POP 156 #define RT_DIAG_DISABLE_CALL_HOST_FROM_DEVICE_WARN 157 #endif /* !defined(__CUDACC__) */ 158 159 /* 160 * RT_DEVICE_NOINLINE may be used for non-performance critical 161 * functions that should not be inlined to minimize the amount 162 * of code that needs to be processed by the device compiler's 163 * optimizer. 164 */ 165 #ifndef __has_attribute 166 #define __has_attribute(x) 0 167 #endif 168 #if __has_attribute(noinline) 169 #define RT_NOINLINE_ATTR __attribute__((noinline)) 170 #else 171 #define RT_NOINLINE_ATTR 172 #endif 173 #if (defined(__CUDACC__) || defined(__CUDA__)) && defined(__CUDA_ARCH__) 174 #define RT_DEVICE_NOINLINE RT_NOINLINE_ATTR 175 #define RT_DEVICE_NOINLINE_HOST_INLINE 176 #else 177 #define RT_DEVICE_NOINLINE 178 #define RT_DEVICE_NOINLINE_HOST_INLINE inline 179 #endif 180 181 /* RT_OPTNONE_ATTR allows disabling optimizations per function. */ 182 #if __has_attribute(optimize) 183 /* GCC style. */ 184 #define RT_OPTNONE_ATTR __attribute__((optimize("O0"))) 185 #elif __has_attribute(optnone) 186 /* Clang style. */ 187 #define RT_OPTNONE_ATTR __attribute__((optnone)) 188 #else 189 #define RT_OPTNONE_ATTR 190 #endif 191 192 #endif /* !FORTRAN_RUNTIME_API_ATTRS_H_ */ 193