Warning: This file is not a C or C++ file. It does not have highlighting.
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 | #endif /* !FORTRAN_RUNTIME_API_ATTRS_H_ */ |
160 |
Warning: This file is not a C or C++ file. It does not have highlighting.