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_sad_u8, "UiUiUiUi" , "nc" ) |
144 | BUILTIN(__builtin_amdgcn_msad_u8, "UiUiUiUi" , "nc" ) |
145 | BUILTIN(__builtin_amdgcn_sad_hi_u8, "UiUiUiUi" , "nc" ) |
146 | BUILTIN(__builtin_amdgcn_sad_u16, "UiUiUiUi" , "nc" ) |
147 | BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi" , "nc" ) |
148 | BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi" , "nc" ) |
149 | BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui" , "nc" ) |
150 | |
151 | //===----------------------------------------------------------------------===// |
152 | // Ballot builtins. |
153 | //===----------------------------------------------------------------------===// |
154 | |
155 | TARGET_BUILTIN(__builtin_amdgcn_ballot_w32, "ZUib" , "nc" , "wavefrontsize32" ) |
156 | BUILTIN(__builtin_amdgcn_ballot_w64, "WUib" , "nc" ) |
157 | |
158 | // Deprecated intrinsics in favor of __builtin_amdgn_ballot_{w32|w64} |
159 | BUILTIN(__builtin_amdgcn_uicmp, "WUiUiUiIi" , "nc" ) |
160 | BUILTIN(__builtin_amdgcn_uicmpl, "WUiWUiWUiIi" , "nc" ) |
161 | BUILTIN(__builtin_amdgcn_sicmp, "WUiiiIi" , "nc" ) |
162 | BUILTIN(__builtin_amdgcn_sicmpl, "WUiWiWiIi" , "nc" ) |
163 | BUILTIN(__builtin_amdgcn_fcmp, "WUiddIi" , "nc" ) |
164 | BUILTIN(__builtin_amdgcn_fcmpf, "WUiffIi" , "nc" ) |
165 | |
166 | //===----------------------------------------------------------------------===// |
167 | // Flat addressing builtins. |
168 | //===----------------------------------------------------------------------===// |
169 | BUILTIN(__builtin_amdgcn_is_shared, "bvC*0" , "nc" ) |
170 | BUILTIN(__builtin_amdgcn_is_private, "bvC*0" , "nc" ) |
171 | |
172 | //===----------------------------------------------------------------------===// |
173 | // GWS builtins. |
174 | //===----------------------------------------------------------------------===// |
175 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_init, "vUiUi" , "n" , "gws" ) |
176 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_barrier, "vUiUi" , "n" , "gws" ) |
177 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_v, "vUi" , "n" , "gws" ) |
178 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_br, "vUiUi" , "n" , "gws" ) |
179 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_p, "vUi" , "n" , "gws" ) |
180 | |
181 | //===----------------------------------------------------------------------===// |
182 | // CI+ only builtins. |
183 | //===----------------------------------------------------------------------===// |
184 | TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v" , "n" , "ci-insts" ) |
185 | TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v" , "n" , "ci-insts" ) |
186 | TARGET_BUILTIN(__builtin_amdgcn_ds_gws_sema_release_all, "vUi" , "n" , "ci-insts" ) |
187 | |
188 | //===----------------------------------------------------------------------===// |
189 | // Interpolation builtins. |
190 | //===----------------------------------------------------------------------===// |
191 | BUILTIN(__builtin_amdgcn_interp_p1_f16, "ffUiUibUi" , "nc" ) |
192 | BUILTIN(__builtin_amdgcn_interp_p2_f16, "hffUiUibUi" , "nc" ) |
193 | BUILTIN(__builtin_amdgcn_interp_p1, "ffUiUiUi" , "nc" ) |
194 | BUILTIN(__builtin_amdgcn_interp_p2, "fffUiUiUi" , "nc" ) |
195 | BUILTIN(__builtin_amdgcn_interp_mov, "fUiUiUiUi" , "nc" ) |
196 | |
197 | //===----------------------------------------------------------------------===// |
198 | // VI+ only builtins. |
199 | //===----------------------------------------------------------------------===// |
200 | |
201 | TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh" , "nc" , "16-bit-insts" ) |
202 | TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh" , "nc" , "16-bit-insts" ) |
203 | TARGET_BUILTIN(__builtin_amdgcn_sqrth, "hh" , "nc" , "16-bit-insts" ) |
204 | TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh" , "nc" , "16-bit-insts" ) |
205 | TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh" , "nc" , "16-bit-insts" ) |
206 | TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh" , "nc" , "16-bit-insts" ) |
207 | TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi" , "nc" , "16-bit-insts" ) |
208 | TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh" , "nc" , "16-bit-insts" ) |
209 | TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh" , "nc" , "16-bit-insts" ) |
210 | TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh" , "nc" , "16-bit-insts" ) |
211 | TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi" , "nc" , "16-bit-insts" ) |
212 | TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "WUi" , "n" , "s-memrealtime" ) |
213 | TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb" , "nc" , "dpp" ) |
214 | TARGET_BUILTIN(__builtin_amdgcn_update_dpp, "iiiIiIiIiIb" , "nc" , "dpp" ) |
215 | TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v" , "n" , "gfx8-insts" ) |
216 | TARGET_BUILTIN(__builtin_amdgcn_perm, "UiUiUiUi" , "nc" , "gfx8-insts" ) |
217 | |
218 | //===----------------------------------------------------------------------===// |
219 | // GFX9+ only builtins. |
220 | //===----------------------------------------------------------------------===// |
221 | |
222 | TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh" , "nc" , "gfx9-insts" ) |
223 | |
224 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
225 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1f" , "t" , "atomic-fadd-rtn-insts" ) |
226 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2f16, "V2hV2h*1V2h" , "t" , "atomic-buffer-global-pk-add-f16-insts" ) |
227 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
228 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmax_f64, "dd*1d" , "t" , "gfx90a-insts" ) |
229 | |
230 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
231 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmin_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
232 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fmax_f64, "dd*0d" , "t" , "gfx90a-insts" ) |
233 | |
234 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f64, "dd*3d" , "t" , "gfx90a-insts" ) |
235 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_f32, "ff*3f" , "t" , "gfx8-insts" ) |
236 | |
237 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_f32, "ff*0f" , "t" , "gfx940-insts" ) |
238 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2f16, "V2hV2h*0V2h" , "t" , "atomic-flat-pk-add-16-insts" ) |
239 | TARGET_BUILTIN(__builtin_amdgcn_flat_atomic_fadd_v2bf16, "V2sV2s*0V2s" , "t" , "atomic-flat-pk-add-16-insts" ) |
240 | TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s" , "t" , "atomic-global-pk-add-bf16-inst" ) |
241 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s" , "t" , "atomic-ds-pk-add-16-insts" ) |
242 | TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2f16, "V2hV2h*3V2h" , "t" , "atomic-ds-pk-add-16-insts" ) |
243 | |
244 | //===----------------------------------------------------------------------===// |
245 | // Deep learning builtins. |
246 | //===----------------------------------------------------------------------===// |
247 | |
248 | TARGET_BUILTIN(__builtin_amdgcn_fdot2, "fV2hV2hfIb" , "nc" , "dot10-insts" ) |
249 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_f16_f16, "hV2hV2hh" , "nc" , "dot9-insts" ) |
250 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_bf16_bf16, "sV2sV2ss" , "nc" , "dot9-insts" ) |
251 | TARGET_BUILTIN(__builtin_amdgcn_fdot2_f32_bf16, "fV2sV2sfIb" , "nc" , "dot9-insts" ) |
252 | TARGET_BUILTIN(__builtin_amdgcn_sdot2, "SiV2SsV2SsSiIb" , "nc" , "dot2-insts" ) |
253 | TARGET_BUILTIN(__builtin_amdgcn_udot2, "UiV2UsV2UsUiIb" , "nc" , "dot2-insts" ) |
254 | TARGET_BUILTIN(__builtin_amdgcn_sdot4, "SiSiSiSiIb" , "nc" , "dot1-insts" ) |
255 | TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb" , "nc" , "dot7-insts" ) |
256 | TARGET_BUILTIN(__builtin_amdgcn_sudot4, "iIbiIbiiIb" , "nc" , "dot8-insts" ) |
257 | TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb" , "nc" , "dot1-insts" ) |
258 | TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb" , "nc" , "dot7-insts" ) |
259 | TARGET_BUILTIN(__builtin_amdgcn_sudot8, "iIbiIbiiIb" , "nc" , "dot8-insts" ) |
260 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_bf8, "fUiUif" , "nc" , "dot11-insts" ) |
261 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_fp8, "fUiUif" , "nc" , "dot11-insts" ) |
262 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_fp8_fp8, "fUiUif" , "nc" , "dot11-insts" ) |
263 | TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif" , "nc" , "dot11-insts" ) |
264 | |
265 | //===----------------------------------------------------------------------===// |
266 | // GFX10+ only builtins. |
267 | //===----------------------------------------------------------------------===// |
268 | TARGET_BUILTIN(__builtin_amdgcn_permlane16, "UiUiUiUiUiIbIb" , "nc" , "gfx10-insts" ) |
269 | TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb" , "nc" , "gfx10-insts" ) |
270 | TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi" , "nc" , "gfx10-insts" ) |
271 | TARGET_BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "vIs" , "n" , "gfx10-insts" ) |
272 | |
273 | //===----------------------------------------------------------------------===// |
274 | // Raytracing builtins. |
275 | // By default the 1st argument is i32 and the 4/5-th arguments are float4. |
276 | // Postfix l indicates the 1st argument is i64. |
277 | // Postfix h indicates the 4/5-th arguments are half4. |
278 | //===----------------------------------------------------------------------===// |
279 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray, "V4UiUifV4fV4fV4fV4Ui" , "nc" , "gfx10-insts" ) |
280 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_h, "V4UiUifV4fV4hV4hV4Ui" , "nc" , "gfx10-insts" ) |
281 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_l, "V4UiWUifV4fV4fV4fV4Ui" , "nc" , "gfx10-insts" ) |
282 | TARGET_BUILTIN(__builtin_amdgcn_image_bvh_intersect_ray_lh, "V4UiWUifV4fV4hV4hV4Ui" , "nc" , "gfx10-insts" ) |
283 | |
284 | |
285 | //===----------------------------------------------------------------------===// |
286 | // GFX11+ only builtins. |
287 | //===----------------------------------------------------------------------===// |
288 | |
289 | // TODO: This is a no-op in wave32. Should the builtin require wavefrontsize64? |
290 | TARGET_BUILTIN(__builtin_amdgcn_permlane64, "UiUi" , "nc" , "gfx11-insts" ) |
291 | TARGET_BUILTIN(__builtin_amdgcn_s_wait_event_export_ready, "v" , "n" , "gfx11-insts" ) |
292 | |
293 | //===----------------------------------------------------------------------===// |
294 | // WMMA builtins. |
295 | // Postfix w32 indicates the builtin requires wavefront size of 32. |
296 | // Postfix w64 indicates the builtin requires wavefront size of 64. |
297 | //===----------------------------------------------------------------------===// |
298 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32, "V8fV16hV16hV8f" , "nc" , "gfx11-insts,wavefrontsize32" ) |
299 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32, "V8fV16sV16sV8f" , "nc" , "gfx11-insts,wavefrontsize32" ) |
300 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32, "V16hV16hV16hV16hIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
301 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32, "V16sV16sV16sV16sIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
302 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32, "V16hV16hV16hV16hIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
303 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32, "V16sV16sV16sV16sIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
304 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32, "V8iIbV4iIbV4iV8iIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
305 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx11-insts,wavefrontsize32" ) |
306 | |
307 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64, "V4fV16hV16hV4f" , "nc" , "gfx11-insts,wavefrontsize64" ) |
308 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64, "V4fV16sV16sV4f" , "nc" , "gfx11-insts,wavefrontsize64" ) |
309 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64, "V8hV16hV16hV8hIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
310 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64, "V8sV16sV16sV8sIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
311 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64, "V8hV16hV16hV8hIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
312 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64, "V8sV16sV16sV8sIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
313 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64, "V4iIbV4iIbV4iV4iIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
314 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64, "V4iIbV2iIbV2iV4iIb" , "nc" , "gfx11-insts,wavefrontsize64" ) |
315 | |
316 | TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtn, "UiUIi" , "n" , "gfx11-insts" ) |
317 | TARGET_BUILTIN(__builtin_amdgcn_s_sendmsg_rtnl, "UWiUIi" , "n" , "gfx11-insts" ) |
318 | |
319 | TARGET_BUILTIN(__builtin_amdgcn_ds_bvh_stack_rtn, "V2UiUiUiV4UiIi" , "n" , "gfx11-insts" ) |
320 | |
321 | //===----------------------------------------------------------------------===// |
322 | // Special builtins. |
323 | //===----------------------------------------------------------------------===// |
324 | BUILTIN(__builtin_amdgcn_read_exec, "WUi" , "nc" ) |
325 | BUILTIN(__builtin_amdgcn_read_exec_lo, "Ui" , "nc" ) |
326 | BUILTIN(__builtin_amdgcn_read_exec_hi, "Ui" , "nc" ) |
327 | |
328 | BUILTIN(__builtin_amdgcn_endpgm, "v" , "nr" ) |
329 | |
330 | BUILTIN(__builtin_amdgcn_get_fpenv, "WUi" , "n" ) |
331 | BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi" , "n" ) |
332 | |
333 | //===----------------------------------------------------------------------===// |
334 | // R600-NI only builtins. |
335 | //===----------------------------------------------------------------------===// |
336 | |
337 | BUILTIN(__builtin_r600_implicitarg_ptr, "Uc*7" , "nc" ) |
338 | |
339 | BUILTIN(__builtin_r600_read_tgid_x, "Ui" , "nc" ) |
340 | BUILTIN(__builtin_r600_read_tgid_y, "Ui" , "nc" ) |
341 | BUILTIN(__builtin_r600_read_tgid_z, "Ui" , "nc" ) |
342 | |
343 | BUILTIN(__builtin_r600_read_tidig_x, "Ui" , "nc" ) |
344 | BUILTIN(__builtin_r600_read_tidig_y, "Ui" , "nc" ) |
345 | BUILTIN(__builtin_r600_read_tidig_z, "Ui" , "nc" ) |
346 | |
347 | BUILTIN(__builtin_r600_recipsqrt_ieee, "dd" , "nc" ) |
348 | BUILTIN(__builtin_r600_recipsqrt_ieeef, "ff" , "nc" ) |
349 | |
350 | //===----------------------------------------------------------------------===// |
351 | // MFMA builtins. |
352 | //===----------------------------------------------------------------------===// |
353 | |
354 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x1f32, "V32fffV32fIiIiIi" , "nc" , "mai-insts" ) |
355 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x1f32, "V16fffV16fIiIiIi" , "nc" , "mai-insts" ) |
356 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x1f32, "V4fffV4fIiIiIi" , "nc" , "mai-insts" ) |
357 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2f32, "V16fffV16fIiIiIi" , "nc" , "mai-insts" ) |
358 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f32, "V4fffV4fIiIiIi" , "nc" , "mai-insts" ) |
359 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4f16, "V32fV4hV4hV32fIiIiIi" , "nc" , "mai-insts" ) |
360 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4f16, "V16fV4hV4hV16fIiIiIi" , "nc" , "mai-insts" ) |
361 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x4f16, "V4fV4hV4hV4fIiIiIi" , "nc" , "mai-insts" ) |
362 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x8f16, "V16fV4hV4hV16fIiIiIi" , "nc" , "mai-insts" ) |
363 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x16f16, "V4fV4hV4hV4fIiIiIi" , "nc" , "mai-insts" ) |
364 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x4i8, "V32iiiV32iIiIiIi" , "nc" , "mai-insts" ) |
365 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x4i8, "V16iiiV16iIiIiIi" , "nc" , "mai-insts" ) |
366 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_4x4x4i8, "V4iiiV4iIiIiIi" , "nc" , "mai-insts" ) |
367 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x8i8, "V16iiiV16iIiIiIi" , "nc" , "mai-insts" ) |
368 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x16i8, "V4iiiV4iIiIiIi" , "nc" , "mai-insts" ) |
369 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x2bf16, "V32fV2sV2sV32fIiIiIi" , "nc" , "mai-insts" ) |
370 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x2bf16, "V16fV2sV2sV16fIiIiIi" , "nc" , "mai-insts" ) |
371 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x2bf16, "V4fV2sV2sV4fIiIiIi" , "nc" , "mai-insts" ) |
372 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4bf16, "V16fV2sV2sV16fIiIiIi" , "nc" , "mai-insts" ) |
373 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8bf16, "V4fV2sV2sV4fIiIiIi" , "nc" , "mai-insts" ) |
374 | |
375 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4bf16_1k, "V32fV4sV4sV32fIiIiIi" , "nc" , "mai-insts" ) |
376 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x4bf16_1k, "V16fV4sV4sV16fIiIiIi" , "nc" , "mai-insts" ) |
377 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_4x4x4bf16_1k, "V4fV4sV4sV4fIiIiIi" , "nc" , "mai-insts" ) |
378 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x8bf16_1k, "V16fV4sV4sV16fIiIiIi" , "nc" , "mai-insts" ) |
379 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x16bf16_1k, "V4fV4sV4sV4fIiIiIi" , "nc" , "mai-insts" ) |
380 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f64_16x16x4f64, "V4dddV4dIiIiIi" , "nc" , "mai-insts" ) |
381 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f64_4x4x4f64, "ddddIiIiIi" , "nc" , "mai-insts" ) |
382 | |
383 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_16x16x32_i8, "V4iWiWiV4iIiIiIi" , "nc" , "mai-insts" ) |
384 | TARGET_BUILTIN(__builtin_amdgcn_mfma_i32_32x32x16_i8, "V16iWiWiV16iIiIiIi" , "nc" , "mai-insts" ) |
385 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x8_xf32, "V4fV2fV2fV4fIiIiIi" , "nc" , "mai-insts" ) |
386 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x4_xf32, "V16fV2fV2fV16fIiIiIi" , "nc" , "mai-insts" ) |
387 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
388 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
389 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
390 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8, "V4fWiWiV4fIiIiIi" , "nc" , "fp8-insts" ) |
391 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
392 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
393 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
394 | TARGET_BUILTIN(__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8, "V16fWiWiV16fIiIiIi" , "nc" , "fp8-insts" ) |
395 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_f16, "V4fV4hV8hV4fiIiIi" , "nc" , "mai-insts" ) |
396 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_f16, "V16fV4hV8hV16fiIiIi" , "nc" , "mai-insts" ) |
397 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x32_bf16, "V4fV4sV8sV4fiIiIi" , "nc" , "mai-insts" ) |
398 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x16_bf16, "V16fV4sV8sV16fiIiIi" , "nc" , "mai-insts" ) |
399 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_16x16x64_i8, "V4iV2iV4iV4iiIiIi" , "nc" , "mai-insts" ) |
400 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_i32_32x32x32_i8, "V16iV2iV4iV16iiIiIi" , "nc" , "mai-insts" ) |
401 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
402 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
403 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
404 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8, "V4fV2iV4iV4fiIiIi" , "nc" , "fp8-insts" ) |
405 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf8_bf8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
406 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
407 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
408 | TARGET_BUILTIN(__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8, "V16fV2iV4iV16fiIiIi" , "nc" , "fp8-insts" ) |
409 | |
410 | TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_bf8, "fiIi" , "nc" , "fp8-conversion-insts" ) |
411 | TARGET_BUILTIN(__builtin_amdgcn_cvt_f32_fp8, "fiIi" , "nc" , "fp8-conversion-insts" ) |
412 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_bf8, "V2fiIb" , "nc" , "fp8-conversion-insts" ) |
413 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f32_fp8, "V2fiIb" , "nc" , "fp8-conversion-insts" ) |
414 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_bf8_f32, "iffiIb" , "nc" , "fp8-conversion-insts" ) |
415 | TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb" , "nc" , "fp8-conversion-insts" ) |
416 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi" , "nc" , "fp8-conversion-insts" ) |
417 | TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi" , "nc" , "fp8-conversion-insts" ) |
418 | |
419 | //===----------------------------------------------------------------------===// |
420 | // GFX12+ only builtins. |
421 | //===----------------------------------------------------------------------===// |
422 | |
423 | TARGET_BUILTIN(__builtin_amdgcn_s_sleep_var, "vUi" , "n" , "gfx12-insts" ) |
424 | TARGET_BUILTIN(__builtin_amdgcn_permlane16_var, "UiUiUiUiIbIb" , "nc" , "gfx12-insts" ) |
425 | TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb" , "nc" , "gfx12-insts" ) |
426 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi" , "n" , "gfx12-insts" ) |
427 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vi" , "n" , "gfx12-insts" ) |
428 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs" , "n" , "gfx12-insts" ) |
429 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi" , "n" , "gfx12-insts" ) |
430 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst_var, "bi" , "n" , "gfx12-insts" ) |
431 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vii" , "n" , "gfx12-insts" ) |
432 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vi" , "n" , "gfx12-insts" ) |
433 | TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vi" , "n" , "gfx12-insts" ) |
434 | TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "b" , "n" , "gfx12-insts" ) |
435 | TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii" , "n" , "gfx12-insts" ) |
436 | |
437 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b64_v2i32, "V2iV2i*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
438 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v8i16, "V8sV8s*1" , "nc" , "gfx12-insts,wavefrontsize32" ) |
439 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b64_i32, "ii*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
440 | TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16, "V4sV4s*1" , "nc" , "gfx12-insts,wavefrontsize64" ) |
441 | |
442 | //===----------------------------------------------------------------------===// |
443 | // WMMA builtins. |
444 | // Postfix w32 indicates the builtin requires wavefront size of 32. |
445 | // Postfix w64 indicates the builtin requires wavefront size of 64. |
446 | // |
447 | // Some of these are very similar to their GFX11 counterparts, but they don't |
448 | // require replication of the A,B matrices, so they use fewer vector elements. |
449 | // Therefore, we add an "_gfx12" suffix to distinguish them from the existing |
450 | // builtins. |
451 | //===----------------------------------------------------------------------===// |
452 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12, "V8fV8hV8hV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
453 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12, "V8fV8sV8sV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
454 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12, "V8hV8hV8hV8h" , "nc" , "gfx12-insts,wavefrontsize32" ) |
455 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12, "V8sV8sV8sV8s" , "nc" , "gfx12-insts,wavefrontsize32" ) |
456 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
457 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12, "V8iIbiIbiV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
458 | // These are gfx12-only, but for consistency with the other WMMA variants we're |
459 | // keeping the "_gfx12" suffix. |
460 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
461 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
462 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
463 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12, "V8fV2iV2iV8f" , "nc" , "gfx12-insts,wavefrontsize32" ) |
464 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12, "V8iIbV2iIbV2iV8iIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
465 | |
466 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12, "V4fV4hV4hV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
467 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12, "V4fV4sV4sV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
468 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12, "V4hV4hV4hV4h" , "nc" , "gfx12-insts,wavefrontsize64" ) |
469 | TARGET_BUILTIN(__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12, "V4sV4sV4sV4s" , "nc" , "gfx12-insts,wavefrontsize64" ) |
470 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
471 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
472 | // These are gfx12-only, but for consistency with the other WMMA variants we're |
473 | // keeping the "_gfx12" suffix. |
474 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
475 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
476 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
477 | TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12, "V4fiiV4f" , "nc" , "gfx12-insts,wavefrontsize64" ) |
478 | TARGET_BUILTIN(__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12, "V4iIbiIbiV4iIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
479 | |
480 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32, "V8fV8hV16hV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
481 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32, "V8fV8sV16sV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
482 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32, "V8hV8hV16hV8hs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
483 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32, "V8sV8sV16sV8ss" , "nc" , "gfx12-insts,wavefrontsize32" ) |
484 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32, "V8iIbV2iIbV4iV8isIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
485 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32, "V8iIbiIbV2iV8isIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
486 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32, "V8iIbV2iIbV4iV8isIb" , "nc" , "gfx12-insts,wavefrontsize32" ) |
487 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32, "V8fV2iV4iV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
488 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32, "V8fV2iV4iV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
489 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32, "V8fV2iV4iV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
490 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32, "V8fV2iV4iV8fs" , "nc" , "gfx12-insts,wavefrontsize32" ) |
491 | |
492 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64, "V4fV4hV8hV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
493 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64, "V4fV4sV8sV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
494 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64, "V4hV4hV8hV4hs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
495 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64, "V4sV4sV8sV4ss" , "nc" , "gfx12-insts,wavefrontsize64" ) |
496 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64, "V4iIbiIbV2iV4isIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
497 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64, "V4iIbiIbiV4isIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
498 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64, "V4iIbiIbV2iV4isIb" , "nc" , "gfx12-insts,wavefrontsize64" ) |
499 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64, "V4fiV2iV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
500 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64, "V4fiV2iV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
501 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64, "V4fiV2iV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
502 | TARGET_BUILTIN(__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64, "V4fiV2iV4fs" , "nc" , "gfx12-insts,wavefrontsize64" ) |
503 | |
504 | #undef BUILTIN |
505 | #undef TARGET_BUILTIN |
506 | |