1 | //===--- BuiltinsPTX.def - PTX Builtin function database ----*- 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 | // This file defines the PTX-specific builtin function database. Users of |
10 | // this file must define the BUILTIN macro to make use of this information. |
11 | // |
12 | //===----------------------------------------------------------------------===// |
13 | |
14 | // The format of this database matches clang/Basic/Builtins.def. |
15 | |
16 | #if defined(BUILTIN) && !defined(TARGET_BUILTIN) |
17 | # define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) |
18 | #endif |
19 | |
20 | #pragma push_macro("SM_53") |
21 | #pragma push_macro("SM_70") |
22 | #pragma push_macro("SM_72") |
23 | #pragma push_macro("SM_75") |
24 | #pragma push_macro("SM_80") |
25 | #pragma push_macro("SM_86") |
26 | #pragma push_macro("SM_87") |
27 | #pragma push_macro("SM_89") |
28 | #pragma push_macro("SM_90") |
29 | #pragma push_macro("SM_90a") |
30 | #define SM_90a "sm_90a" |
31 | #define SM_90 "sm_90|" SM_90a |
32 | #define SM_89 "sm_89|" SM_90 |
33 | #define SM_87 "sm_87|" SM_89 |
34 | #define SM_86 "sm_86|" SM_87 |
35 | #define SM_80 "sm_80|" SM_86 |
36 | #define SM_75 "sm_75|" SM_80 |
37 | #define SM_72 "sm_72|" SM_75 |
38 | #define SM_70 "sm_70|" SM_72 |
39 | |
40 | #pragma push_macro("SM_60") |
41 | #define SM_60 "sm_60|sm_61|sm_62|" SM_70 |
42 | #define SM_53 "sm_53|" SM_60 |
43 | |
44 | #pragma push_macro("PTX42") |
45 | #pragma push_macro("PTX60") |
46 | #pragma push_macro("PTX61") |
47 | #pragma push_macro("PTX62") |
48 | #pragma push_macro("PTX63") |
49 | #pragma push_macro("PTX64") |
50 | #pragma push_macro("PTX65") |
51 | #pragma push_macro("PTX70") |
52 | #pragma push_macro("PTX71") |
53 | #pragma push_macro("PTX72") |
54 | #pragma push_macro("PTX73") |
55 | #pragma push_macro("PTX74") |
56 | #pragma push_macro("PTX75") |
57 | #pragma push_macro("PTX76") |
58 | #pragma push_macro("PTX77") |
59 | #pragma push_macro("PTX78") |
60 | #pragma push_macro("PTX80") |
61 | #pragma push_macro("PTX81") |
62 | #pragma push_macro("PTX82") |
63 | #pragma push_macro("PTX83") |
64 | #define PTX83 "ptx83" |
65 | #define PTX82 "ptx82|" PTX83 |
66 | #define PTX81 "ptx81|" PTX82 |
67 | #define PTX80 "ptx80|" PTX81 |
68 | #define PTX78 "ptx78|" PTX80 |
69 | #define PTX77 "ptx77|" PTX78 |
70 | #define PTX76 "ptx76|" PTX77 |
71 | #define PTX75 "ptx75|" PTX76 |
72 | #define PTX74 "ptx74|" PTX75 |
73 | #define PTX73 "ptx73|" PTX74 |
74 | #define PTX72 "ptx72|" PTX73 |
75 | #define PTX71 "ptx71|" PTX72 |
76 | #define PTX70 "ptx70|" PTX71 |
77 | #define PTX65 "ptx65|" PTX70 |
78 | #define PTX64 "ptx64|" PTX65 |
79 | #define PTX63 "ptx63|" PTX64 |
80 | #define PTX62 "ptx62|" PTX63 |
81 | #define PTX61 "ptx61|" PTX62 |
82 | #define PTX60 "ptx60|" PTX61 |
83 | #define PTX42 "ptx42|" PTX60 |
84 | |
85 | #pragma push_macro("AND") |
86 | #define AND(a, b) "(" a "),(" b ")" |
87 | |
88 | // Special Registers |
89 | |
90 | BUILTIN(__nvvm_read_ptx_sreg_tid_x, "i" , "nc" ) |
91 | BUILTIN(__nvvm_read_ptx_sreg_tid_y, "i" , "nc" ) |
92 | BUILTIN(__nvvm_read_ptx_sreg_tid_z, "i" , "nc" ) |
93 | BUILTIN(__nvvm_read_ptx_sreg_tid_w, "i" , "nc" ) |
94 | |
95 | BUILTIN(__nvvm_read_ptx_sreg_ntid_x, "i" , "nc" ) |
96 | BUILTIN(__nvvm_read_ptx_sreg_ntid_y, "i" , "nc" ) |
97 | BUILTIN(__nvvm_read_ptx_sreg_ntid_z, "i" , "nc" ) |
98 | BUILTIN(__nvvm_read_ptx_sreg_ntid_w, "i" , "nc" ) |
99 | |
100 | BUILTIN(__nvvm_read_ptx_sreg_ctaid_x, "i" , "nc" ) |
101 | BUILTIN(__nvvm_read_ptx_sreg_ctaid_y, "i" , "nc" ) |
102 | BUILTIN(__nvvm_read_ptx_sreg_ctaid_z, "i" , "nc" ) |
103 | BUILTIN(__nvvm_read_ptx_sreg_ctaid_w, "i" , "nc" ) |
104 | |
105 | BUILTIN(__nvvm_read_ptx_sreg_nctaid_x, "i" , "nc" ) |
106 | BUILTIN(__nvvm_read_ptx_sreg_nctaid_y, "i" , "nc" ) |
107 | BUILTIN(__nvvm_read_ptx_sreg_nctaid_z, "i" , "nc" ) |
108 | BUILTIN(__nvvm_read_ptx_sreg_nctaid_w, "i" , "nc" ) |
109 | |
110 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_x, "i" , "nc" , AND(SM_90, PTX78)) |
111 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_y, "i" , "nc" , AND(SM_90, PTX78)) |
112 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_z, "i" , "nc" , AND(SM_90, PTX78)) |
113 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_clusterid_w, "i" , "nc" , AND(SM_90, PTX78)) |
114 | |
115 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_x, "i" , "nc" , AND(SM_90, PTX78)) |
116 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_y, "i" , "nc" , AND(SM_90, PTX78)) |
117 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_z, "i" , "nc" , AND(SM_90, PTX78)) |
118 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_nclusterid_w, "i" , "nc" , AND(SM_90, PTX78)) |
119 | |
120 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_x, "i" , "nc" , AND(SM_90, PTX78)) |
121 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_y, "i" , "nc" , AND(SM_90, PTX78)) |
122 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_z, "i" , "nc" , AND(SM_90, PTX78)) |
123 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctaid_w, "i" , "nc" , AND(SM_90, PTX78)) |
124 | |
125 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_x, "i" , "nc" , AND(SM_90, PTX78)) |
126 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_y, "i" , "nc" , AND(SM_90, PTX78)) |
127 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_z, "i" , "nc" , AND(SM_90, PTX78)) |
128 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctaid_w, "i" , "nc" , AND(SM_90, PTX78)) |
129 | |
130 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_ctarank, "i" , "nc" , AND(SM_90, PTX78)) |
131 | TARGET_BUILTIN(__nvvm_read_ptx_sreg_cluster_nctarank, "i" , "nc" , AND(SM_90, PTX78)) |
132 | |
133 | TARGET_BUILTIN(__nvvm_is_explicit_cluster, "b" , "nc" , AND(SM_90, PTX78)) |
134 | |
135 | BUILTIN(__nvvm_read_ptx_sreg_laneid, "i" , "nc" ) |
136 | BUILTIN(__nvvm_read_ptx_sreg_warpid, "i" , "nc" ) |
137 | BUILTIN(__nvvm_read_ptx_sreg_nwarpid, "i" , "nc" ) |
138 | |
139 | BUILTIN(__nvvm_read_ptx_sreg_smid, "i" , "nc" ) |
140 | BUILTIN(__nvvm_read_ptx_sreg_nsmid, "i" , "nc" ) |
141 | BUILTIN(__nvvm_read_ptx_sreg_gridid, "i" , "nc" ) |
142 | |
143 | BUILTIN(__nvvm_read_ptx_sreg_lanemask_eq, "i" , "nc" ) |
144 | BUILTIN(__nvvm_read_ptx_sreg_lanemask_le, "i" , "nc" ) |
145 | BUILTIN(__nvvm_read_ptx_sreg_lanemask_lt, "i" , "nc" ) |
146 | BUILTIN(__nvvm_read_ptx_sreg_lanemask_ge, "i" , "nc" ) |
147 | BUILTIN(__nvvm_read_ptx_sreg_lanemask_gt, "i" , "nc" ) |
148 | |
149 | BUILTIN(__nvvm_read_ptx_sreg_clock, "i" , "n" ) |
150 | BUILTIN(__nvvm_read_ptx_sreg_clock64, "LLi" , "n" ) |
151 | BUILTIN(__nvvm_read_ptx_sreg_globaltimer, "LLi" , "n" ) |
152 | |
153 | BUILTIN(__nvvm_read_ptx_sreg_pm0, "i" , "n" ) |
154 | BUILTIN(__nvvm_read_ptx_sreg_pm1, "i" , "n" ) |
155 | BUILTIN(__nvvm_read_ptx_sreg_pm2, "i" , "n" ) |
156 | BUILTIN(__nvvm_read_ptx_sreg_pm3, "i" , "n" ) |
157 | |
158 | // MISC |
159 | |
160 | BUILTIN(__nvvm_prmt, "UiUiUiUi" , "" ) |
161 | BUILTIN(__nvvm_exit, "v" , "r" ) |
162 | BUILTIN(__nvvm_reflect, "UicC*" , "r" ) |
163 | TARGET_BUILTIN(__nvvm_nanosleep, "vUi" , "n" , AND(SM_70, PTX63)) |
164 | |
165 | // Min Max |
166 | |
167 | TARGET_BUILTIN(__nvvm_fmin_f16, "hhh" , "" , AND(SM_80, PTX70)) |
168 | TARGET_BUILTIN(__nvvm_fmin_ftz_f16, "hhh" , "" , AND(SM_80, PTX70)) |
169 | TARGET_BUILTIN(__nvvm_fmin_nan_f16, "hhh" , "" , AND(SM_80, PTX70)) |
170 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16, "hhh" , "" , AND(SM_80, PTX70)) |
171 | TARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
172 | TARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
173 | TARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
174 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f16, "hhh" , "" , |
175 | AND(SM_86, PTX72)) |
176 | TARGET_BUILTIN(__nvvm_fmin_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
177 | TARGET_BUILTIN(__nvvm_fmin_ftz_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
178 | TARGET_BUILTIN(__nvvm_fmin_nan_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
179 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
180 | TARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
181 | AND(SM_86, PTX72)) |
182 | TARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
183 | AND(SM_86, PTX72)) |
184 | TARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
185 | AND(SM_86, PTX72)) |
186 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
187 | AND(SM_86, PTX72)) |
188 | TARGET_BUILTIN(__nvvm_fmin_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
189 | TARGET_BUILTIN(__nvvm_fmin_ftz_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
190 | TARGET_BUILTIN(__nvvm_fmin_nan_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
191 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
192 | TARGET_BUILTIN(__nvvm_fmin_xorsign_abs_bf16, "yyy" , "" , AND(SM_86, PTX72)) |
193 | TARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_bf16, "yyy" , "" , |
194 | AND(SM_86, PTX72)) |
195 | TARGET_BUILTIN(__nvvm_fmin_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
196 | TARGET_BUILTIN(__nvvm_fmin_ftz_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
197 | TARGET_BUILTIN(__nvvm_fmin_nan_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
198 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
199 | TARGET_BUILTIN(__nvvm_fmin_xorsign_abs_bf16x2, "V2yV2yV2y" , "" , |
200 | AND(SM_86, PTX72)) |
201 | TARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_bf16x2, "V2yV2yV2y" , "" , |
202 | AND(SM_86, PTX72)) |
203 | BUILTIN(__nvvm_fmin_f, "fff" , "" ) |
204 | BUILTIN(__nvvm_fmin_ftz_f, "fff" , "" ) |
205 | TARGET_BUILTIN(__nvvm_fmin_nan_f, "fff" , "" , AND(SM_80, PTX70)) |
206 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_f, "fff" , "" , AND(SM_80, PTX70)) |
207 | TARGET_BUILTIN(__nvvm_fmin_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
208 | TARGET_BUILTIN(__nvvm_fmin_ftz_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
209 | TARGET_BUILTIN(__nvvm_fmin_nan_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
210 | TARGET_BUILTIN(__nvvm_fmin_ftz_nan_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
211 | BUILTIN(__nvvm_fmin_d, "ddd" , "" ) |
212 | |
213 | TARGET_BUILTIN(__nvvm_fmax_f16, "hhh" , "" , AND(SM_80, PTX70)) |
214 | TARGET_BUILTIN(__nvvm_fmax_ftz_f16, "hhh" , "" , AND(SM_80, PTX70)) |
215 | TARGET_BUILTIN(__nvvm_fmax_nan_f16, "hhh" , "" , AND(SM_80, PTX70)) |
216 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16, "hhh" , "" , AND(SM_80, PTX70)) |
217 | TARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
218 | TARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
219 | TARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f16, "hhh" , "" , AND(SM_86, PTX72)) |
220 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f16, "hhh" , "" , |
221 | AND(SM_86, PTX72)) |
222 | TARGET_BUILTIN(__nvvm_fmax_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
223 | TARGET_BUILTIN(__nvvm_fmax_ftz_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
224 | TARGET_BUILTIN(__nvvm_fmax_nan_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
225 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f16x2, "V2hV2hV2h" , "" , AND(SM_80, PTX70)) |
226 | TARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
227 | AND(SM_86, PTX72)) |
228 | TARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
229 | AND(SM_86, PTX72)) |
230 | TARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
231 | AND(SM_86, PTX72)) |
232 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f16x2, "V2hV2hV2h" , "" , |
233 | AND(SM_86, PTX72)) |
234 | TARGET_BUILTIN(__nvvm_fmax_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
235 | TARGET_BUILTIN(__nvvm_fmax_ftz_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
236 | TARGET_BUILTIN(__nvvm_fmax_nan_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
237 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_bf16, "yyy" , "" , AND(SM_80, PTX70)) |
238 | TARGET_BUILTIN(__nvvm_fmax_xorsign_abs_bf16, "yyy" , "" , AND(SM_86, PTX72)) |
239 | TARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_bf16, "yyy" , "" , |
240 | AND(SM_86, PTX72)) |
241 | TARGET_BUILTIN(__nvvm_fmax_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
242 | TARGET_BUILTIN(__nvvm_fmax_ftz_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
243 | TARGET_BUILTIN(__nvvm_fmax_nan_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
244 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_bf16x2, "V2yV2yV2y" , "" , AND(SM_80, PTX70)) |
245 | TARGET_BUILTIN(__nvvm_fmax_xorsign_abs_bf16x2, "V2yV2yV2y" , "" , |
246 | AND(SM_86, PTX72)) |
247 | TARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_bf16x2, "V2yV2yV2y" , "" , |
248 | AND(SM_86, PTX72)) |
249 | BUILTIN(__nvvm_fmax_f, "fff" , "" ) |
250 | BUILTIN(__nvvm_fmax_ftz_f, "fff" , "" ) |
251 | TARGET_BUILTIN(__nvvm_fmax_nan_f, "fff" , "" , AND(SM_80, PTX70)) |
252 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_f, "fff" , "" , AND(SM_80, PTX70)) |
253 | TARGET_BUILTIN(__nvvm_fmax_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
254 | TARGET_BUILTIN(__nvvm_fmax_ftz_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
255 | TARGET_BUILTIN(__nvvm_fmax_nan_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
256 | TARGET_BUILTIN(__nvvm_fmax_ftz_nan_xorsign_abs_f, "fff" , "" , AND(SM_86, PTX72)) |
257 | BUILTIN(__nvvm_fmax_d, "ddd" , "" ) |
258 | |
259 | // Multiplication |
260 | |
261 | BUILTIN(__nvvm_mulhi_i, "iii" , "" ) |
262 | BUILTIN(__nvvm_mulhi_ui, "UiUiUi" , "" ) |
263 | BUILTIN(__nvvm_mulhi_ll, "LLiLLiLLi" , "" ) |
264 | BUILTIN(__nvvm_mulhi_ull, "ULLiULLiULLi" , "" ) |
265 | |
266 | BUILTIN(__nvvm_mul_rn_ftz_f, "fff" , "" ) |
267 | BUILTIN(__nvvm_mul_rn_f, "fff" , "" ) |
268 | BUILTIN(__nvvm_mul_rz_ftz_f, "fff" , "" ) |
269 | BUILTIN(__nvvm_mul_rz_f, "fff" , "" ) |
270 | BUILTIN(__nvvm_mul_rm_ftz_f, "fff" , "" ) |
271 | BUILTIN(__nvvm_mul_rm_f, "fff" , "" ) |
272 | BUILTIN(__nvvm_mul_rp_ftz_f, "fff" , "" ) |
273 | BUILTIN(__nvvm_mul_rp_f, "fff" , "" ) |
274 | |
275 | BUILTIN(__nvvm_mul_rn_d, "ddd" , "" ) |
276 | BUILTIN(__nvvm_mul_rz_d, "ddd" , "" ) |
277 | BUILTIN(__nvvm_mul_rm_d, "ddd" , "" ) |
278 | BUILTIN(__nvvm_mul_rp_d, "ddd" , "" ) |
279 | |
280 | BUILTIN(__nvvm_mul24_i, "iii" , "" ) |
281 | BUILTIN(__nvvm_mul24_ui, "UiUiUi" , "" ) |
282 | |
283 | // Div |
284 | |
285 | BUILTIN(__nvvm_div_approx_ftz_f, "fff" , "" ) |
286 | BUILTIN(__nvvm_div_approx_f, "fff" , "" ) |
287 | |
288 | BUILTIN(__nvvm_div_rn_ftz_f, "fff" , "" ) |
289 | BUILTIN(__nvvm_div_rn_f, "fff" , "" ) |
290 | BUILTIN(__nvvm_div_rz_ftz_f, "fff" , "" ) |
291 | BUILTIN(__nvvm_div_rz_f, "fff" , "" ) |
292 | BUILTIN(__nvvm_div_rm_ftz_f, "fff" , "" ) |
293 | BUILTIN(__nvvm_div_rm_f, "fff" , "" ) |
294 | BUILTIN(__nvvm_div_rp_ftz_f, "fff" , "" ) |
295 | BUILTIN(__nvvm_div_rp_f, "fff" , "" ) |
296 | |
297 | BUILTIN(__nvvm_div_rn_d, "ddd" , "" ) |
298 | BUILTIN(__nvvm_div_rz_d, "ddd" , "" ) |
299 | BUILTIN(__nvvm_div_rm_d, "ddd" , "" ) |
300 | BUILTIN(__nvvm_div_rp_d, "ddd" , "" ) |
301 | |
302 | // Sad |
303 | |
304 | BUILTIN(__nvvm_sad_i, "iiii" , "" ) |
305 | BUILTIN(__nvvm_sad_ui, "UiUiUiUi" , "" ) |
306 | |
307 | // Floor, Ceil |
308 | |
309 | BUILTIN(__nvvm_floor_ftz_f, "ff" , "" ) |
310 | BUILTIN(__nvvm_floor_f, "ff" , "" ) |
311 | BUILTIN(__nvvm_floor_d, "dd" , "" ) |
312 | |
313 | BUILTIN(__nvvm_ceil_ftz_f, "ff" , "" ) |
314 | BUILTIN(__nvvm_ceil_f, "ff" , "" ) |
315 | BUILTIN(__nvvm_ceil_d, "dd" , "" ) |
316 | |
317 | // Abs |
318 | |
319 | BUILTIN(__nvvm_fabs_ftz_f, "ff" , "" ) |
320 | BUILTIN(__nvvm_fabs_f, "ff" , "" ) |
321 | BUILTIN(__nvvm_fabs_d, "dd" , "" ) |
322 | |
323 | // Round |
324 | |
325 | BUILTIN(__nvvm_round_ftz_f, "ff" , "" ) |
326 | BUILTIN(__nvvm_round_f, "ff" , "" ) |
327 | BUILTIN(__nvvm_round_d, "dd" , "" ) |
328 | |
329 | // Trunc |
330 | |
331 | BUILTIN(__nvvm_trunc_ftz_f, "ff" , "" ) |
332 | BUILTIN(__nvvm_trunc_f, "ff" , "" ) |
333 | BUILTIN(__nvvm_trunc_d, "dd" , "" ) |
334 | |
335 | // Saturate |
336 | |
337 | BUILTIN(__nvvm_saturate_ftz_f, "ff" , "" ) |
338 | BUILTIN(__nvvm_saturate_f, "ff" , "" ) |
339 | BUILTIN(__nvvm_saturate_d, "dd" , "" ) |
340 | |
341 | // Exp2, Log2 |
342 | |
343 | BUILTIN(__nvvm_ex2_approx_ftz_f, "ff" , "" ) |
344 | BUILTIN(__nvvm_ex2_approx_f, "ff" , "" ) |
345 | BUILTIN(__nvvm_ex2_approx_d, "dd" , "" ) |
346 | TARGET_BUILTIN(__nvvm_ex2_approx_f16, "hh" , "" , AND(SM_75, PTX70)) |
347 | TARGET_BUILTIN(__nvvm_ex2_approx_f16x2, "V2hV2h" , "" , AND(SM_75, PTX70)) |
348 | |
349 | BUILTIN(__nvvm_lg2_approx_ftz_f, "ff" , "" ) |
350 | BUILTIN(__nvvm_lg2_approx_f, "ff" , "" ) |
351 | BUILTIN(__nvvm_lg2_approx_d, "dd" , "" ) |
352 | |
353 | // Sin, Cos |
354 | |
355 | BUILTIN(__nvvm_sin_approx_ftz_f, "ff" , "" ) |
356 | BUILTIN(__nvvm_sin_approx_f, "ff" , "" ) |
357 | |
358 | BUILTIN(__nvvm_cos_approx_ftz_f, "ff" , "" ) |
359 | BUILTIN(__nvvm_cos_approx_f, "ff" , "" ) |
360 | |
361 | // Fma |
362 | |
363 | TARGET_BUILTIN(__nvvm_fma_rn_f16, "hhhh" , "" , AND(SM_53, PTX42)) |
364 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_f16, "hhhh" , "" , AND(SM_53, PTX42)) |
365 | TARGET_BUILTIN(__nvvm_fma_rn_sat_f16, "hhhh" , "" , AND(SM_53, PTX42)) |
366 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16, "hhhh" , "" , AND(SM_53, PTX42)) |
367 | TARGET_BUILTIN(__nvvm_fma_rn_relu_f16, "hhhh" , "" , AND(SM_80, PTX70)) |
368 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16, "hhhh" , "" , AND(SM_80, PTX70)) |
369 | TARGET_BUILTIN(__nvvm_fma_rn_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_53, PTX42)) |
370 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_53, PTX42)) |
371 | TARGET_BUILTIN(__nvvm_fma_rn_sat_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_53, PTX42)) |
372 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_sat_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_53, PTX42)) |
373 | TARGET_BUILTIN(__nvvm_fma_rn_relu_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_80, PTX70)) |
374 | TARGET_BUILTIN(__nvvm_fma_rn_ftz_relu_f16x2, "V2hV2hV2hV2h" , "" , AND(SM_80, PTX70)) |
375 | TARGET_BUILTIN(__nvvm_fma_rn_bf16, "yyyy" , "" , AND(SM_80, PTX70)) |
376 | TARGET_BUILTIN(__nvvm_fma_rn_relu_bf16, "yyyy" , "" , AND(SM_80, PTX70)) |
377 | TARGET_BUILTIN(__nvvm_fma_rn_bf16x2, "V2yV2yV2yV2y" , "" , AND(SM_80, PTX70)) |
378 | TARGET_BUILTIN(__nvvm_fma_rn_relu_bf16x2, "V2yV2yV2yV2y" , "" , AND(SM_80, PTX70)) |
379 | BUILTIN(__nvvm_fma_rn_ftz_f, "ffff" , "" ) |
380 | BUILTIN(__nvvm_fma_rn_f, "ffff" , "" ) |
381 | BUILTIN(__nvvm_fma_rz_ftz_f, "ffff" , "" ) |
382 | BUILTIN(__nvvm_fma_rz_f, "ffff" , "" ) |
383 | BUILTIN(__nvvm_fma_rm_ftz_f, "ffff" , "" ) |
384 | BUILTIN(__nvvm_fma_rm_f, "ffff" , "" ) |
385 | BUILTIN(__nvvm_fma_rp_ftz_f, "ffff" , "" ) |
386 | BUILTIN(__nvvm_fma_rp_f, "ffff" , "" ) |
387 | BUILTIN(__nvvm_fma_rn_d, "dddd" , "" ) |
388 | BUILTIN(__nvvm_fma_rz_d, "dddd" , "" ) |
389 | BUILTIN(__nvvm_fma_rm_d, "dddd" , "" ) |
390 | BUILTIN(__nvvm_fma_rp_d, "dddd" , "" ) |
391 | |
392 | // Rcp |
393 | |
394 | BUILTIN(__nvvm_rcp_rn_ftz_f, "ff" , "" ) |
395 | BUILTIN(__nvvm_rcp_rn_f, "ff" , "" ) |
396 | BUILTIN(__nvvm_rcp_rz_ftz_f, "ff" , "" ) |
397 | BUILTIN(__nvvm_rcp_rz_f, "ff" , "" ) |
398 | BUILTIN(__nvvm_rcp_rm_ftz_f, "ff" , "" ) |
399 | BUILTIN(__nvvm_rcp_rm_f, "ff" , "" ) |
400 | BUILTIN(__nvvm_rcp_rp_ftz_f, "ff" , "" ) |
401 | BUILTIN(__nvvm_rcp_rp_f, "ff" , "" ) |
402 | |
403 | BUILTIN(__nvvm_rcp_rn_d, "dd" , "" ) |
404 | BUILTIN(__nvvm_rcp_rz_d, "dd" , "" ) |
405 | BUILTIN(__nvvm_rcp_rm_d, "dd" , "" ) |
406 | BUILTIN(__nvvm_rcp_rp_d, "dd" , "" ) |
407 | |
408 | BUILTIN(__nvvm_rcp_approx_ftz_f, "ff" , "" ) |
409 | BUILTIN(__nvvm_rcp_approx_ftz_d, "dd" , "" ) |
410 | |
411 | // Sqrt |
412 | |
413 | BUILTIN(__nvvm_sqrt_rn_ftz_f, "ff" , "" ) |
414 | BUILTIN(__nvvm_sqrt_rn_f, "ff" , "" ) |
415 | BUILTIN(__nvvm_sqrt_rz_ftz_f, "ff" , "" ) |
416 | BUILTIN(__nvvm_sqrt_rz_f, "ff" , "" ) |
417 | BUILTIN(__nvvm_sqrt_rm_ftz_f, "ff" , "" ) |
418 | BUILTIN(__nvvm_sqrt_rm_f, "ff" , "" ) |
419 | BUILTIN(__nvvm_sqrt_rp_ftz_f, "ff" , "" ) |
420 | BUILTIN(__nvvm_sqrt_rp_f, "ff" , "" ) |
421 | BUILTIN(__nvvm_sqrt_approx_ftz_f, "ff" , "" ) |
422 | BUILTIN(__nvvm_sqrt_approx_f, "ff" , "" ) |
423 | |
424 | BUILTIN(__nvvm_sqrt_rn_d, "dd" , "" ) |
425 | BUILTIN(__nvvm_sqrt_rz_d, "dd" , "" ) |
426 | BUILTIN(__nvvm_sqrt_rm_d, "dd" , "" ) |
427 | BUILTIN(__nvvm_sqrt_rp_d, "dd" , "" ) |
428 | |
429 | // Rsqrt |
430 | |
431 | BUILTIN(__nvvm_rsqrt_approx_ftz_f, "ff" , "" ) |
432 | BUILTIN(__nvvm_rsqrt_approx_f, "ff" , "" ) |
433 | BUILTIN(__nvvm_rsqrt_approx_d, "dd" , "" ) |
434 | |
435 | // Add |
436 | |
437 | BUILTIN(__nvvm_add_rn_ftz_f, "fff" , "" ) |
438 | BUILTIN(__nvvm_add_rn_f, "fff" , "" ) |
439 | BUILTIN(__nvvm_add_rz_ftz_f, "fff" , "" ) |
440 | BUILTIN(__nvvm_add_rz_f, "fff" , "" ) |
441 | BUILTIN(__nvvm_add_rm_ftz_f, "fff" , "" ) |
442 | BUILTIN(__nvvm_add_rm_f, "fff" , "" ) |
443 | BUILTIN(__nvvm_add_rp_ftz_f, "fff" , "" ) |
444 | BUILTIN(__nvvm_add_rp_f, "fff" , "" ) |
445 | |
446 | BUILTIN(__nvvm_add_rn_d, "ddd" , "" ) |
447 | BUILTIN(__nvvm_add_rz_d, "ddd" , "" ) |
448 | BUILTIN(__nvvm_add_rm_d, "ddd" , "" ) |
449 | BUILTIN(__nvvm_add_rp_d, "ddd" , "" ) |
450 | |
451 | // Convert |
452 | |
453 | BUILTIN(__nvvm_d2f_rn_ftz, "fd" , "" ) |
454 | BUILTIN(__nvvm_d2f_rn, "fd" , "" ) |
455 | BUILTIN(__nvvm_d2f_rz_ftz, "fd" , "" ) |
456 | BUILTIN(__nvvm_d2f_rz, "fd" , "" ) |
457 | BUILTIN(__nvvm_d2f_rm_ftz, "fd" , "" ) |
458 | BUILTIN(__nvvm_d2f_rm, "fd" , "" ) |
459 | BUILTIN(__nvvm_d2f_rp_ftz, "fd" , "" ) |
460 | BUILTIN(__nvvm_d2f_rp, "fd" , "" ) |
461 | |
462 | BUILTIN(__nvvm_d2i_rn, "id" , "" ) |
463 | BUILTIN(__nvvm_d2i_rz, "id" , "" ) |
464 | BUILTIN(__nvvm_d2i_rm, "id" , "" ) |
465 | BUILTIN(__nvvm_d2i_rp, "id" , "" ) |
466 | |
467 | BUILTIN(__nvvm_d2ui_rn, "Uid" , "" ) |
468 | BUILTIN(__nvvm_d2ui_rz, "Uid" , "" ) |
469 | BUILTIN(__nvvm_d2ui_rm, "Uid" , "" ) |
470 | BUILTIN(__nvvm_d2ui_rp, "Uid" , "" ) |
471 | |
472 | BUILTIN(__nvvm_i2d_rn, "di" , "" ) |
473 | BUILTIN(__nvvm_i2d_rz, "di" , "" ) |
474 | BUILTIN(__nvvm_i2d_rm, "di" , "" ) |
475 | BUILTIN(__nvvm_i2d_rp, "di" , "" ) |
476 | |
477 | BUILTIN(__nvvm_ui2d_rn, "dUi" , "" ) |
478 | BUILTIN(__nvvm_ui2d_rz, "dUi" , "" ) |
479 | BUILTIN(__nvvm_ui2d_rm, "dUi" , "" ) |
480 | BUILTIN(__nvvm_ui2d_rp, "dUi" , "" ) |
481 | |
482 | BUILTIN(__nvvm_f2i_rn_ftz, "if" , "" ) |
483 | BUILTIN(__nvvm_f2i_rn, "if" , "" ) |
484 | BUILTIN(__nvvm_f2i_rz_ftz, "if" , "" ) |
485 | BUILTIN(__nvvm_f2i_rz, "if" , "" ) |
486 | BUILTIN(__nvvm_f2i_rm_ftz, "if" , "" ) |
487 | BUILTIN(__nvvm_f2i_rm, "if" , "" ) |
488 | BUILTIN(__nvvm_f2i_rp_ftz, "if" , "" ) |
489 | BUILTIN(__nvvm_f2i_rp, "if" , "" ) |
490 | |
491 | BUILTIN(__nvvm_f2ui_rn_ftz, "Uif" , "" ) |
492 | BUILTIN(__nvvm_f2ui_rn, "Uif" , "" ) |
493 | BUILTIN(__nvvm_f2ui_rz_ftz, "Uif" , "" ) |
494 | BUILTIN(__nvvm_f2ui_rz, "Uif" , "" ) |
495 | BUILTIN(__nvvm_f2ui_rm_ftz, "Uif" , "" ) |
496 | BUILTIN(__nvvm_f2ui_rm, "Uif" , "" ) |
497 | BUILTIN(__nvvm_f2ui_rp_ftz, "Uif" , "" ) |
498 | BUILTIN(__nvvm_f2ui_rp, "Uif" , "" ) |
499 | |
500 | BUILTIN(__nvvm_i2f_rn, "fi" , "" ) |
501 | BUILTIN(__nvvm_i2f_rz, "fi" , "" ) |
502 | BUILTIN(__nvvm_i2f_rm, "fi" , "" ) |
503 | BUILTIN(__nvvm_i2f_rp, "fi" , "" ) |
504 | |
505 | BUILTIN(__nvvm_ui2f_rn, "fUi" , "" ) |
506 | BUILTIN(__nvvm_ui2f_rz, "fUi" , "" ) |
507 | BUILTIN(__nvvm_ui2f_rm, "fUi" , "" ) |
508 | BUILTIN(__nvvm_ui2f_rp, "fUi" , "" ) |
509 | |
510 | BUILTIN(__nvvm_lohi_i2d, "dii" , "" ) |
511 | |
512 | BUILTIN(__nvvm_d2i_lo, "id" , "" ) |
513 | BUILTIN(__nvvm_d2i_hi, "id" , "" ) |
514 | |
515 | BUILTIN(__nvvm_f2ll_rn_ftz, "LLif" , "" ) |
516 | BUILTIN(__nvvm_f2ll_rn, "LLif" , "" ) |
517 | BUILTIN(__nvvm_f2ll_rz_ftz, "LLif" , "" ) |
518 | BUILTIN(__nvvm_f2ll_rz, "LLif" , "" ) |
519 | BUILTIN(__nvvm_f2ll_rm_ftz, "LLif" , "" ) |
520 | BUILTIN(__nvvm_f2ll_rm, "LLif" , "" ) |
521 | BUILTIN(__nvvm_f2ll_rp_ftz, "LLif" , "" ) |
522 | BUILTIN(__nvvm_f2ll_rp, "LLif" , "" ) |
523 | |
524 | BUILTIN(__nvvm_f2ull_rn_ftz, "ULLif" , "" ) |
525 | BUILTIN(__nvvm_f2ull_rn, "ULLif" , "" ) |
526 | BUILTIN(__nvvm_f2ull_rz_ftz, "ULLif" , "" ) |
527 | BUILTIN(__nvvm_f2ull_rz, "ULLif" , "" ) |
528 | BUILTIN(__nvvm_f2ull_rm_ftz, "ULLif" , "" ) |
529 | BUILTIN(__nvvm_f2ull_rm, "ULLif" , "" ) |
530 | BUILTIN(__nvvm_f2ull_rp_ftz, "ULLif" , "" ) |
531 | BUILTIN(__nvvm_f2ull_rp, "ULLif" , "" ) |
532 | |
533 | BUILTIN(__nvvm_d2ll_rn, "LLid" , "" ) |
534 | BUILTIN(__nvvm_d2ll_rz, "LLid" , "" ) |
535 | BUILTIN(__nvvm_d2ll_rm, "LLid" , "" ) |
536 | BUILTIN(__nvvm_d2ll_rp, "LLid" , "" ) |
537 | |
538 | BUILTIN(__nvvm_d2ull_rn, "ULLid" , "" ) |
539 | BUILTIN(__nvvm_d2ull_rz, "ULLid" , "" ) |
540 | BUILTIN(__nvvm_d2ull_rm, "ULLid" , "" ) |
541 | BUILTIN(__nvvm_d2ull_rp, "ULLid" , "" ) |
542 | |
543 | BUILTIN(__nvvm_ll2f_rn, "fLLi" , "" ) |
544 | BUILTIN(__nvvm_ll2f_rz, "fLLi" , "" ) |
545 | BUILTIN(__nvvm_ll2f_rm, "fLLi" , "" ) |
546 | BUILTIN(__nvvm_ll2f_rp, "fLLi" , "" ) |
547 | |
548 | BUILTIN(__nvvm_ull2f_rn, "fULLi" , "" ) |
549 | BUILTIN(__nvvm_ull2f_rz, "fULLi" , "" ) |
550 | BUILTIN(__nvvm_ull2f_rm, "fULLi" , "" ) |
551 | BUILTIN(__nvvm_ull2f_rp, "fULLi" , "" ) |
552 | |
553 | BUILTIN(__nvvm_ll2d_rn, "dLLi" , "" ) |
554 | BUILTIN(__nvvm_ll2d_rz, "dLLi" , "" ) |
555 | BUILTIN(__nvvm_ll2d_rm, "dLLi" , "" ) |
556 | BUILTIN(__nvvm_ll2d_rp, "dLLi" , "" ) |
557 | |
558 | BUILTIN(__nvvm_ull2d_rn, "dULLi" , "" ) |
559 | BUILTIN(__nvvm_ull2d_rz, "dULLi" , "" ) |
560 | BUILTIN(__nvvm_ull2d_rm, "dULLi" , "" ) |
561 | BUILTIN(__nvvm_ull2d_rp, "dULLi" , "" ) |
562 | |
563 | BUILTIN(__nvvm_f2h_rn_ftz, "Usf" , "" ) |
564 | BUILTIN(__nvvm_f2h_rn, "Usf" , "" ) |
565 | |
566 | TARGET_BUILTIN(__nvvm_ff2bf16x2_rn, "V2yff" , "" , AND(SM_80,PTX70)) |
567 | TARGET_BUILTIN(__nvvm_ff2bf16x2_rn_relu, "V2yff" , "" , AND(SM_80,PTX70)) |
568 | TARGET_BUILTIN(__nvvm_ff2bf16x2_rz, "V2yff" , "" , AND(SM_80,PTX70)) |
569 | TARGET_BUILTIN(__nvvm_ff2bf16x2_rz_relu, "V2yff" , "" , AND(SM_80,PTX70)) |
570 | |
571 | TARGET_BUILTIN(__nvvm_ff2f16x2_rn, "V2hff" , "" , AND(SM_80,PTX70)) |
572 | TARGET_BUILTIN(__nvvm_ff2f16x2_rn_relu, "V2hff" , "" , AND(SM_80,PTX70)) |
573 | TARGET_BUILTIN(__nvvm_ff2f16x2_rz, "V2hff" , "" , AND(SM_80,PTX70)) |
574 | TARGET_BUILTIN(__nvvm_ff2f16x2_rz_relu, "V2hff" , "" , AND(SM_80,PTX70)) |
575 | |
576 | TARGET_BUILTIN(__nvvm_f2bf16_rn, "yf" , "" , AND(SM_80,PTX70)) |
577 | TARGET_BUILTIN(__nvvm_f2bf16_rn_relu, "yf" , "" , AND(SM_80,PTX70)) |
578 | TARGET_BUILTIN(__nvvm_f2bf16_rz, "yf" , "" , AND(SM_80,PTX70)) |
579 | TARGET_BUILTIN(__nvvm_f2bf16_rz_relu, "yf" , "" , AND(SM_80,PTX70)) |
580 | |
581 | TARGET_BUILTIN(__nvvm_f2tf32_rna, "ZUif" , "" , AND(SM_80,PTX70)) |
582 | |
583 | // Bitcast |
584 | |
585 | BUILTIN(__nvvm_bitcast_f2i, "if" , "" ) |
586 | BUILTIN(__nvvm_bitcast_i2f, "fi" , "" ) |
587 | |
588 | BUILTIN(__nvvm_bitcast_ll2d, "dLLi" , "" ) |
589 | BUILTIN(__nvvm_bitcast_d2ll, "LLid" , "" ) |
590 | |
591 | // FNS |
592 | TARGET_BUILTIN(__nvvm_fns, "UiUiUii" , "n" , PTX60) |
593 | |
594 | // Sync |
595 | |
596 | BUILTIN(__syncthreads, "v" , "" ) |
597 | BUILTIN(__nvvm_bar0_popc, "ii" , "" ) |
598 | BUILTIN(__nvvm_bar0_and, "ii" , "" ) |
599 | BUILTIN(__nvvm_bar0_or, "ii" , "" ) |
600 | BUILTIN(__nvvm_bar_sync, "vi" , "n" ) |
601 | TARGET_BUILTIN(__nvvm_bar_warp_sync, "vUi" , "n" , PTX60) |
602 | TARGET_BUILTIN(__nvvm_barrier_sync, "vUi" , "n" , PTX60) |
603 | TARGET_BUILTIN(__nvvm_barrier_sync_cnt, "vUiUi" , "n" , PTX60) |
604 | |
605 | TARGET_BUILTIN(__nvvm_barrier_cluster_arrive, "v" , "n" , AND(SM_90,PTX78)) |
606 | TARGET_BUILTIN(__nvvm_barrier_cluster_arrive_relaxed, "v" , "n" , AND(SM_90,PTX80)) |
607 | TARGET_BUILTIN(__nvvm_barrier_cluster_wait, "v" , "n" , AND(SM_90,PTX78)) |
608 | TARGET_BUILTIN(__nvvm_fence_sc_cluster, "v" , "n" , AND(SM_90,PTX78)) |
609 | |
610 | // Shuffle |
611 | |
612 | BUILTIN(__nvvm_shfl_down_i32, "iiii" , "" ) |
613 | BUILTIN(__nvvm_shfl_down_f32, "ffii" , "" ) |
614 | BUILTIN(__nvvm_shfl_up_i32, "iiii" , "" ) |
615 | BUILTIN(__nvvm_shfl_up_f32, "ffii" , "" ) |
616 | BUILTIN(__nvvm_shfl_bfly_i32, "iiii" , "" ) |
617 | BUILTIN(__nvvm_shfl_bfly_f32, "ffii" , "" ) |
618 | BUILTIN(__nvvm_shfl_idx_i32, "iiii" , "" ) |
619 | BUILTIN(__nvvm_shfl_idx_f32, "ffii" , "" ) |
620 | |
621 | TARGET_BUILTIN(__nvvm_shfl_sync_down_i32, "iUiiii" , "" , PTX60) |
622 | TARGET_BUILTIN(__nvvm_shfl_sync_down_f32, "fUifii" , "" , PTX60) |
623 | TARGET_BUILTIN(__nvvm_shfl_sync_up_i32, "iUiiii" , "" , PTX60) |
624 | TARGET_BUILTIN(__nvvm_shfl_sync_up_f32, "fUifii" , "" , PTX60) |
625 | TARGET_BUILTIN(__nvvm_shfl_sync_bfly_i32, "iUiiii" , "" , PTX60) |
626 | TARGET_BUILTIN(__nvvm_shfl_sync_bfly_f32, "fUifii" , "" , PTX60) |
627 | TARGET_BUILTIN(__nvvm_shfl_sync_idx_i32, "iUiiii" , "" , PTX60) |
628 | TARGET_BUILTIN(__nvvm_shfl_sync_idx_f32, "fUifii" , "" , PTX60) |
629 | |
630 | // Vote |
631 | BUILTIN(__nvvm_vote_all, "bb" , "" ) |
632 | BUILTIN(__nvvm_vote_any, "bb" , "" ) |
633 | BUILTIN(__nvvm_vote_uni, "bb" , "" ) |
634 | BUILTIN(__nvvm_vote_ballot, "Uib" , "" ) |
635 | |
636 | TARGET_BUILTIN(__nvvm_vote_all_sync, "bUib" , "" , PTX60) |
637 | TARGET_BUILTIN(__nvvm_vote_any_sync, "bUib" , "" , PTX60) |
638 | TARGET_BUILTIN(__nvvm_vote_uni_sync, "bUib" , "" , PTX60) |
639 | TARGET_BUILTIN(__nvvm_vote_ballot_sync, "UiUib" , "" , PTX60) |
640 | |
641 | // Mask |
642 | TARGET_BUILTIN(__nvvm_activemask, "Ui" , "n" , PTX62) |
643 | |
644 | // Match |
645 | TARGET_BUILTIN(__nvvm_match_any_sync_i32, "UiUiUi" , "" , AND(SM_70,PTX60)) |
646 | TARGET_BUILTIN(__nvvm_match_any_sync_i64, "UiUiWi" , "" , AND(SM_70,PTX60)) |
647 | // These return a pair {value, predicate}, which requires custom lowering. |
648 | TARGET_BUILTIN(__nvvm_match_all_sync_i32p, "UiUiUii*" , "" , AND(SM_70,PTX60)) |
649 | TARGET_BUILTIN(__nvvm_match_all_sync_i64p, "UiUiWii*" , "" , AND(SM_70,PTX60)) |
650 | |
651 | // Redux |
652 | TARGET_BUILTIN(__nvvm_redux_sync_add, "iii" , "" , AND(SM_80,PTX70)) |
653 | TARGET_BUILTIN(__nvvm_redux_sync_min, "iii" , "" , AND(SM_80,PTX70)) |
654 | TARGET_BUILTIN(__nvvm_redux_sync_max, "iii" , "" , AND(SM_80,PTX70)) |
655 | TARGET_BUILTIN(__nvvm_redux_sync_umin, "UiUii" , "" , AND(SM_80,PTX70)) |
656 | TARGET_BUILTIN(__nvvm_redux_sync_umax, "UiUii" , "" , AND(SM_80,PTX70)) |
657 | TARGET_BUILTIN(__nvvm_redux_sync_and, "iii" , "" , AND(SM_80,PTX70)) |
658 | TARGET_BUILTIN(__nvvm_redux_sync_xor, "iii" , "" , AND(SM_80,PTX70)) |
659 | TARGET_BUILTIN(__nvvm_redux_sync_or, "iii" , "" , AND(SM_80,PTX70)) |
660 | |
661 | // Membar |
662 | |
663 | BUILTIN(__nvvm_membar_cta, "v" , "" ) |
664 | BUILTIN(__nvvm_membar_gl, "v" , "" ) |
665 | BUILTIN(__nvvm_membar_sys, "v" , "" ) |
666 | |
667 | // mbarrier |
668 | |
669 | TARGET_BUILTIN(__nvvm_mbarrier_init, "vWi*i" , "" , AND(SM_80,PTX70)) |
670 | TARGET_BUILTIN(__nvvm_mbarrier_init_shared, "vWi*3i" , "" , AND(SM_80,PTX70)) |
671 | |
672 | TARGET_BUILTIN(__nvvm_mbarrier_inval, "vWi*" , "" , AND(SM_80,PTX70)) |
673 | TARGET_BUILTIN(__nvvm_mbarrier_inval_shared, "vWi*3" , "" , AND(SM_80,PTX70)) |
674 | |
675 | TARGET_BUILTIN(__nvvm_mbarrier_arrive, "WiWi*" , "" , AND(SM_80,PTX70)) |
676 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_shared, "WiWi*3" , "" , AND(SM_80,PTX70)) |
677 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete, "WiWi*i" , "" , AND(SM_80,PTX70)) |
678 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_noComplete_shared, "WiWi*3i" , "" , AND(SM_80,PTX70)) |
679 | |
680 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop, "WiWi*" , "" , AND(SM_80,PTX70)) |
681 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_shared, "WiWi*3" , "" , AND(SM_80,PTX70)) |
682 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete, "WiWi*i" , "" , AND(SM_80,PTX70)) |
683 | TARGET_BUILTIN(__nvvm_mbarrier_arrive_drop_noComplete_shared, "WiWi*3i" , "" , AND(SM_80,PTX70)) |
684 | |
685 | TARGET_BUILTIN(__nvvm_mbarrier_test_wait, "bWi*Wi" , "" , AND(SM_80,PTX70)) |
686 | TARGET_BUILTIN(__nvvm_mbarrier_test_wait_shared, "bWi*3Wi" , "" , AND(SM_80,PTX70)) |
687 | |
688 | TARGET_BUILTIN(__nvvm_mbarrier_pending_count, "iWi" , "" , AND(SM_80,PTX70)) |
689 | |
690 | // Memcpy, Memset |
691 | |
692 | BUILTIN(__nvvm_memcpy, "vUc*Uc*zi" ,"" ) |
693 | BUILTIN(__nvvm_memset, "vUc*Uczi" ,"" ) |
694 | |
695 | // Image |
696 | |
697 | BUILTIN(__builtin_ptx_read_image2Dfi_, "V4fiiii" , "" ) |
698 | BUILTIN(__builtin_ptx_read_image2Dff_, "V4fiiff" , "" ) |
699 | BUILTIN(__builtin_ptx_read_image2Dii_, "V4iiiii" , "" ) |
700 | BUILTIN(__builtin_ptx_read_image2Dif_, "V4iiiff" , "" ) |
701 | |
702 | BUILTIN(__builtin_ptx_read_image3Dfi_, "V4fiiiiii" , "" ) |
703 | BUILTIN(__builtin_ptx_read_image3Dff_, "V4fiiffff" , "" ) |
704 | BUILTIN(__builtin_ptx_read_image3Dii_, "V4iiiiiii" , "" ) |
705 | BUILTIN(__builtin_ptx_read_image3Dif_, "V4iiiffff" , "" ) |
706 | |
707 | BUILTIN(__builtin_ptx_write_image2Df_, "viiiffff" , "" ) |
708 | BUILTIN(__builtin_ptx_write_image2Di_, "viiiiiii" , "" ) |
709 | BUILTIN(__builtin_ptx_write_image2Dui_, "viiiUiUiUiUi" , "" ) |
710 | BUILTIN(__builtin_ptx_get_image_depthi_, "ii" , "" ) |
711 | BUILTIN(__builtin_ptx_get_image_heighti_, "ii" , "" ) |
712 | BUILTIN(__builtin_ptx_get_image_widthi_, "ii" , "" ) |
713 | BUILTIN(__builtin_ptx_get_image_channel_data_typei_, "ii" , "" ) |
714 | BUILTIN(__builtin_ptx_get_image_channel_orderi_, "ii" , "" ) |
715 | |
716 | // Atomic |
717 | // |
718 | // We need the atom intrinsics because |
719 | // - they are used in converging analysis |
720 | // - they are used in address space analysis and optimization |
721 | // So it does not hurt to expose them as builtins. |
722 | // |
723 | BUILTIN(__nvvm_atom_add_gen_i, "iiD*i" , "n" ) |
724 | TARGET_BUILTIN(__nvvm_atom_cta_add_gen_i, "iiD*i" , "n" , SM_60) |
725 | TARGET_BUILTIN(__nvvm_atom_sys_add_gen_i, "iiD*i" , "n" , SM_60) |
726 | BUILTIN(__nvvm_atom_add_gen_l, "LiLiD*Li" , "n" ) |
727 | TARGET_BUILTIN(__nvvm_atom_cta_add_gen_l, "LiLiD*Li" , "n" , SM_60) |
728 | TARGET_BUILTIN(__nvvm_atom_sys_add_gen_l, "LiLiD*Li" , "n" , SM_60) |
729 | BUILTIN(__nvvm_atom_add_gen_ll, "LLiLLiD*LLi" , "n" ) |
730 | TARGET_BUILTIN(__nvvm_atom_cta_add_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
731 | TARGET_BUILTIN(__nvvm_atom_sys_add_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
732 | BUILTIN(__nvvm_atom_add_gen_f, "ffD*f" , "n" ) |
733 | TARGET_BUILTIN(__nvvm_atom_cta_add_gen_f, "ffD*f" , "n" , SM_60) |
734 | TARGET_BUILTIN(__nvvm_atom_sys_add_gen_f, "ffD*f" , "n" , SM_60) |
735 | TARGET_BUILTIN(__nvvm_atom_add_gen_d, "ddD*d" , "n" , SM_60) |
736 | TARGET_BUILTIN(__nvvm_atom_cta_add_gen_d, "ddD*d" , "n" , SM_60) |
737 | TARGET_BUILTIN(__nvvm_atom_sys_add_gen_d, "ddD*d" , "n" , SM_60) |
738 | |
739 | BUILTIN(__nvvm_atom_sub_gen_i, "iiD*i" , "n" ) |
740 | BUILTIN(__nvvm_atom_sub_gen_l, "LiLiD*Li" , "n" ) |
741 | BUILTIN(__nvvm_atom_sub_gen_ll, "LLiLLiD*LLi" , "n" ) |
742 | |
743 | BUILTIN(__nvvm_atom_xchg_gen_i, "iiD*i" , "n" ) |
744 | TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_i, "iiD*i" , "n" , SM_60) |
745 | TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_i, "iiD*i" , "n" , SM_60) |
746 | BUILTIN(__nvvm_atom_xchg_gen_l, "LiLiD*Li" , "n" ) |
747 | TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_l, "LiLiD*Li" , "n" , SM_60) |
748 | TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_l, "LiLiD*Li" , "n" , SM_60) |
749 | BUILTIN(__nvvm_atom_xchg_gen_ll, "LLiLLiD*LLi" , "n" ) |
750 | TARGET_BUILTIN(__nvvm_atom_cta_xchg_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
751 | TARGET_BUILTIN(__nvvm_atom_sys_xchg_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
752 | |
753 | BUILTIN(__nvvm_atom_max_gen_i, "iiD*i" , "n" ) |
754 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_i, "iiD*i" , "n" , SM_60) |
755 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_i, "iiD*i" , "n" , SM_60) |
756 | BUILTIN(__nvvm_atom_max_gen_ui, "UiUiD*Ui" , "n" ) |
757 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
758 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
759 | BUILTIN(__nvvm_atom_max_gen_l, "LiLiD*Li" , "n" ) |
760 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_l, "LiLiD*Li" , "n" , SM_60) |
761 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_l, "LiLiD*Li" , "n" , SM_60) |
762 | BUILTIN(__nvvm_atom_max_gen_ul, "ULiULiD*ULi" , "n" ) |
763 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ul, "ULiULiD*ULi" , "n" , SM_60) |
764 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ul, "ULiULiD*ULi" , "n" , SM_60) |
765 | BUILTIN(__nvvm_atom_max_gen_ll, "LLiLLiD*LLi" , "n" ) |
766 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
767 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
768 | BUILTIN(__nvvm_atom_max_gen_ull, "ULLiULLiD*ULLi" , "n" ) |
769 | TARGET_BUILTIN(__nvvm_atom_cta_max_gen_ull, "ULLiULLiD*ULLi" , "n" , SM_60) |
770 | TARGET_BUILTIN(__nvvm_atom_sys_max_gen_ull, "ULLiULLiD*ULLi" , "n" , SM_60) |
771 | |
772 | BUILTIN(__nvvm_atom_min_gen_i, "iiD*i" , "n" ) |
773 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_i, "iiD*i" , "n" , SM_60) |
774 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_i, "iiD*i" , "n" , SM_60) |
775 | BUILTIN(__nvvm_atom_min_gen_ui, "UiUiD*Ui" , "n" ) |
776 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
777 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
778 | BUILTIN(__nvvm_atom_min_gen_l, "LiLiD*Li" , "n" ) |
779 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_l, "LiLiD*Li" , "n" , SM_60) |
780 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_l, "LiLiD*Li" , "n" , SM_60) |
781 | BUILTIN(__nvvm_atom_min_gen_ul, "ULiULiD*ULi" , "n" ) |
782 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ul, "ULiULiD*ULi" , "n" , SM_60) |
783 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ul, "ULiULiD*ULi" , "n" , SM_60) |
784 | BUILTIN(__nvvm_atom_min_gen_ll, "LLiLLiD*LLi" , "n" ) |
785 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
786 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
787 | BUILTIN(__nvvm_atom_min_gen_ull, "ULLiULLiD*ULLi" , "n" ) |
788 | TARGET_BUILTIN(__nvvm_atom_cta_min_gen_ull, "ULLiULLiD*ULLi" , "n" , SM_60) |
789 | TARGET_BUILTIN(__nvvm_atom_sys_min_gen_ull, "ULLiULLiD*ULLi" , "n" , SM_60) |
790 | |
791 | BUILTIN(__nvvm_atom_inc_gen_ui, "UiUiD*Ui" , "n" ) |
792 | TARGET_BUILTIN(__nvvm_atom_cta_inc_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
793 | TARGET_BUILTIN(__nvvm_atom_sys_inc_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
794 | BUILTIN(__nvvm_atom_dec_gen_ui, "UiUiD*Ui" , "n" ) |
795 | TARGET_BUILTIN(__nvvm_atom_cta_dec_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
796 | TARGET_BUILTIN(__nvvm_atom_sys_dec_gen_ui, "UiUiD*Ui" , "n" , SM_60) |
797 | |
798 | BUILTIN(__nvvm_atom_and_gen_i, "iiD*i" , "n" ) |
799 | TARGET_BUILTIN(__nvvm_atom_cta_and_gen_i, "iiD*i" , "n" , SM_60) |
800 | TARGET_BUILTIN(__nvvm_atom_sys_and_gen_i, "iiD*i" , "n" , SM_60) |
801 | BUILTIN(__nvvm_atom_and_gen_l, "LiLiD*Li" , "n" ) |
802 | TARGET_BUILTIN(__nvvm_atom_cta_and_gen_l, "LiLiD*Li" , "n" , SM_60) |
803 | TARGET_BUILTIN(__nvvm_atom_sys_and_gen_l, "LiLiD*Li" , "n" , SM_60) |
804 | BUILTIN(__nvvm_atom_and_gen_ll, "LLiLLiD*LLi" , "n" ) |
805 | TARGET_BUILTIN(__nvvm_atom_cta_and_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
806 | TARGET_BUILTIN(__nvvm_atom_sys_and_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
807 | |
808 | BUILTIN(__nvvm_atom_or_gen_i, "iiD*i" , "n" ) |
809 | TARGET_BUILTIN(__nvvm_atom_cta_or_gen_i, "iiD*i" , "n" , SM_60) |
810 | TARGET_BUILTIN(__nvvm_atom_sys_or_gen_i, "iiD*i" , "n" , SM_60) |
811 | BUILTIN(__nvvm_atom_or_gen_l, "LiLiD*Li" , "n" ) |
812 | TARGET_BUILTIN(__nvvm_atom_cta_or_gen_l, "LiLiD*Li" , "n" , SM_60) |
813 | TARGET_BUILTIN(__nvvm_atom_sys_or_gen_l, "LiLiD*Li" , "n" , SM_60) |
814 | BUILTIN(__nvvm_atom_or_gen_ll, "LLiLLiD*LLi" , "n" ) |
815 | TARGET_BUILTIN(__nvvm_atom_cta_or_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
816 | TARGET_BUILTIN(__nvvm_atom_sys_or_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
817 | |
818 | BUILTIN(__nvvm_atom_xor_gen_i, "iiD*i" , "n" ) |
819 | TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_i, "iiD*i" , "n" , SM_60) |
820 | TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_i, "iiD*i" , "n" , SM_60) |
821 | BUILTIN(__nvvm_atom_xor_gen_l, "LiLiD*Li" , "n" ) |
822 | TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_l, "LiLiD*Li" , "n" , SM_60) |
823 | TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_l, "LiLiD*Li" , "n" , SM_60) |
824 | BUILTIN(__nvvm_atom_xor_gen_ll, "LLiLLiD*LLi" , "n" ) |
825 | TARGET_BUILTIN(__nvvm_atom_cta_xor_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
826 | TARGET_BUILTIN(__nvvm_atom_sys_xor_gen_ll, "LLiLLiD*LLi" , "n" , SM_60) |
827 | |
828 | BUILTIN(__nvvm_atom_cas_gen_i, "iiD*ii" , "n" ) |
829 | TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_i, "iiD*ii" , "n" , SM_60) |
830 | TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_i, "iiD*ii" , "n" , SM_60) |
831 | BUILTIN(__nvvm_atom_cas_gen_l, "LiLiD*LiLi" , "n" ) |
832 | TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_l, "LiLiD*LiLi" , "n" , SM_60) |
833 | TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_l, "LiLiD*LiLi" , "n" , SM_60) |
834 | BUILTIN(__nvvm_atom_cas_gen_ll, "LLiLLiD*LLiLLi" , "n" ) |
835 | TARGET_BUILTIN(__nvvm_atom_cta_cas_gen_ll, "LLiLLiD*LLiLLi" , "n" , SM_60) |
836 | TARGET_BUILTIN(__nvvm_atom_sys_cas_gen_ll, "LLiLLiD*LLiLLi" , "n" , SM_60) |
837 | |
838 | // Compiler Error Warn |
839 | BUILTIN(__nvvm_compiler_error, "vcC*4" , "n" ) |
840 | BUILTIN(__nvvm_compiler_warn, "vcC*4" , "n" ) |
841 | |
842 | BUILTIN(__nvvm_ldu_c, "ccC*" , "" ) |
843 | BUILTIN(__nvvm_ldu_sc, "ScScC*" , "" ) |
844 | BUILTIN(__nvvm_ldu_s, "ssC*" , "" ) |
845 | BUILTIN(__nvvm_ldu_i, "iiC*" , "" ) |
846 | BUILTIN(__nvvm_ldu_l, "LiLiC*" , "" ) |
847 | BUILTIN(__nvvm_ldu_ll, "LLiLLiC*" , "" ) |
848 | |
849 | BUILTIN(__nvvm_ldu_uc, "UcUcC*" , "" ) |
850 | BUILTIN(__nvvm_ldu_us, "UsUsC*" , "" ) |
851 | BUILTIN(__nvvm_ldu_ui, "UiUiC*" , "" ) |
852 | BUILTIN(__nvvm_ldu_ul, "ULiULiC*" , "" ) |
853 | BUILTIN(__nvvm_ldu_ull, "ULLiULLiC*" , "" ) |
854 | |
855 | BUILTIN(__nvvm_ldu_h, "hhC*" , "" ) |
856 | BUILTIN(__nvvm_ldu_f, "ffC*" , "" ) |
857 | BUILTIN(__nvvm_ldu_d, "ddC*" , "" ) |
858 | |
859 | BUILTIN(__nvvm_ldu_c2, "E2cE2cC*" , "" ) |
860 | BUILTIN(__nvvm_ldu_sc2, "E2ScE2ScC*" , "" ) |
861 | BUILTIN(__nvvm_ldu_c4, "E4cE4cC*" , "" ) |
862 | BUILTIN(__nvvm_ldu_sc4, "E4ScE4ScC*" , "" ) |
863 | BUILTIN(__nvvm_ldu_s2, "E2sE2sC*" , "" ) |
864 | BUILTIN(__nvvm_ldu_s4, "E4sE4sC*" , "" ) |
865 | BUILTIN(__nvvm_ldu_i2, "E2iE2iC*" , "" ) |
866 | BUILTIN(__nvvm_ldu_i4, "E4iE4iC*" , "" ) |
867 | BUILTIN(__nvvm_ldu_l2, "E2LiE2LiC*" , "" ) |
868 | BUILTIN(__nvvm_ldu_ll2, "E2LLiE2LLiC*" , "" ) |
869 | |
870 | BUILTIN(__nvvm_ldu_uc2, "E2UcE2UcC*" , "" ) |
871 | BUILTIN(__nvvm_ldu_uc4, "E4UcE4UcC*" , "" ) |
872 | BUILTIN(__nvvm_ldu_us2, "E2UsE2UsC*" , "" ) |
873 | BUILTIN(__nvvm_ldu_us4, "E4UsE4UsC*" , "" ) |
874 | BUILTIN(__nvvm_ldu_ui2, "E2UiE2UiC*" , "" ) |
875 | BUILTIN(__nvvm_ldu_ui4, "E4UiE4UiC*" , "" ) |
876 | BUILTIN(__nvvm_ldu_ul2, "E2ULiE2ULiC*" , "" ) |
877 | BUILTIN(__nvvm_ldu_ull2, "E2ULLiE2ULLiC*" , "" ) |
878 | |
879 | BUILTIN(__nvvm_ldu_h2, "E2hE2hC*" , "" ) |
880 | BUILTIN(__nvvm_ldu_f2, "E2fE2fC*" , "" ) |
881 | BUILTIN(__nvvm_ldu_f4, "E4fE4fC*" , "" ) |
882 | BUILTIN(__nvvm_ldu_d2, "E2dE2dC*" , "" ) |
883 | |
884 | BUILTIN(__nvvm_ldg_c, "ccC*" , "" ) |
885 | BUILTIN(__nvvm_ldg_sc, "ScScC*" , "" ) |
886 | BUILTIN(__nvvm_ldg_s, "ssC*" , "" ) |
887 | BUILTIN(__nvvm_ldg_i, "iiC*" , "" ) |
888 | BUILTIN(__nvvm_ldg_l, "LiLiC*" , "" ) |
889 | BUILTIN(__nvvm_ldg_ll, "LLiLLiC*" , "" ) |
890 | |
891 | BUILTIN(__nvvm_ldg_uc, "UcUcC*" , "" ) |
892 | BUILTIN(__nvvm_ldg_us, "UsUsC*" , "" ) |
893 | BUILTIN(__nvvm_ldg_ui, "UiUiC*" , "" ) |
894 | BUILTIN(__nvvm_ldg_ul, "ULiULiC*" , "" ) |
895 | BUILTIN(__nvvm_ldg_ull, "ULLiULLiC*" , "" ) |
896 | |
897 | BUILTIN(__nvvm_ldg_h, "hhC*" , "" ) |
898 | BUILTIN(__nvvm_ldg_f, "ffC*" , "" ) |
899 | BUILTIN(__nvvm_ldg_d, "ddC*" , "" ) |
900 | |
901 | BUILTIN(__nvvm_ldg_c2, "E2cE2cC*" , "" ) |
902 | BUILTIN(__nvvm_ldg_sc2, "E2ScE2ScC*" , "" ) |
903 | BUILTIN(__nvvm_ldg_c4, "E4cE4cC*" , "" ) |
904 | BUILTIN(__nvvm_ldg_sc4, "E4ScE4ScC*" , "" ) |
905 | BUILTIN(__nvvm_ldg_s2, "E2sE2sC*" , "" ) |
906 | BUILTIN(__nvvm_ldg_s4, "E4sE4sC*" , "" ) |
907 | BUILTIN(__nvvm_ldg_i2, "E2iE2iC*" , "" ) |
908 | BUILTIN(__nvvm_ldg_i4, "E4iE4iC*" , "" ) |
909 | BUILTIN(__nvvm_ldg_l2, "E2LiE2LiC*" , "" ) |
910 | BUILTIN(__nvvm_ldg_ll2, "E2LLiE2LLiC*" , "" ) |
911 | |
912 | BUILTIN(__nvvm_ldg_uc2, "E2UcE2UcC*" , "" ) |
913 | BUILTIN(__nvvm_ldg_uc4, "E4UcE4UcC*" , "" ) |
914 | BUILTIN(__nvvm_ldg_us2, "E2UsE2UsC*" , "" ) |
915 | BUILTIN(__nvvm_ldg_us4, "E4UsE4UsC*" , "" ) |
916 | BUILTIN(__nvvm_ldg_ui2, "E2UiE2UiC*" , "" ) |
917 | BUILTIN(__nvvm_ldg_ui4, "E4UiE4UiC*" , "" ) |
918 | BUILTIN(__nvvm_ldg_ul2, "E2ULiE2ULiC*" , "" ) |
919 | BUILTIN(__nvvm_ldg_ull2, "E2ULLiE2ULLiC*" , "" ) |
920 | |
921 | BUILTIN(__nvvm_ldg_h2, "E2hE2hC*" , "" ) |
922 | BUILTIN(__nvvm_ldg_f2, "E2fE2fC*" , "" ) |
923 | BUILTIN(__nvvm_ldg_f4, "E4fE4fC*" , "" ) |
924 | BUILTIN(__nvvm_ldg_d2, "E2dE2dC*" , "" ) |
925 | |
926 | // Address space predicates. |
927 | BUILTIN(__nvvm_isspacep_const, "bvC*" , "nc" ) |
928 | BUILTIN(__nvvm_isspacep_global, "bvC*" , "nc" ) |
929 | BUILTIN(__nvvm_isspacep_local, "bvC*" , "nc" ) |
930 | BUILTIN(__nvvm_isspacep_shared, "bvC*" , "nc" ) |
931 | TARGET_BUILTIN(__nvvm_isspacep_shared_cluster,"bvC*" , "nc" , AND(SM_90,PTX78)) |
932 | |
933 | // Builtins to support WMMA instructions on sm_70 |
934 | TARGET_BUILTIN(__hmma_m16n16k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_70,PTX60)) |
935 | TARGET_BUILTIN(__hmma_m16n16k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_70,PTX60)) |
936 | TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX60)) |
937 | TARGET_BUILTIN(__hmma_m16n16k16_ld_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX60)) |
938 | TARGET_BUILTIN(__hmma_m16n16k16_st_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX60)) |
939 | TARGET_BUILTIN(__hmma_m16n16k16_st_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX60)) |
940 | |
941 | TARGET_BUILTIN(__hmma_m32n8k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
942 | TARGET_BUILTIN(__hmma_m32n8k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
943 | TARGET_BUILTIN(__hmma_m32n8k16_ld_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
944 | TARGET_BUILTIN(__hmma_m32n8k16_ld_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX61)) |
945 | TARGET_BUILTIN(__hmma_m32n8k16_st_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
946 | TARGET_BUILTIN(__hmma_m32n8k16_st_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX61)) |
947 | |
948 | TARGET_BUILTIN(__hmma_m8n32k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
949 | TARGET_BUILTIN(__hmma_m8n32k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
950 | TARGET_BUILTIN(__hmma_m8n32k16_ld_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
951 | TARGET_BUILTIN(__hmma_m8n32k16_ld_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX61)) |
952 | TARGET_BUILTIN(__hmma_m8n32k16_st_c_f16, "vi*iC*UiIi" , "" , AND(SM_70,PTX61)) |
953 | TARGET_BUILTIN(__hmma_m8n32k16_st_c_f32, "vf*fC*UiIi" , "" , AND(SM_70,PTX61)) |
954 | |
955 | TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f16, "vi*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX60)) |
956 | TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f16, "vf*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX60)) |
957 | TARGET_BUILTIN(__hmma_m16n16k16_mma_f32f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX60)) |
958 | TARGET_BUILTIN(__hmma_m16n16k16_mma_f16f32, "vi*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX60)) |
959 | |
960 | TARGET_BUILTIN(__hmma_m32n8k16_mma_f16f16, "vi*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX61)) |
961 | TARGET_BUILTIN(__hmma_m32n8k16_mma_f32f16, "vf*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX61)) |
962 | TARGET_BUILTIN(__hmma_m32n8k16_mma_f32f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX61)) |
963 | TARGET_BUILTIN(__hmma_m32n8k16_mma_f16f32, "vi*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX61)) |
964 | |
965 | TARGET_BUILTIN(__hmma_m8n32k16_mma_f16f16, "vi*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX61)) |
966 | TARGET_BUILTIN(__hmma_m8n32k16_mma_f32f16, "vf*iC*iC*iC*IiIi" , "" , AND(SM_70,PTX61)) |
967 | TARGET_BUILTIN(__hmma_m8n32k16_mma_f32f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX61)) |
968 | TARGET_BUILTIN(__hmma_m8n32k16_mma_f16f32, "vi*iC*iC*fC*IiIi" , "" , AND(SM_70,PTX61)) |
969 | |
970 | // Builtins to support integer and sub-integer WMMA instructions on sm_72/sm_75 |
971 | TARGET_BUILTIN(__bmma_m8n8k128_ld_a_b1, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
972 | TARGET_BUILTIN(__bmma_m8n8k128_ld_b_b1, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
973 | TARGET_BUILTIN(__bmma_m8n8k128_ld_c, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
974 | TARGET_BUILTIN(__bmma_m8n8k128_mma_and_popc_b1, "vi*iC*iC*iC*Ii" , "" , AND(SM_80,PTX71)) |
975 | TARGET_BUILTIN(__bmma_m8n8k128_mma_xor_popc_b1, "vi*iC*iC*iC*Ii" , "" , AND(SM_75,PTX63)) |
976 | TARGET_BUILTIN(__bmma_m8n8k128_st_c_i32, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
977 | TARGET_BUILTIN(__imma_m16n16k16_ld_a_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
978 | TARGET_BUILTIN(__imma_m16n16k16_ld_a_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
979 | TARGET_BUILTIN(__imma_m16n16k16_ld_b_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
980 | TARGET_BUILTIN(__imma_m16n16k16_ld_b_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
981 | TARGET_BUILTIN(__imma_m16n16k16_ld_c, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
982 | TARGET_BUILTIN(__imma_m16n16k16_mma_s8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
983 | TARGET_BUILTIN(__imma_m16n16k16_mma_u8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
984 | TARGET_BUILTIN(__imma_m16n16k16_st_c_i32, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
985 | TARGET_BUILTIN(__imma_m32n8k16_ld_a_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
986 | TARGET_BUILTIN(__imma_m32n8k16_ld_a_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
987 | TARGET_BUILTIN(__imma_m32n8k16_ld_b_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
988 | TARGET_BUILTIN(__imma_m32n8k16_ld_b_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
989 | TARGET_BUILTIN(__imma_m32n8k16_ld_c, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
990 | TARGET_BUILTIN(__imma_m32n8k16_mma_s8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
991 | TARGET_BUILTIN(__imma_m32n8k16_mma_u8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
992 | TARGET_BUILTIN(__imma_m32n8k16_st_c_i32, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
993 | TARGET_BUILTIN(__imma_m8n32k16_ld_a_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
994 | TARGET_BUILTIN(__imma_m8n32k16_ld_a_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
995 | TARGET_BUILTIN(__imma_m8n32k16_ld_b_s8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
996 | TARGET_BUILTIN(__imma_m8n32k16_ld_b_u8, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
997 | TARGET_BUILTIN(__imma_m8n32k16_ld_c, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
998 | TARGET_BUILTIN(__imma_m8n32k16_mma_s8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
999 | TARGET_BUILTIN(__imma_m8n32k16_mma_u8, "vi*iC*iC*iC*IiIi" , "" , AND(SM_72,PTX63)) |
1000 | TARGET_BUILTIN(__imma_m8n32k16_st_c_i32, "vi*iC*UiIi" , "" , AND(SM_72,PTX63)) |
1001 | TARGET_BUILTIN(__imma_m8n8k32_ld_a_s4, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1002 | TARGET_BUILTIN(__imma_m8n8k32_ld_a_u4, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1003 | TARGET_BUILTIN(__imma_m8n8k32_ld_b_s4, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1004 | TARGET_BUILTIN(__imma_m8n8k32_ld_b_u4, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1005 | TARGET_BUILTIN(__imma_m8n8k32_ld_c, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1006 | TARGET_BUILTIN(__imma_m8n8k32_mma_s4, "vi*iC*iC*iC*IiIi" , "" , AND(SM_75,PTX63)) |
1007 | TARGET_BUILTIN(__imma_m8n8k32_mma_u4, "vi*iC*iC*iC*IiIi" , "" , AND(SM_75,PTX63)) |
1008 | TARGET_BUILTIN(__imma_m8n8k32_st_c_i32, "vi*iC*UiIi" , "" , AND(SM_75,PTX63)) |
1009 | |
1010 | // Builtins to support double and alternate float WMMA instructions on sm_80 |
1011 | TARGET_BUILTIN(__dmma_m8n8k4_ld_a, "vd*dC*UiIi" , "" , AND(SM_80,PTX70)) |
1012 | TARGET_BUILTIN(__dmma_m8n8k4_ld_b, "vd*dC*UiIi" , "" , AND(SM_80,PTX70)) |
1013 | TARGET_BUILTIN(__dmma_m8n8k4_ld_c, "vd*dC*UiIi" , "" , AND(SM_80,PTX70)) |
1014 | TARGET_BUILTIN(__dmma_m8n8k4_st_c_f64, "vd*dC*UiIi" , "" , AND(SM_80,PTX70)) |
1015 | TARGET_BUILTIN(__dmma_m8n8k4_mma_f64, "vd*dC*dC*dC*IiIi" , "" , AND(SM_80,PTX70)) |
1016 | |
1017 | TARGET_BUILTIN(__mma_bf16_m16n16k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1018 | TARGET_BUILTIN(__mma_bf16_m16n16k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1019 | TARGET_BUILTIN(__mma_bf16_m16n16k16_mma_f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_80,PTX70)) |
1020 | TARGET_BUILTIN(__mma_bf16_m8n32k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1021 | TARGET_BUILTIN(__mma_bf16_m8n32k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1022 | TARGET_BUILTIN(__mma_bf16_m8n32k16_mma_f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_80,PTX70)) |
1023 | TARGET_BUILTIN(__mma_bf16_m32n8k16_ld_a, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1024 | TARGET_BUILTIN(__mma_bf16_m32n8k16_ld_b, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1025 | TARGET_BUILTIN(__mma_bf16_m32n8k16_mma_f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_80,PTX70)) |
1026 | |
1027 | TARGET_BUILTIN(__mma_tf32_m16n16k8_ld_a, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1028 | TARGET_BUILTIN(__mma_tf32_m16n16k8_ld_b, "vi*iC*UiIi" , "" , AND(SM_80,PTX70)) |
1029 | TARGET_BUILTIN(__mma_tf32_m16n16k8_ld_c, "vf*fC*UiIi" , "" , AND(SM_80,PTX70)) |
1030 | TARGET_BUILTIN(__mma_m16n16k8_st_c_f32, "vf*fC*UiIi" , "" , AND(SM_80,PTX70)) |
1031 | TARGET_BUILTIN(__mma_tf32_m16n16k8_mma_f32, "vf*iC*iC*fC*IiIi" , "" , AND(SM_80,PTX70)) |
1032 | |
1033 | // Async Copy |
1034 | TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive, "vWi*" , "" , AND(SM_80,PTX70)) |
1035 | TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_shared, "vWi*3" , "" , AND(SM_80,PTX70)) |
1036 | TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc, "vWi*" , "" , AND(SM_80,PTX70)) |
1037 | TARGET_BUILTIN(__nvvm_cp_async_mbarrier_arrive_noinc_shared, "vWi*3" , "" , AND(SM_80,PTX70)) |
1038 | |
1039 | TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_4, "vv*3vC*1." , "" , AND(SM_80,PTX70)) |
1040 | TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_8, "vv*3vC*1." , "" , AND(SM_80,PTX70)) |
1041 | TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1." , "" , AND(SM_80,PTX70)) |
1042 | TARGET_BUILTIN(__nvvm_cp_async_cg_shared_global_16, "vv*3vC*1." , "" , AND(SM_80,PTX70)) |
1043 | |
1044 | TARGET_BUILTIN(__nvvm_cp_async_commit_group, "v" , "" , AND(SM_80,PTX70)) |
1045 | TARGET_BUILTIN(__nvvm_cp_async_wait_group, "vIi" , "" , AND(SM_80,PTX70)) |
1046 | TARGET_BUILTIN(__nvvm_cp_async_wait_all, "v" , "" , AND(SM_80,PTX70)) |
1047 | |
1048 | |
1049 | // bf16, bf16x2 abs, neg |
1050 | TARGET_BUILTIN(__nvvm_abs_bf16, "yy" , "" , AND(SM_80,PTX70)) |
1051 | TARGET_BUILTIN(__nvvm_abs_bf16x2, "V2yV2y" , "" , AND(SM_80,PTX70)) |
1052 | TARGET_BUILTIN(__nvvm_neg_bf16, "yy" , "" , AND(SM_80,PTX70)) |
1053 | TARGET_BUILTIN(__nvvm_neg_bf16x2, "V2yV2y" , "" , AND(SM_80,PTX70)) |
1054 | |
1055 | TARGET_BUILTIN(__nvvm_mapa, "v*v*i" , "" , AND(SM_90, PTX78)) |
1056 | TARGET_BUILTIN(__nvvm_mapa_shared_cluster, "v*3v*3i" , "" , AND(SM_90, PTX78)) |
1057 | TARGET_BUILTIN(__nvvm_getctarank, "iv*" , "" , AND(SM_90, PTX78)) |
1058 | TARGET_BUILTIN(__nvvm_getctarank_shared_cluster, "iv*3" , "" , AND(SM_90,PTX78)) |
1059 | |
1060 | #undef BUILTIN |
1061 | #undef TARGET_BUILTIN |
1062 | #pragma pop_macro("AND") |
1063 | #pragma pop_macro("SM_53") |
1064 | #pragma pop_macro("SM_60") |
1065 | #pragma pop_macro("SM_70") |
1066 | #pragma pop_macro("SM_72") |
1067 | #pragma pop_macro("SM_75") |
1068 | #pragma pop_macro("SM_80") |
1069 | #pragma pop_macro("SM_86") |
1070 | #pragma pop_macro("SM_87") |
1071 | #pragma pop_macro("SM_89") |
1072 | #pragma pop_macro("SM_90") |
1073 | #pragma pop_macro("SM_90a") |
1074 | #pragma pop_macro("PTX42") |
1075 | #pragma pop_macro("PTX60") |
1076 | #pragma pop_macro("PTX61") |
1077 | #pragma pop_macro("PTX62") |
1078 | #pragma pop_macro("PTX63") |
1079 | #pragma pop_macro("PTX64") |
1080 | #pragma pop_macro("PTX65") |
1081 | #pragma pop_macro("PTX70") |
1082 | #pragma pop_macro("PTX71") |
1083 | #pragma pop_macro("PTX72") |
1084 | #pragma pop_macro("PTX73") |
1085 | #pragma pop_macro("PTX74") |
1086 | #pragma pop_macro("PTX75") |
1087 | #pragma pop_macro("PTX76") |
1088 | #pragma pop_macro("PTX77") |
1089 | #pragma pop_macro("PTX78") |
1090 | #pragma pop_macro("PTX80") |
1091 | #pragma pop_macro("PTX81") |
1092 | #pragma pop_macro("PTX82") |
1093 | #pragma pop_macro("PTX83") |
1094 | |