1 | //==- BuiltinsAMDGPU.def - AMDGPU 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 AMDGPU-specific builtin function database. Users of |
10 | // this file must define the BUILTIN macro to make use of this information. |
11 | // |
12 | // Note: (unsigned) long int type should be avoided in builtin definitions |
13 | // since it has different size on Linux (64 bit) and Windows (32 bit). |
14 | // (unsigned) long long int type should also be avoided, which is 64 bit for |
15 | // C/C++/HIP but is 128 bit for OpenCL. Use `W` as width modifier in builtin |
16 | // definitions since it is fixed for 64 bit. |
17 | //===----------------------------------------------------------------------===// |
18 | |
19 | // The format of this database matches clang/Basic/Builtins.def. |
20 | |
21 | #if defined(BUILTIN) && !defined(TARGET_BUILTIN) |
22 | # define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS) |
23 | #endif |
24 | //===----------------------------------------------------------------------===// |
25 | // SI+ only builtins. |
26 | //===----------------------------------------------------------------------===// |
27 | |
28 | BUILTIN(__builtin_amdgcn_dispatch_ptr, "v*4" , "nc" ) |
29 | BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "v*4" , "nc" ) |
30 | BUILTIN(__builtin_amdgcn_implicitarg_ptr, "v*4" , "nc" ) |
31 | BUILTIN(__builtin_amdgcn_queue_ptr, "v*4" , "nc" ) |
32 | |
33 | BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui" , "nc" ) |
34 | BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui" , "nc" ) |
35 | BUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui" , "nc" ) |
36 | |
37 | BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui" , "nc" ) |
38 | BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui" , "nc" ) |
39 | BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui" , "nc" ) |
40 | |
41 | BUILTIN(__builtin_amdgcn_workgroup_size_x, "Us" , "nc" ) |
42 | BUILTIN(__builtin_amdgcn_workgroup_size_y, "Us" , "nc" ) |
43 | BUILTIN(__builtin_amdgcn_workgroup_size_z, "Us" , "nc" ) |
44 | |
45 | BUILTIN(__builtin_amdgcn_grid_size_x, "Ui" , "nc" ) |
46 | BUILTIN(__builtin_amdgcn_grid_size_y, "Ui" , "nc" ) |
47 | BUILTIN(__builtin_amdgcn_grid_size_z, "Ui" , "nc" ) |
48 | |
49 | BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi" , "nc" ) |
50 | BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi" , "nc" ) |
51 | |
52 | TARGET_BUILTIN(__builtin_amdgcn_s_memtime, "WUi" , "n" , "s-memtime-inst" ) |
53 | |
54 | //===----------------------------------------------------------------------===// |
55 | // Instruction builtins. |
56 | //===----------------------------------------------------------------------===// |
57 | BUILTIN(__builtin_amdgcn_s_getreg, "UiIi" , "n" ) |
58 | BUILTIN(__builtin_amdgcn_s_setreg, "vIiUi" , "n" ) |
59 | BUILTIN(__builtin_amdgcn_s_getpc, "WUi" , "n" ) |
60 | BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi" , "n" ) |
61 | BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi" , "n" ) |
62 | BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi" , "n" ) |
63 | BUILTIN(__builtin_amdgcn_s_barrier, "v" , "n" ) |
64 | BUILTIN(__builtin_amdgcn_s_ttracedata, "vi" , "n" ) |
65 | BUILTIN(__builtin_amdgcn_wave_barrier, "v" , "n" ) |
66 | BUILTIN(__builtin_amdgcn_sched_barrier, "vIi" , "n" ) |
67 | BUILTIN(__builtin_amdgcn_sched_group_barrier, "vIiIiIi" , "n" ) |
68 | BUILTIN(__builtin_amdgcn_iglp_opt, "vIi" , "n" ) |
69 | BUILTIN(__builtin_amdgcn_s_dcache_inv, "v" , "n" ) |
70 | BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v" , "n" ) |
71 | BUILTIN(__builtin_amdgcn_fence, "vUicC*." , "n" ) |
72 | BUILTIN(__builtin_amdgcn_groupstaticsize, "Ui" , "n" ) |
73 | BUILTIN(__builtin_amdgcn_wavefrontsize, "Ui" , "nc" ) |
74 | |
75 | BUILTIN(__builtin_amdgcn_atomic_inc32, "UZiUZiD*UZiUicC*" , "n" ) |
76 | BUILTIN(__builtin_amdgcn_atomic_inc64, "UWiUWiD*UWiUicC*" , "n" ) |
77 | |
78 | BUILTIN(__builtin_amdgcn_atomic_dec32, "UZiUZiD*UZiUicC*" , "n" ) |
79 | BUILTIN(__builtin_amdgcn_atomic_dec64, "UWiUWiD*UWiUicC*" , "n" ) |
80 | |
81 | // FIXME: Need to disallow constant address space. |
82 | BUILTIN(__builtin_amdgcn_div_scale, "dddbb*" , "n" ) |
83 | BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*" , "n" ) |
84 | BUILTIN(__builtin_amdgcn_div_fmas, "ddddb" , "nc" ) |
85 | BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb" , "nc" ) |
86 | BUILTIN(__builtin_amdgcn_div_fixup, "dddd" , "nc" ) |
87 | BUILTIN(__builtin_amdgcn_div_fixupf, "ffff" , "nc" ) |
88 | BUILTIN(__builtin_amdgcn_trig_preop, "ddi" , "nc" ) |
89 | BUILTIN(__builtin_amdgcn_trig_preopf, "ffi" , "nc" ) |
90 | BUILTIN(__builtin_amdgcn_rcp, "dd" , "nc" ) |
91 | BUILTIN(__builtin_amdgcn_rcpf, "ff" , "nc" ) |
92 | BUILTIN(__builtin_amdgcn_sqrt, "dd" , "nc" ) |
93 | BUILTIN(__builtin_amdgcn_sqrtf, "ff" , "nc" ) |
94 | BUILTIN(__builtin_amdgcn_rsq, "dd" , "nc" ) |
95 | BUILTIN(__builtin_amdgcn_rsqf, "ff" , "nc" ) |
96 | BUILTIN(__builtin_amdgcn_rsq_clamp, "dd" , "nc" ) |
97 | BUILTIN(__builtin_amdgcn_rsq_clampf, "ff" , "nc" ) |
98 | BUILTIN(__builtin_amdgcn_sinf, "ff" , "nc" ) |
99 | BUILTIN(__builtin_amdgcn_cosf, "ff" , "nc" ) |
100 | BUILTIN(__builtin_amdgcn_logf, "ff" , "nc" ) |
101 | BUILTIN(__builtin_amdgcn_exp2f, "ff" , "nc" ) |
102 | BUILTIN(__builtin_amdgcn_log_clampf, "ff" , "nc" ) |
103 | BUILTIN(__builtin_amdgcn_ldexp, "ddi" , "nc" ) |
104 | BUILTIN(__builtin_amdgcn_ldexpf, "ffi" , "nc" ) |
105 | BUILTIN(__builtin_amdgcn_frexp_mant, "dd" , "nc" ) |
106 | BUILTIN(__builtin_amdgcn_frexp_mantf, "ff" , "nc" ) |
107 | BUILTIN(__builtin_amdgcn_frexp_exp, "id" , "nc" ) |
108 | BUILTIN(__builtin_amdgcn_frexp_expf, "if" , "nc" ) |
109 | BUILTIN(__builtin_amdgcn_fract, "dd" , "nc" ) |
110 | BUILTIN(__builtin_amdgcn_fractf, "ff" , "nc" ) |
111 | BUILTIN(__builtin_amdgcn_lerp, "UiUiUiUi" , "nc" ) |
112 | BUILTIN(__builtin_amdgcn_class, "bdi" , "nc" ) |
113 | BUILTIN(__builtin_amdgcn_classf, "bfi" , "nc" ) |
114 | BUILTIN(__builtin_amdgcn_cubeid, "ffff" , "nc" ) |
115 | BUILTIN(__builtin_amdgcn_cubesc, "ffff" , "nc" ) |
116 | BUILTIN(__builtin_amdgcn_cubetc, "ffff" , "nc" ) |
117 | BUILTIN(__builtin_amdgcn_cubema, "ffff" , "nc" ) |
118 | BUILTIN(__builtin_amdgcn_s_sleep, "vIi" , "n" ) |
119 | BUILTIN(__builtin_amdgcn_s_incperflevel, "vIi" , "n" ) |
120 | BUILTIN(__builtin_amdgcn_s_decperflevel, "vIi" , "n" ) |
121 | BUILTIN(__builtin_amdgcn_s_setprio, "vIs" , "n" ) |
122 | BUILTIN(__builtin_amdgcn_ds_swizzle, "iiIi" , "nc" ) |
123 | BUILTIN(__builtin_amdgcn_ds_permute, "iii" , "nc" ) |
124 | BUILTIN(__builtin_amdgcn_ds_bpermute, "iii" , "nc" ) |
125 | BUILTIN(__builtin_amdgcn_readfirstlane, "ii" , "nc" ) |
126 | BUILTIN(__builtin_amdgcn_readlane, "iii" , "nc" ) |
127 | BUILTIN(__builtin_amdgcn_fmed3f, "ffff" , "nc" ) |
128 | BUILTIN(__builtin_amdgcn_ds_faddf, "ff*3fIiIiIb" , "n" ) |
129 | BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb" , "n" ) |
130 | BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb" , "n" ) |
131 | BUILTIN(__builtin_amdgcn_ds_append, "ii*3" , "n" ) |
132 | BUILTIN(__builtin_amdgcn_ds_consume, "ii*3" , "n" ) |
133 | BUILTIN(__builtin_amdgcn_alignbit, "UiUiUiUi" , "nc" ) |
134 | BUILTIN(__builtin_amdgcn_alignbyte, "UiUiUiUi" , "nc" ) |
135 | BUILTIN(__builtin_amdgcn_ubfe, "UiUiUiUi" , "nc" ) |
136 | BUILTIN(__builtin_amdgcn_sbfe, "UiUiUiUi" , "nc" ) |
137 | BUILTIN(__builtin_amdgcn_cvt_pkrtz, "E2hff" , "nc" ) |
138 | BUILTIN(__builtin_amdgcn_cvt_pknorm_i16, "E2sff" , "nc" ) |
139 | BUILTIN(__builtin_amdgcn_cvt_pknorm_u16, "E2Usff" , "nc" ) |
140 | BUILTIN(__builtin_amdgcn_cvt_pk_i16, "E2sii" , "nc" ) |
141 | BUILTIN(__builtin_amdgcn_cvt_pk_u16, "E2UsUiUi" , "nc" ) |
142 | BUILTIN(__builtin_amdgcn_cvt_pk_u8_f32, "UifUiUi" , "nc" ) |
143 | BUILTIN(__builtin_amdgcn_cvt_off_f32_i4, "fi" , "nc" ) |
144 | BUILTIN(__builtin_amdgcn_sad_u8, "UiUiUiUi" , "nc" ) |
145 | BUILTIN(__builtin_amdgcn_msad_u8, "UiUiUiUi" , "nc" ) |
146 | BUILTIN(__builtin_amdgcn_sad_hi_u8, "UiUiUiUi" , "nc" ) |
147 | BUILTIN(__builtin_amdgcn_sad_u16, "UiUiUiUi" , "nc" ) |
148 | BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi" , "nc" ) |
149 | BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi" , "nc" ) |
150 | BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui" , "nc" ) |
151 | |
152 | BUILTIN(__builtin_amdgcn_make_buffer_rsrc, "Qbv*sii" , "nc" ) |
153 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b8, "vUcQbiiIi" , "n" ) |
154 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b16, "vUsQbiiIi" , "n" ) |
155 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b32, "vUiQbiiIi" , "n" ) |
156 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b64, "vV2UiQbiiIi" , "n" ) |
157 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b96, "vV3UiQbiiIi" , "n" ) |
158 | BUILTIN(__builtin_amdgcn_raw_buffer_store_b128, "vV4UiQbiiIi" , "n" ) |
159 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b8, "UcQbiiIi" , "n" ) |
160 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b16, "UsQbiiIi" , "n" ) |
161 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b32, "UiQbiiIi" , "n" ) |
162 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b64, "V2UiQbiiIi" , "n" ) |
163 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b96, "V3UiQbiiIi" , "n" ) |
164 | BUILTIN(__builtin_amdgcn_raw_buffer_load_b128, "V4UiQbiiIi" , "n" ) |
165 | |
166 | TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi" , "t" , "vmem-to-lds-load-insts" ) |
167 | |
168 | //===----------------------------------------------------------------------===// |
169 | // Ballot builtins. |
170 | //===----------------------------------------------------------------------===// |
171 | |
172 | TARGET_BUILTIN(__builtin_amdgcn_ballot_w32, "ZUib" , "nc" , "wavefrontsize32" ) |
173 | BUILTIN(__builtin_amdgcn_ballot_w64, "WUib" , "nc" ) |
174 | |
175 | // Deprecated intrinsics in favor of __builtin_amdgn_ballot_{w32|w64} |
176 | BUILTIN(__builtin_amdgcn_uicmp, "WUiUiUiIi" , "nc" ) |
177 | BUILTIN(__builtin_amdgcn_uicmpl, "WUiWUiWUiIi" , "nc" ) |
178 | BUILTIN(__builtin_amdgcn_sicmp, "WUiiiIi" , "nc" ) |
179 | BUILTIN(__builtin_amdgcn_sicmpl, "WUiWiWiIi" , "nc" ) |
180 | BUILTIN(__builtin_amdgcn_fcmp, "WUiddIi" , "nc" ) |
181 | BUILTIN(__builtin_amdgcn_fcmpf, "WUiffIi" , "nc" ) |
182 | |
183 | //===----------------------------------------------------------------------===// |
184 | // Flat addressing builtins. |
185 | //===----------------------------------------------------------------------===// |
186 | BUILTIN(__builtin_amdgcn_is_shared, "bvC*0" , "nc" ) |
187 | BUILTIN(__builtin_amdgcn_is_private, "bvC*0" , "nc" ) |
188 | |
189 | //===----------------------------------------------------------------------===// |
190 | // GWS builtins. |
191 | //===----------------------------------------------------------------------===// |
192 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi" , "n" , "gws" ) |
193 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi" , "n" , "gws" ) |
194 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi" , "n" , "gws" ) |
195 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi" , "n" , "gws" ) |
196 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi" , "n" , "gws" ) |
197 | |
198 | //===----------------------------------------------------------------------===// |
199 | // CI+ only builtins. |
200 | //===----------------------------------------------------------------------===// |
201 | TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v" , "n" , "ci-insts" ) |
202 | TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v" , "n" , "ci-insts" ) |
203 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi" , "n" , "ci-insts" ) |
204 | |
205 | //===----------------------------------------------------------------------===// |
206 | // Interpolation builtins. |
207 | //===----------------------------------------------------------------------===// |
208 | BUILTIN(__builtin_amdgcn_interp_p1_f16, "ffUiUibUi" , "nc" ) |
209 | BUILTIN(__builtin_amdgcn_interp_p2_f16, "hffUiUibUi" , "nc" ) |
210 | BUILTIN(__builtin_amdgcn_interp_p1, "ffUiUiUi" , "nc" ) |
211 | BUILTIN(__builtin_amdgcn_interp_p2, "fffUiUiUi" , "nc" ) |
212 | BUILTIN(__builtin_amdgcn_interp_mov, "fUiUiUiUi" , "nc" ) |
213 | |
214 | //===----------------------------------------------------------------------===// |
215 | // VI+ only builtins. |
216 | //===----------------------------------------------------------------------===// |
217 | |
218 | TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh" , "nc" , "16-bit-insts" ) |
219 | TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh" , "nc" , "16-bit-insts" ) |
220 | TARGET_BUILTIN(__builtin_amdgcn_sqrth, "hh" , "nc" , "16-bit-insts" ) |
221 | TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh" , "nc" , "16-bit-insts" ) |
222 | TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh" , "nc" , "16-bit-insts" ) |
223 | TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh" , "nc" , "16-bit-insts" ) |
224 | TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi" , "nc" , "16-bit-insts" ) |
225 | TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh" , "nc" , "16-bit-insts" ) |
226 | TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh" , "nc" , "16-bit-insts" ) |
227 | TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh" , "nc" , "16-bit-insts" ) |
228 | TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi" , "nc" , "16-bit-insts" ) |
229 | TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "WUi" , "n" , "s-memrealtime" ) |
230 | TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb" , "nct" , "dpp" ) |
231 | TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb" , "nct" , "dpp" ) |
232 | TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v" , "n" , "gfx8-insts" ) |
233 | TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi" , "nc" , "gfx8-insts" ) |
234 | |
235 | //===----------------------------------------------------------------------===// |
236 | // GFX9+ only builtins. |
237 | //===----------------------------------------------------------------------===// |
238 | |
239 | TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh" , "nc" , "gfx9-insts" ) |
240 | |
241 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
242 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f" , "t" , "atomic-fadd-rtn-insts" ) |
243 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h" , "t" , "atomic-buffer-global-pk-add-f16-insts" ) |
244 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
245 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
246 | |
247 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
248 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmin_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
249 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
250 | |
251 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d" , "t" , "gfx90a-insts" ) |
252 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f" , "t" , "gfx8-insts" ) |
253 | |
254 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f" , "t" , "gfx940-insts" ) |
255 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h" , "t" , "atomic-flat-pk-add-16-insts" ) |
256 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s" , "t" , "atomic-flat-pk-add-16-insts" ) |
257 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s" , "t" , "atomic-global-pk-add-bf16-inst" ) |
258 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s" , "t" , "atomic-ds-pk-add-16-insts" ) |
259 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h" , "t" , "atomic-ds-pk-add-16-insts" ) |
260 | TARGET_BUILTIN(__builtin_amdgcn_load_to_lds, "vv*v*3IUiIiIUi" , "" , "vmem-to-lds-load-insts" ) |
261 | TARGET_BUILTIN(__builtin_amdgcn_global_load_lds, "vv*1v*3IUiIiIUi" , "t" , "vmem-to-lds-load-insts" ) |
262 | |
263 | //===----------------------------------------------------------------------===// |
264 | // Deep learning builtins. |
265 | //===----------------------------------------------------------------------===// |
266 | |
267 | TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb" , "nc" , "dot10-insts" ) |
268 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_f16_f16, "hV2hV2hh" , "nc" , "dot9-insts" ) |
269 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_bf16_bf16, "sV2sV2ss" , "nc" , "dot9-insts" ) |
270 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_f32_bf16, "fV2sV2sfIb" , "nc" , "dot12-insts" ) |
271 | TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb" , "nc" , "dot2-insts" ) |
272 | TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb" , "nc" , "dot2-insts" ) |
273 | TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb" , "nc" , "dot1-insts" ) |
274 | TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb" , "nc" , "dot7-insts" ) |
275 | TARGET_BUILTIN(__builtin_amdgcn_sudot4, "iIbiIbiiIb" , "nc" , "dot8-insts" ) |
276 | TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb" , "nc" , "dot1-insts" ) |
277 | TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb" , "nc" , "dot7-insts" ) |
278 | TARGET_BUILTIN(__builtin_amdgcn_sudot8, "iIbiIbiiIb" , "nc" , "dot8-insts" ) |
279 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_bf8, "fUiUif" , "nc" , "dot11-insts" ) |
280 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_fp8, "fUiUif" , "nc" , "dot11-insts" ) |
281 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_fp8, "fUiUif" , "nc" , "dot11-insts" ) |
282 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif" , "nc" , "dot11-insts" ) |
283 | TARGET_BUILTIN(__builtin_amdgcn_fdot2c_f32_bf16, "fV2yV2yfIb" , "nc" , "dot13-insts" ) |
284 | |
285 | //===----------------------------------------------------------------------===// |
286 | // GFX10+ only builtins. |
287 | //===----------------------------------------------------------------------===// |
288 | TARGET_BUILTIN(__builtin_amdgcn_permlane16, "UiUiUiUiUiIbIb" , "nc" , "gfx10-insts" ) |
289 | TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb" , "nc" , "gfx10-insts" ) |
290 | TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi" , "nct" , "gfx10-insts" ) |
291 | TARGET_BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "vIs" , "n" , "gfx10-insts" ) |
292 | |
293 | //===----------------------------------------------------------------------===// |
294 | // Raytracing builtins. |
295 | // By default the 1st argument is i32 and the 4/5-th arguments are float4. |
296 | // Postfix l indicates the 1st argument is i64. |
297 | // Postfix h indicates the 4/5-th arguments are half4. |
298 | //===----------------------------------------------------------------------===// |
299 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray, "V4UiUifV4fV4fV4fV4Ui" , "nc" , "gfx10-insts" ) |
300 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_h, "V4UiUifV4fV4hV4hV4Ui" , "nc" , "gfx10-insts" ) |
301 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_l, "V4UiWUifV4fV4fV4fV4Ui" , "nc" , "gfx10-insts" ) |
302 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_lh, "V4UiWUifV4fV4hV4hV4Ui" , "nc" , "gfx10-insts" ) |
303 | |
304 | |
305 | //===----------------------------------------------------------------------===// |
306 | // GFX11+ only builtins. |
307 | //===----------------------------------------------------------------------===// |
308 | |
309 | // TODO: This is a no-op in wave32. Should the builtin require wavefrontsize64? |
310 | TARGET_BUILTIN(__builtin_amdgcn_permlane64, "UiUi" , "nc" , "gfx11-insts" ) |
311 | TARGET_BUILTIN(__builtin_amdgcn_s_wait_event_export_ready, "v" , "n" , "gfx11-insts" ) |
312 | |
313 | //===----------------------------------------------------------------------===// |
314 | // WMMA builtins. |
315 | // Postfix w32 indicates the builtin requires wavefront size of 32. |
316 | // Postfix w64 indicates the builtin requires wavefront size of 64. |
317 | //===----------------------------------------------------------------------===// |
318 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, "V8fV16hV16hV8f" , "nc" , "gfx11-insts,wavefrontsize32" ) |
319 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32, "V8fV16sV16sV8f" , "nc" , "gfx11-insts,wavefrontsize32" ) |
320 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32, "V16hV16hV16hV16hIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
321 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32, "V16sV16sV16sV16sIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
322 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32, "V16hV16hV16hV16hIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
323 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32, "V16sV16sV16sV16sIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
324 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32, "V8iIbV4iIbV4iV8iIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
325 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
326 | |
327 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64, "V4fV16hV16hV4f" , "nc" , "gfx11-insts,wavefrontsize64" ) |
328 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64, "V4fV16sV16sV4f" , "nc" , "gfx11-insts,wavefrontsize64" ) |
329 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64, "V8hV16hV16hV8hIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
330 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64, "V8sV16sV16sV8sIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
331 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64, "V8hV16hV16hV8hIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
332 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64, "V8sV16sV16sV8sIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
333 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64, "V4iIbV4iIbV4iV4iIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
334 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64, "V4iIbV2iIbV2iV4iIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
335 | |
336 | TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtn, "UiUIi" , "n" , "gfx11-insts" ) |
337 | TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtnl, "UWiUIi" , "n" , "gfx11-insts" ) |
338 | |
339 | TARGET_BUILTIN(__builtin_amdgcn_ds_bvh_stack_rtn, "V2UiUiUiV4UiIi" , "n" , "gfx11-insts" ) |
340 | |
341 | //===----------------------------------------------------------------------===// |
342 | // Special builtins. |
343 | //===----------------------------------------------------------------------===// |
344 | BUILTIN(__builtin_amdgcn_read_exec, "WUi" , "nc" ) |
345 | BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui" , "nc" ) |
346 | BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui" , "nc" ) |
347 | |
348 | BUILTIN(__builtin_amdgcn_endpgm, "v" , "nr" ) |
349 | |
350 | BUILTIN(__builtin_amdgcn_get_fpenv, "WUi" , "n" ) |
351 | BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi" , "n" ) |
352 | |
353 | //===----------------------------------------------------------------------===// |
354 | // R600-NI only builtins. |
355 | //===----------------------------------------------------------------------===// |
356 | |
357 | BUILTIN(__builtin_r600_implicitarg_ptr, "Uc*7" , "nc" ) |
358 | |
359 | BUILTIN(__builtin_r600_read_tgid_x, "Ui" , "nc" ) |
360 | BUILTIN(__builtin_r600_read_tgid_y, "Ui" , "nc" ) |
361 | BUILTIN(__builtin_r600_read_tgid_z, "Ui" , "nc" ) |
362 | |
363 | BUILTIN(__builtin_r600_read_tidig_x, "Ui" , "nc" ) |
364 | BUILTIN(__builtin_r600_read_tidig_y, "Ui" , "nc" ) |
365 | BUILTIN(__builtin_r600_read_tidig_z, "Ui" , "nc" ) |
366 | |
367 | BUILTIN(__builtin_r600_recipsqrt_ieee, "dd" , "nc" ) |
368 | BUILTIN(__builtin_r600_recipsqrt_ieeef, "ff" , "nc" ) |
369 | |
370 | //===----------------------------------------------------------------------===// |
371 | // MFMA builtins. |
372 | //===----------------------------------------------------------------------===// |
373 | |
374 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x1f32, "V32fffV32fIiIiIi" , "nc" , "mai-insts" ) |
375 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x1f32, "V16fffV16fIiIiIi" , "nc" , "mai-insts" ) |
376 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x1f32, "V4fffV4fIiIiIi" , "nc" , "mai-insts" ) |
377 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2f32, "V16fffV16fIiIiIi" , "nc" , "mai-insts" ) |
378 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f32, "V4fffV4fIiIiIi" , "nc" , "mai-insts" ) |
379 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4f16, "V32fV4hV4hV32fIiIiIi" , "nc" , "mai-insts" ) |
380 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f16, "V16fV4hV4hV16fIiIiIi" , "nc" , "mai-insts" ) |
381 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x4f16, "V4fV4hV4hV4fIiIiIi" , "nc" , "mai-insts" ) |
382 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x8f16, "V16fV4hV4hV16fIiIiIi" , "nc" , "mai-insts" ) |
383 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x16f16, "V4fV4hV4hV4fIiIiIi" , "nc" , "mai-insts" ) |
384 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x4i8, "V32iiiV32iIiIiIi" , "nc" , "mai-insts" ) |
385 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x4i8, "V16iiiV16iIiIiIi" , "nc" , "mai-insts" ) |
386 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_4x4x4i8, "V4iiiV4iIiIiIi" , "nc" , "mai-insts" ) |
387 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x8i8, "V16iiiV16iIiIiIi" , "nc" , "mai-insts" ) |
388 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x16i8, "V4iiiV4iIiIiIi" , "nc" , "mai-insts" ) |
389 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2bf16, "V32fV2sV2sV32fIiIiIi" , "nc" , "mai-insts" ) |
390 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x2bf16, "V16fV2sV2sV16fIiIiIi" , "nc" , "mai-insts" ) |
391 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x2bf16, "V4fV2sV2sV4fIiIiIi" , "nc" , "mai-insts" ) |
392 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4bf16, "V16fV2sV2sV16fIiIiIi" , "nc" , "mai-insts" ) |
393 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8bf16, "V4fV2sV2sV4fIiIiIi" , "nc" , "mai-insts" ) |
394 | |
395 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4bf16_1k, "V32fV4sV4sV32fIiIiIi" , "nc" , "mai-insts" ) |
396 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4bf16_1k, "V16fV4sV4sV16fIiIiIi" , "nc" , "mai-insts" ) |
397 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x4bf16_1k, "V4fV4sV4sV4fIiIiIi" , "nc" , "mai-insts" ) |
398 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x8bf16_1k, "V16fV4sV4sV16fIiIiIi" , "nc" , "mai-insts" ) |
399 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x16bf16_1k, "V4fV4sV4sV4fIiIiIi" , "nc" , "mai-insts" ) |
400 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f64_16x16x4f64, "V4dddV4dIiIiIi" , "nc" , "mai-insts" ) |
401 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f64_4x4x4f64, "ddddIiIiIi" , "nc" , "mai-insts" ) |
402 | |
403 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x32_i8, "V4iWiWiV4iIiIiIi" , "nc" , "mai-insts" ) |
404 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x16_i8, "V16iWiWiV16iIiIiIi" , "nc" , "mai-insts" ) |
405 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8_xf32, "V4fV2fV2fV4fIiIiIi" , "nc" , "mai-insts" ) |
406 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4_xf32, "V16fV2fV2fV16fIiIiIi" , "nc" , "mai-insts" ) |
407 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
408 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
409 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
410 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
411 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
412 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
413 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
414 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
415 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_f16, "V4fV4hV8hV4fiIiIi" , "nc" , "mai-insts" ) |
416 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_f16, "V16fV4hV8hV16fiIiIi" , "nc" , "mai-insts" ) |
417 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_bf16, "V4fV4sV8sV4fiIiIi" , "nc" , "mai-insts" ) |
418 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_bf16, "V16fV4sV8sV16fiIiIi" , "nc" , "mai-insts" ) |
419 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x64_i8, "V4iV2iV4iV4iiIiIi" , "nc" , "mai-insts" ) |
420 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_32x32x32_i8, "V16iV2iV4iV16iiIiIi" , "nc" , "mai-insts" ) |
421 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
422 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
423 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
424 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
425 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf8_bf8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
426 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
427 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
428 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
429 | |
430 | TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_bf8, "fiIi" , "nc" , "fp8-conversion-insts" ) |
431 | TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_fp8, "fiIi" , "nc" , "fp8-conversion-insts" ) |
432 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_bf8, "V2fiIb" , "nc" , "fp8-conversion-insts" ) |
433 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_fp8, "V2fiIb" , "nc" , "fp8-conversion-insts" ) |
434 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_bf8_f32, "iffiIb" , "nc" , "fp8-conversion-insts" ) |
435 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb" , "nc" , "fp8-conversion-insts" ) |
436 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi" , "nc" , "fp8-conversion-insts" ) |
437 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi" , "nc" , "fp8-conversion-insts" ) |
438 | |
439 | //===----------------------------------------------------------------------===// |
440 | // GFX950 only builtins. |
441 | //===----------------------------------------------------------------------===// |
442 | TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4, "V4fV8ZiV8ZiV4fIiIiIiiIii" , "nc" , "gfx950-insts" ) |
443 | TARGET_BUILTIN(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4, "V16fV8ZiV8ZiV16fIiIiIiiIii" , "nc" , "gfx950-insts" ) |
444 | |
445 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_f16, "V4fV8hV8hV4fIiIiIi" , "nc" , "gfx950-insts" ) |
446 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf16, "V4fV8yV8yV4fIiIiIi" , "nc" , "gfx950-insts" ) |
447 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_f16, "V16fV8hV8hV16fIiIiIi" , "nc" , "gfx950-insts" ) |
448 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_bf16, "V16fV8yV8yV16fIiIiIi" , "nc" , "gfx950-insts" ) |
449 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x64_i8, "V4iV4iV4iV4iIiIiIi" , "nc" , "gfx950-insts" ) |
450 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x32_i8, "V16iV4iV4iV16iIiIiIi" , "nc" , "gfx950-insts" ) |
451 | |
452 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_f16, "V4fV8hV16hV4fiIiIi" , "nc" , "gfx950-insts" ) |
453 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_f16, "V16fV8hV16hV16fiIiIi" , "nc" , "gfx950-insts" ) |
454 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf16, "V4fV8yV16yV4fiIiIi" , "nc" , "gfx950-insts" ) |
455 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf16, "V16fV8yV16yV16fiIiIi" , "nc" , "gfx950-insts" ) |
456 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x128_i8, "V4iV4iV8iV4iiIiIi" , "nc" , "gfx950-insts" ) |
457 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_32x32x64_i8, "V16iV4iV8iV16iiIiIi" , "nc" , "gfx950-insts" ) |
458 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8, "V4fV4iV8iV4fiIiIi" , "nc" , "gfx950-insts" ) |
459 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8, "V4fV4iV8iV4fiIiIi" , "nc" , "gfx950-insts" ) |
460 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8, "V4fV4iV8iV4fiIiIi" , "nc" , "gfx950-insts" ) |
461 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8, "V4fV4iV8iV4fiIiIi" , "nc" , "gfx950-insts" ) |
462 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8, "V16fV4iV8iV16fiIiIi" , "nc" , "gfx950-insts" ) |
463 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8, "V16fV4iV8iV16fiIiIi" , "nc" , "gfx950-insts" ) |
464 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8, "V16fV4iV8iV16fiIiIi" , "nc" , "gfx950-insts" ) |
465 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8, "V16fV4iV8iV16fiIiIi" , "nc" , "gfx950-insts" ) |
466 | |
467 | TARGET_BUILTIN(__builtin_amdgcn_permlane16_swap, "V2UiUiUiIbIb" , "nc" , "permlane16-swap" ) |
468 | TARGET_BUILTIN(__builtin_amdgcn_permlane32_swap, "V2UiUiUiIbIb" , "nc" , "permlane32-swap" ) |
469 | |
470 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr4_b64_v2i32, "V2iV2i*3" , "nc" , "gfx950-insts" ) |
471 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr6_b96_v3i32, "V3iV3i*3" , "nc" , "gfx950-insts" ) |
472 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr8_b64_v2i32, "V2iV2i*3" , "nc" , "gfx950-insts" ) |
473 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr16_b64_v4i16, "V4sV4s*3" , "nc" , "gfx950-insts" ) |
474 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr16_b64_v4f16, "V4hV4h*3" , "nc" , "gfx950-insts" ) |
475 | TARGET_BUILTIN(__builtin_amdgcn_ds_read_tr16_b64_v4bf16, "V4yV4y*3" , "nc" , "gfx950-insts" ) |
476 | |
477 | TARGET_BUILTIN(__builtin_amdgcn_ashr_pk_i8_i32, "UsUiUiUi" , "nc" , "ashr-pk-insts" ) |
478 | TARGET_BUILTIN(__builtin_amdgcn_ashr_pk_u8_i32, "UsUiUiUi" , "nc" , "ashr-pk-insts" ) |
479 | |
480 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32, "V6UiV16fV16ff" , "nc" , "gfx950-insts" ) |
481 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32, "V6UiV16fV16ff" , "nc" , "gfx950-insts" ) |
482 | |
483 | //===----------------------------------------------------------------------===// |
484 | // GFX12+ only builtins. |
485 | //===----------------------------------------------------------------------===// |
486 | |
487 | TARGET_BUILTIN(__builtin_amdgcn_s_sleep_var, "vUi" , "n" , "gfx12-insts" ) |
488 | TARGET_BUILTIN(__builtin_amdgcn_permlane16_var, "UiUiUiUiIbIb" , "nc" , "gfx12-insts" ) |
489 | TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb" , "nc" , "gfx12-insts" ) |
490 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi" , "n" , "gfx12-insts" ) |
491 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vv*i" , "n" , "gfx12-insts" ) |
492 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs" , "n" , "gfx12-insts" ) |
493 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi" , "n" , "gfx12-insts" ) |
494 | TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii" , "n" , "gfx12-insts" ) |
495 | TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*" , "n" , "gfx12-insts" ) |
496 | TARGET_BUILTIN(__builtin_amdgcn_s_prefetch_data, "vvC*Ui" , "nc" , "gfx12-insts" ) |
497 | TARGET_BUILTIN(__builtin_amdgcn_s_buffer_prefetch_data, "vQbIiUi" , "nc" , "gfx12-insts" ) |
498 | |
499 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b64_v2i32, "V2iV2i*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
500 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v8i16, "V8sV8s*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
501 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v8f16, "V8hV8h*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
502 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v8bf16, "V8yV8y*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
503 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b64_i32, "ii*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
504 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16, "V4sV4s*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
505 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4f16, "V4hV4h*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
506 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4bf16, "V4yV4y*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
507 | |
508 | TARGET_BUILTIN(__builtin_amdgcn_ds_bpermute_fi_b32, "iii" , "nc" , "gfx12-insts" ) |
509 | |
510 | // For the following two builtins, the second and third return values of the |
511 | // intrinsics are returned through the last two pointer-type function arguments. |
512 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh8_intersect_ray, "V10UiWUifUcV3fV3fUiV4UiV3f*V3f*" , "nc" , "gfx12-insts" ) |
513 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_dual_intersect_ray, "V10UiWUifUcV3fV3fV2UiV4UiV3f*V3f*" , "nc" , "gfx12-insts" ) |
514 | |
515 | TARGET_BUILTIN(__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn, "V2UiUiUiV4UiIi" , "n" , "gfx11-insts" ) |
516 | TARGET_BUILTIN(__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn, "V2UiUiUiV8UiIi" , "n" , "gfx12-insts" ) |
517 | |
518 | // The intrinsic returns {i64, i32}, the builtin returns <2 x i64>. |
519 | // The second return value of the intrinsic is zext'ed. |
520 | TARGET_BUILTIN(__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn, "V2WUiUiUiV8UiIi" , "n" , "gfx12-insts" ) |
521 | |
522 | //===----------------------------------------------------------------------===// |
523 | // WMMA builtins. |
524 | // Postfix w32 indicates the builtin requires wavefront size of 32. |
525 | // Postfix w64 indicates the builtin requires wavefront size of 64. |
526 | // |
527 | // Some of these are very similar to their GFX11 counterparts, but they don't |
528 | // require replication of the A,B matrices, so they use fewer vector elements. |
529 | // Therefore, we add an "_gfx12" suffix to distinguish them from the existing |
530 | // builtins. |
531 | //===----------------------------------------------------------------------===// |
532 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12, "V8fV8hV8hV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
533 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12, "V8fV8sV8sV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
534 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12, "V8hV8hV8hV8h" , "nc" , "gfx12-insts,wavefrontsize32" ) |
535 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12, "V8sV8sV8sV8s" , "nc" , "gfx12-insts,wavefrontsize32" ) |
536 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
537 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12, "V8iIbiIbiV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
538 | // These are gfx12-only, but for consistency with the other WMMA variants we're |
539 | // keeping the "_gfx12" suffix. |
540 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
541 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
542 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
543 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
544 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
545 | |
546 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12, "V4fV4hV4hV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
547 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12, "V4fV4sV4sV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
548 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12, "V4hV4hV4hV4h" , "nc" , "gfx12-insts,wavefrontsize64" ) |
549 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12, "V4sV4sV4sV4s" , "nc" , "gfx12-insts,wavefrontsize64" ) |
550 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
551 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
552 | // These are gfx12-only, but for consistency with the other WMMA variants we're |
553 | // keeping the "_gfx12" suffix. |
554 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
555 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
556 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
557 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
558 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
559 | |
560 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32, "V8fV8hV16hV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
561 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32, "V8fV8sV16sV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
562 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32, "V8hV8hV16hV8hi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
563 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32, "V8sV8sV16sV8si" , "nc" , "gfx12-insts,wavefrontsize32" ) |
564 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32, "V8iIbV2iIbV4iV8iiIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
565 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32, "V8iIbiIbV2iV8iiIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
566 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32, "V8iIbV2iIbV4iV8iiIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
567 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32, "V8fV2iV4iV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
568 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32, "V8fV2iV4iV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
569 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32, "V8fV2iV4iV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
570 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32, "V8fV2iV4iV8fi" , "nc" , "gfx12-insts,wavefrontsize32" ) |
571 | |
572 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64, "V4fV4hV8hV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
573 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64, "V4fV4sV8sV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
574 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64, "V4hV4hV8hV4hi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
575 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64, "V4sV4sV8sV4si" , "nc" , "gfx12-insts,wavefrontsize64" ) |
576 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64, "V4iIbiIbV2iV4iiIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
577 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64, "V4iIbiIbiV4iiIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
578 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64, "V4iIbiIbV2iV4iiIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
579 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64, "V4fiV2iV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
580 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64, "V4fiV2iV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
581 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
582 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fi" , "nc" , "gfx12-insts,wavefrontsize64" ) |
583 | |
584 | TARGET_BUILTIN(__builtin_amdgcn_prng_b32, "UiUi" , "nc" , "prng-inst" ) |
585 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_fp6_f16, "V6UiV32hf" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
586 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf6_f16, "V6UiV32hf" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
587 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_fp6_bf16, "V6UiV32yf" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
588 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16, "V6UiV32yf" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
589 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_f16_fp8, "V2hV2hifIiIb" , "nc" , "fp8-cvt-scale-insts" ) |
590 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_f16_bf8, "V2hV2hifIiIb" , "nc" , "bf8-cvt-scale-insts" ) |
591 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_f32_fp8, "fifIi" , "nc" , "fp8-cvt-scale-insts" ) |
592 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_f32_bf8, "fifIi" , "nc" , "bf8-cvt-scale-insts" ) |
593 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_f32, "V2sV2sfffIb" , "nc" , "fp8-cvt-scale-insts" ) |
594 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_f32, "V2sV2sfffIb" , "nc" , "bf8-cvt-scale-insts" ) |
595 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_fp8, "V2fUifIb" , "nc" , "fp8-cvt-scale-insts" ) |
596 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_bf8, "V2fUifIb" , "nc" , "bf8-cvt-scale-insts" ) |
597 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_f16, "V2sV2sV2hfIb" , "nc" , "fp8-cvt-scale-insts" ) |
598 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp8_bf16, "V2sV2sV2yfIb" , "nc" , "fp8-cvt-scale-insts" ) |
599 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_f16, "V2sV2sV2hfIb" , "nc" , "bf8-cvt-scale-insts" ) |
600 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf8_bf16, "V2sV2sV2yfIb" , "nc" , "bf8-cvt-scale-insts" ) |
601 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f32_fp4, "V2fUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
602 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp4_f32, "UiUifffIi" , "nc" , "fp4-cvt-scale-insts" ) |
603 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f16_fp4, "V2hUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
604 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp4, "V2yUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
605 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f32_fp6, "V32fV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
606 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f32_bf6, "V32fV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
607 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f16_fp6, "V32hV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
608 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf16_fp6, "V32yV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
609 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_f16_bf6, "V32hV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
610 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk32_bf16_bf6, "V32yV6Uif" , "nc" , "fp6bf6-cvt-scale-insts" ) |
611 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f16_fp8, "V2hUifIb" , "nc" , "fp8-cvt-scale-insts" ) |
612 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf16_fp8, "V2yUifIb" , "nc" , "fp8-cvt-scale-insts" ) |
613 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_f16_bf8, "V2hUifIb" , "nc" , "bf8-cvt-scale-insts" ) |
614 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_bf16_bf8, "V2yUifIb" , "nc" , "bf8-cvt-scale-insts" ) |
615 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp4_f16, "UiUiV2hfIi" , "nc" , "fp4-cvt-scale-insts" ) |
616 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_pk_fp4_bf16, "UiUiV2yfIi" , "nc" , "fp4-cvt-scale-insts" ) |
617 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f16, "UiUiV2hUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
618 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_bf16, "UiUiV2yUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
619 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32, "UiUiV2fUifIi" , "nc" , "fp4-cvt-scale-insts" ) |
620 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_bf16, "iiyUifIi" , "nc" , "bf8-cvt-scale-insts" ) |
621 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_f16, "iihUifIi" , "nc" , "bf8-cvt-scale-insts" ) |
622 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_bf8_f32, "iifUifIi" , "nc" , "bf8-cvt-scale-insts" ) |
623 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_bf16, "iiyUifIi" , "nc" , "fp8-cvt-scale-insts" ) |
624 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_f16, "iihUifIi" , "nc" , "fp8-cvt-scale-insts" ) |
625 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_fp8_f32, "iifUifIi" , "nc" , "fp8-cvt-scale-insts" ) |
626 | |
627 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_bf16, "V6UiV32yUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
628 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f16, "V6UiV32hUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
629 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_bf6_f32, "V6UiV32fUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
630 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_bf16, "V6UiV32yUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
631 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f16, "V6UiV32hUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
632 | TARGET_BUILTIN(__builtin_amdgcn_cvt_scalef32_sr_pk32_fp6_f32, "V6UiV32fUif" , "nc" , "f16bf16-to-fp6bf6-cvt-scale-insts" ) |
633 | TARGET_BUILTIN(__builtin_amdgcn_bitop3_b32, "iiiiIUi" , "nc" , "bitop3-insts" ) |
634 | TARGET_BUILTIN(__builtin_amdgcn_bitop3_b16, "ssssIUi" , "nc" , "bitop3-insts" ) |
635 | |
636 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf16_f32, "V2yV2yfUiIb" , "nc" , "f32-to-f16bf16-cvt-sr-insts" ) |
637 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_f16_f32, "V2hV2hfUiIb" , "nc" , "f32-to-f16bf16-cvt-sr-insts" ) |
638 | |
639 | #undef BUILTIN |
640 | #undef TARGET_BUILTIN |
641 | |