clang 23.0.0git
CIRGenBuiltinAMDGPU.cpp
Go to the documentation of this file.
1//===---- CIRGenBuiltinAMDGPU.cpp - Emit CIR for AMDGPU builtins ----------===//
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 contains code to emit AMDGPU Builtin calls.
10//
11//===----------------------------------------------------------------------===//
12
13#include "CIRGenFunction.h"
14
15#include "mlir/IR/Value.h"
16#include "clang/Basic/TargetBuiltins.h"
17#include "llvm/Support/ErrorHandling.h"
18
19using namespace clang;
20using namespace clang::CIRGen;
21
22std::optional<mlir::Value>
23CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
24 const CallExpr *expr) {
25 switch (builtinId) {
26 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
27 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
28 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
29 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
30 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
31 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
32 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
33 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
34 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
35 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
36 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
37 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
38 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
39 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
40 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
41 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
42 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
43 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
44 cgm.errorNYI(expr->getSourceRange(),
45 std::string("unimplemented AMDGPU builtin call: ") +
46 getContext().BuiltinInfo.getName(builtinId));
47 return mlir::Value{};
48 }
49 case AMDGPU::BI__builtin_amdgcn_div_scale:
50 case AMDGPU::BI__builtin_amdgcn_div_scalef: {
51 cgm.errorNYI(expr->getSourceRange(),
52 std::string("unimplemented AMDGPU builtin call: ") +
53 getContext().BuiltinInfo.getName(builtinId));
54 return mlir::Value{};
55 }
56 case AMDGPU::BI__builtin_amdgcn_div_fmas:
57 case AMDGPU::BI__builtin_amdgcn_div_fmasf: {
58 cgm.errorNYI(expr->getSourceRange(),
59 std::string("unimplemented AMDGPU builtin call: ") +
60 getContext().BuiltinInfo.getName(builtinId));
61 return mlir::Value{};
62 }
63 case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
64 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
65 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
66 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
67 cgm.errorNYI(expr->getSourceRange(),
68 std::string("unimplemented AMDGPU builtin call: ") +
69 getContext().BuiltinInfo.getName(builtinId));
70 return mlir::Value{};
71 }
72 case AMDGPU::BI__builtin_amdgcn_permlane16:
73 case AMDGPU::BI__builtin_amdgcn_permlanex16:
74 case AMDGPU::BI__builtin_amdgcn_permlane64: {
75 cgm.errorNYI(expr->getSourceRange(),
76 std::string("unimplemented AMDGPU builtin call: ") +
77 getContext().BuiltinInfo.getName(builtinId));
78 return mlir::Value{};
79 }
80 case AMDGPU::BI__builtin_amdgcn_readlane:
81 case AMDGPU::BI__builtin_amdgcn_readfirstlane:
82 case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
83 cgm.errorNYI(expr->getSourceRange(),
84 std::string("unimplemented AMDGPU builtin call: ") +
85 getContext().BuiltinInfo.getName(builtinId));
86 return mlir::Value{};
87 }
88 case AMDGPU::BI__builtin_amdgcn_div_fixup:
89 case AMDGPU::BI__builtin_amdgcn_div_fixupf:
90 case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
91 cgm.errorNYI(expr->getSourceRange(),
92 std::string("unimplemented AMDGPU builtin call: ") +
93 getContext().BuiltinInfo.getName(builtinId));
94 return mlir::Value{};
95 }
96 case AMDGPU::BI__builtin_amdgcn_trig_preop:
97 case AMDGPU::BI__builtin_amdgcn_trig_preopf: {
98 cgm.errorNYI(expr->getSourceRange(),
99 std::string("unimplemented AMDGPU builtin call: ") +
100 getContext().BuiltinInfo.getName(builtinId));
101 return mlir::Value{};
102 }
103 case AMDGPU::BI__builtin_amdgcn_rcp:
104 case AMDGPU::BI__builtin_amdgcn_rcpf:
105 case AMDGPU::BI__builtin_amdgcn_rcph:
106 case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
107 cgm.errorNYI(expr->getSourceRange(),
108 std::string("unimplemented AMDGPU builtin call: ") +
109 getContext().BuiltinInfo.getName(builtinId));
110 return mlir::Value{};
111 }
112 case AMDGPU::BI__builtin_amdgcn_sqrt:
113 case AMDGPU::BI__builtin_amdgcn_sqrtf:
114 case AMDGPU::BI__builtin_amdgcn_sqrth:
115 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
116 cgm.errorNYI(expr->getSourceRange(),
117 std::string("unimplemented AMDGPU builtin call: ") +
118 getContext().BuiltinInfo.getName(builtinId));
119 return mlir::Value{};
120 }
121 case AMDGPU::BI__builtin_amdgcn_rsq:
122 case AMDGPU::BI__builtin_amdgcn_rsqf:
123 case AMDGPU::BI__builtin_amdgcn_rsqh:
124 case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
125 cgm.errorNYI(expr->getSourceRange(),
126 std::string("unimplemented AMDGPU builtin call: ") +
127 getContext().BuiltinInfo.getName(builtinId));
128 return mlir::Value{};
129 }
130 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
131 case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
132 cgm.errorNYI(expr->getSourceRange(),
133 std::string("unimplemented AMDGPU builtin call: ") +
134 getContext().BuiltinInfo.getName(builtinId));
135 return mlir::Value{};
136 }
137 case AMDGPU::BI__builtin_amdgcn_sinf:
138 case AMDGPU::BI__builtin_amdgcn_sinh:
139 case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
140 cgm.errorNYI(expr->getSourceRange(),
141 std::string("unimplemented AMDGPU builtin call: ") +
142 getContext().BuiltinInfo.getName(builtinId));
143 return mlir::Value{};
144 }
145 case AMDGPU::BI__builtin_amdgcn_cosf:
146 case AMDGPU::BI__builtin_amdgcn_cosh:
147 case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
148 cgm.errorNYI(expr->getSourceRange(),
149 std::string("unimplemented AMDGPU builtin call: ") +
150 getContext().BuiltinInfo.getName(builtinId));
151 return mlir::Value{};
152 }
153 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: {
154 cgm.errorNYI(expr->getSourceRange(),
155 std::string("unimplemented AMDGPU builtin call: ") +
156 getContext().BuiltinInfo.getName(builtinId));
157 return mlir::Value{};
158 }
159 case AMDGPU::BI__builtin_amdgcn_logf:
160 case AMDGPU::BI__builtin_amdgcn_log_bf16: {
161 cgm.errorNYI(expr->getSourceRange(),
162 std::string("unimplemented AMDGPU builtin call: ") +
163 getContext().BuiltinInfo.getName(builtinId));
164 return mlir::Value{};
165 }
166 case AMDGPU::BI__builtin_amdgcn_exp2f:
167 case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
168 cgm.errorNYI(expr->getSourceRange(),
169 std::string("unimplemented AMDGPU builtin call: ") +
170 getContext().BuiltinInfo.getName(builtinId));
171 return mlir::Value{};
172 }
173 case AMDGPU::BI__builtin_amdgcn_log_clampf: {
174 cgm.errorNYI(expr->getSourceRange(),
175 std::string("unimplemented AMDGPU builtin call: ") +
176 getContext().BuiltinInfo.getName(builtinId));
177 return mlir::Value{};
178 }
179 case AMDGPU::BI__builtin_amdgcn_ldexp:
180 case AMDGPU::BI__builtin_amdgcn_ldexpf:
181 case AMDGPU::BI__builtin_amdgcn_ldexph: {
182 cgm.errorNYI(expr->getSourceRange(),
183 std::string("unimplemented AMDGPU builtin call: ") +
184 getContext().BuiltinInfo.getName(builtinId));
185 return mlir::Value{};
186 }
187 case AMDGPU::BI__builtin_amdgcn_frexp_mant:
188 case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
189 case AMDGPU::BI__builtin_amdgcn_frexp_manth: {
190 cgm.errorNYI(expr->getSourceRange(),
191 std::string("unimplemented AMDGPU builtin call: ") +
192 getContext().BuiltinInfo.getName(builtinId));
193 return mlir::Value{};
194 }
195 case AMDGPU::BI__builtin_amdgcn_frexp_exp:
196 case AMDGPU::BI__builtin_amdgcn_frexp_expf:
197 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
198 cgm.errorNYI(expr->getSourceRange(),
199 std::string("unimplemented AMDGPU builtin call: ") +
200 getContext().BuiltinInfo.getName(builtinId));
201 return mlir::Value{};
202 }
203 case AMDGPU::BI__builtin_amdgcn_fract:
204 case AMDGPU::BI__builtin_amdgcn_fractf:
205 case AMDGPU::BI__builtin_amdgcn_fracth: {
206 cgm.errorNYI(expr->getSourceRange(),
207 std::string("unimplemented AMDGPU builtin call: ") +
208 getContext().BuiltinInfo.getName(builtinId));
209 return mlir::Value{};
210 }
211 case AMDGPU::BI__builtin_amdgcn_lerp: {
212 cgm.errorNYI(expr->getSourceRange(),
213 std::string("unimplemented AMDGPU builtin call: ") +
214 getContext().BuiltinInfo.getName(builtinId));
215 return mlir::Value{};
216 }
217 case AMDGPU::BI__builtin_amdgcn_ubfe: {
218 cgm.errorNYI(expr->getSourceRange(),
219 std::string("unimplemented AMDGPU builtin call: ") +
220 getContext().BuiltinInfo.getName(builtinId));
221 return mlir::Value{};
222 }
223 case AMDGPU::BI__builtin_amdgcn_sbfe: {
224 cgm.errorNYI(expr->getSourceRange(),
225 std::string("unimplemented AMDGPU builtin call: ") +
226 getContext().BuiltinInfo.getName(builtinId));
227 return mlir::Value{};
228 }
229 case AMDGPU::BI__builtin_amdgcn_ballot_w32:
230 case AMDGPU::BI__builtin_amdgcn_ballot_w64: {
231 cgm.errorNYI(expr->getSourceRange(),
232 std::string("unimplemented AMDGPU builtin call: ") +
233 getContext().BuiltinInfo.getName(builtinId));
234 return mlir::Value{};
235 }
236 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
237 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
238 cgm.errorNYI(expr->getSourceRange(),
239 std::string("unimplemented AMDGPU builtin call: ") +
240 getContext().BuiltinInfo.getName(builtinId));
241 return mlir::Value{};
242 }
243 case AMDGPU::BI__builtin_amdgcn_tanhf:
244 case AMDGPU::BI__builtin_amdgcn_tanhh:
245 case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
246 cgm.errorNYI(expr->getSourceRange(),
247 std::string("unimplemented AMDGPU builtin call: ") +
248 getContext().BuiltinInfo.getName(builtinId));
249 return mlir::Value{};
250 }
251 case AMDGPU::BI__builtin_amdgcn_uicmp:
252 case AMDGPU::BI__builtin_amdgcn_uicmpl:
253 case AMDGPU::BI__builtin_amdgcn_sicmp:
254 case AMDGPU::BI__builtin_amdgcn_sicmpl: {
255 cgm.errorNYI(expr->getSourceRange(),
256 std::string("unimplemented AMDGPU builtin call: ") +
257 getContext().BuiltinInfo.getName(builtinId));
258 return mlir::Value{};
259 }
260 case AMDGPU::BI__builtin_amdgcn_fcmp:
261 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
262 cgm.errorNYI(expr->getSourceRange(),
263 std::string("unimplemented AMDGPU builtin call: ") +
264 getContext().BuiltinInfo.getName(builtinId));
265 return mlir::Value{};
266 }
267 case AMDGPU::BI__builtin_amdgcn_class:
268 case AMDGPU::BI__builtin_amdgcn_classf:
269 case AMDGPU::BI__builtin_amdgcn_classh: {
270 cgm.errorNYI(expr->getSourceRange(),
271 std::string("unimplemented AMDGPU builtin call: ") +
272 getContext().BuiltinInfo.getName(builtinId));
273 return mlir::Value{};
274 }
275 case AMDGPU::BI__builtin_amdgcn_fmed3f:
276 case AMDGPU::BI__builtin_amdgcn_fmed3h: {
277 cgm.errorNYI(expr->getSourceRange(),
278 std::string("unimplemented AMDGPU builtin call: ") +
279 getContext().BuiltinInfo.getName(builtinId));
280 return mlir::Value{};
281 }
282 case AMDGPU::BI__builtin_amdgcn_ds_append:
283 case AMDGPU::BI__builtin_amdgcn_ds_consume: {
284 cgm.errorNYI(expr->getSourceRange(),
285 std::string("unimplemented AMDGPU builtin call: ") +
286 getContext().BuiltinInfo.getName(builtinId));
287 return mlir::Value{};
288 }
289 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
290 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
291 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
292 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
293 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
294 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
295 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
296 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
297 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
298 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
299 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
300 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
301 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
302 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: {
303 cgm.errorNYI(expr->getSourceRange(),
304 std::string("unimplemented AMDGPU builtin call: ") +
305 getContext().BuiltinInfo.getName(builtinId));
306 return mlir::Value{};
307 }
308 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
309 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
310 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
311 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
312 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
313 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: {
314 cgm.errorNYI(expr->getSourceRange(),
315 std::string("unimplemented AMDGPU builtin call: ") +
316 getContext().BuiltinInfo.getName(builtinId));
317 return mlir::Value{};
318 }
319 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
320 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
321 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
322 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
323 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
324 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {
325 cgm.errorNYI(expr->getSourceRange(),
326 std::string("unimplemented AMDGPU builtin call: ") +
327 getContext().BuiltinInfo.getName(builtinId));
328 return mlir::Value{};
329 }
330 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
331 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
332 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
333 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
334 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
335 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
336 cgm.errorNYI(expr->getSourceRange(),
337 std::string("unimplemented AMDGPU builtin call: ") +
338 getContext().BuiltinInfo.getName(builtinId));
339 return mlir::Value{};
340 }
341 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
342 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
343 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
344 cgm.errorNYI(expr->getSourceRange(),
345 std::string("unimplemented AMDGPU builtin call: ") +
346 getContext().BuiltinInfo.getName(builtinId));
347 return mlir::Value{};
348 }
349 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
350 cgm.errorNYI(expr->getSourceRange(),
351 std::string("unimplemented AMDGPU builtin call: ") +
352 getContext().BuiltinInfo.getName(builtinId));
353 return mlir::Value{};
354 }
355 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
356 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
357 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
358 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
359 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
360 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
361 cgm.errorNYI(expr->getSourceRange(),
362 std::string("unimplemented AMDGPU builtin call: ") +
363 getContext().BuiltinInfo.getName(builtinId));
364 return mlir::Value{};
365 }
366 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
367 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
368 cgm.errorNYI(expr->getSourceRange(),
369 std::string("unimplemented AMDGPU builtin call: ") +
370 getContext().BuiltinInfo.getName(builtinId));
371 return mlir::Value{};
372 }
373 case AMDGPU::BI__builtin_amdgcn_read_exec:
374 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
375 case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
376 cgm.errorNYI(expr->getSourceRange(),
377 std::string("unimplemented AMDGPU builtin call: ") +
378 getContext().BuiltinInfo.getName(builtinId));
379 return mlir::Value{};
380 }
381 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
382 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
383 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
384 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
385 cgm.errorNYI(expr->getSourceRange(),
386 std::string("unimplemented AMDGPU builtin call: ") +
387 getContext().BuiltinInfo.getName(builtinId));
388 return mlir::Value{};
389 }
390 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:
391 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
392 cgm.errorNYI(expr->getSourceRange(),
393 std::string("unimplemented AMDGPU builtin call: ") +
394 getContext().BuiltinInfo.getName(builtinId));
395 return mlir::Value{};
396 }
397 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:
398 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
399 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
400 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
401 cgm.errorNYI(expr->getSourceRange(),
402 std::string("unimplemented AMDGPU builtin call: ") +
403 getContext().BuiltinInfo.getName(builtinId));
404 return mlir::Value{};
405 }
406 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
407 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
408 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
409 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
410 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
411 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
412 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
413 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
414 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
415 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
416 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
417 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
418 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
419 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
420 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
421 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
422 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
423 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
424 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
425 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
426 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
427 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
428 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
429 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
430 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: {
431 cgm.errorNYI(expr->getSourceRange(),
432 std::string("unimplemented AMDGPU builtin call: ") +
433 getContext().BuiltinInfo.getName(builtinId));
434 return mlir::Value{};
435 }
436 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
437 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
438 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
439 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
440 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
441 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
442 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
443 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
444 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
445 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
446 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
447 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
448 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
449 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
450 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
451 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
452 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
453 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
454 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
455 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
456 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
457 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
458 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
459 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
460 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
461 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
462 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
463 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
464 cgm.errorNYI(expr->getSourceRange(),
465 std::string("unimplemented AMDGPU builtin call: ") +
466 getContext().BuiltinInfo.getName(builtinId));
467 return mlir::Value{};
468 }
469 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
470 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
471 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
472 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
473 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
474 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
475 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
476 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
477 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
478 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
479 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
480 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
481 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
482 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
483 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
484 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
485 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
486 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
487 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
488 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
489 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
490 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
491 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
492 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
493 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
494 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
495 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
496 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
497 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
498 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
499 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
500 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
501 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
502 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
503 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
504 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
505 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
506 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
507 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
508 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
509 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
510 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
511 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
512 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
513 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
514 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
515 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
516 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
517 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
518 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
519 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
520 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
521 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
522 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32: {
523 cgm.errorNYI(expr->getSourceRange(),
524 std::string("unimplemented AMDGPU builtin call: ") +
525 getContext().BuiltinInfo.getName(builtinId));
526 return mlir::Value{};
527 }
528 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {
529 cgm.errorNYI(expr->getSourceRange(),
530 std::string("unimplemented AMDGPU builtin call: ") +
531 getContext().BuiltinInfo.getName(builtinId));
532 return mlir::Value{};
533 }
534 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
535 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
536 cgm.errorNYI(expr->getSourceRange(),
537 std::string("unimplemented AMDGPU builtin call: ") +
538 getContext().BuiltinInfo.getName(builtinId));
539 return mlir::Value{};
540 }
541 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
542 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
543 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
544 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
545 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
546 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
547 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
548 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
549 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
550 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
551 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
552 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
553 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
554 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
555 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
556 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
557 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
558 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
559 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
560 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:
561 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:
562 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:
563 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:
564 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:
565 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:
566 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:
567 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:
568 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:
569 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:
570 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:
571 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:
572 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:
573 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:
574 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:
575 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:
576 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:
577 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:
578 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: {
579 cgm.errorNYI(expr->getSourceRange(),
580 std::string("unimplemented AMDGPU builtin call: ") +
581 getContext().BuiltinInfo.getName(builtinId));
582 return mlir::Value{};
583 }
584 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:
585 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:
586 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:
587 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:
588 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:
589 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:
590 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:
591 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:
592 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:
593 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:
594 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:
595 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:
596 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:
597 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:
598 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:
599 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:
600 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:
601 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:
602 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:
603 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:
604 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:
605 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {
606 cgm.errorNYI(expr->getSourceRange(),
607 std::string("unimplemented AMDGPU builtin call: ") +
608 getContext().BuiltinInfo.getName(builtinId));
609 return mlir::Value{};
610 }
611 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
612 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
613 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
614 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
615 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
616 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
617 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:
618 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:
619 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:
620 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:
621 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:
622 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:
623 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:
624 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:
625 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:
626 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:
627 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:
628 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:
629 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:
630 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:
631 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:
632 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:
633 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
634 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
635 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
636 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
637 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
638 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
639 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: {
640 cgm.errorNYI(expr->getSourceRange(),
641 std::string("unimplemented AMDGPU builtin call: ") +
642 getContext().BuiltinInfo.getName(builtinId));
643 return mlir::Value{};
644 }
645 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
646 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
647 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
648 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:
649 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:
650 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:
651 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:
652 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:
653 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:
654 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:
655 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:
656 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:
657 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:
658 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
659 cgm.errorNYI(expr->getSourceRange(),
660 std::string("unimplemented AMDGPU builtin call: ") +
661 getContext().BuiltinInfo.getName(builtinId));
662 return mlir::Value{};
663 }
664 // amdgcn workgroup size
665 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
666 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
667 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
668 cgm.errorNYI(expr->getSourceRange(),
669 std::string("unimplemented AMDGPU builtin call: ") +
670 getContext().BuiltinInfo.getName(builtinId));
671 return mlir::Value{};
672 }
673 case AMDGPU::BI__builtin_amdgcn_grid_size_x:
674 case AMDGPU::BI__builtin_amdgcn_grid_size_y:
675 case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
676 cgm.errorNYI(expr->getSourceRange(),
677 std::string("unimplemented AMDGPU builtin call: ") +
678 getContext().BuiltinInfo.getName(builtinId));
679 return mlir::Value{};
680 }
681 case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
682 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
683 cgm.errorNYI(expr->getSourceRange(),
684 std::string("unimplemented AMDGPU builtin call: ") +
685 getContext().BuiltinInfo.getName(builtinId));
686 return mlir::Value{};
687 }
688 case AMDGPU::BI__builtin_amdgcn_alignbit: {
689 cgm.errorNYI(expr->getSourceRange(),
690 std::string("unimplemented AMDGPU builtin call: ") +
691 getContext().BuiltinInfo.getName(builtinId));
692 return mlir::Value{};
693 }
694 case AMDGPU::BI__builtin_amdgcn_fence: {
695 cgm.errorNYI(expr->getSourceRange(),
696 std::string("unimplemented AMDGPU builtin call: ") +
697 getContext().BuiltinInfo.getName(builtinId));
698 return mlir::Value{};
699 }
700 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
701 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
702 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
703 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
704 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
705 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
706 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
707 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:
708 case AMDGPU::BI__builtin_amdgcn_ds_faddf:
709 case AMDGPU::BI__builtin_amdgcn_ds_fminf:
710 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
711 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
712 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
713 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
714 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
715 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
716 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
717 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
718 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
719 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
720 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
721 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
722 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
723 cgm.errorNYI(expr->getSourceRange(),
724 std::string("unimplemented AMDGPU builtin call: ") +
725 getContext().BuiltinInfo.getName(builtinId));
726 return mlir::Value{};
727 }
728 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
729 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
730 cgm.errorNYI(expr->getSourceRange(),
731 std::string("unimplemented AMDGPU builtin call: ") +
732 getContext().BuiltinInfo.getName(builtinId));
733 return mlir::Value{};
734 }
735 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:
736 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {
737 cgm.errorNYI(expr->getSourceRange(),
738 std::string("unimplemented AMDGPU builtin call: ") +
739 getContext().BuiltinInfo.getName(builtinId));
740 return mlir::Value{};
741 }
742 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
743 case AMDGPU::BI__builtin_amdgcn_bitop3_b16: {
744 cgm.errorNYI(expr->getSourceRange(),
745 std::string("unimplemented AMDGPU builtin call: ") +
746 getContext().BuiltinInfo.getName(builtinId));
747 return mlir::Value{};
748 }
749 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
750 cgm.errorNYI(expr->getSourceRange(),
751 std::string("unimplemented AMDGPU builtin call: ") +
752 getContext().BuiltinInfo.getName(builtinId));
753 return mlir::Value{};
754 }
755 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:
756 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
757 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
758 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
759 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
760 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: {
761 cgm.errorNYI(expr->getSourceRange(),
762 std::string("unimplemented AMDGPU builtin call: ") +
763 getContext().BuiltinInfo.getName(builtinId));
764 return mlir::Value{};
765 }
766 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
767 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
768 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
769 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
770 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
771 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
772 cgm.errorNYI(expr->getSourceRange(),
773 std::string("unimplemented AMDGPU builtin call: ") +
774 getContext().BuiltinInfo.getName(builtinId));
775 return mlir::Value{};
776 }
777 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
778 cgm.errorNYI(expr->getSourceRange(),
779 std::string("unimplemented AMDGPU builtin call: ") +
780 getContext().BuiltinInfo.getName(builtinId));
781 return mlir::Value{};
782 }
783 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
784 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
785 cgm.errorNYI(expr->getSourceRange(),
786 std::string("unimplemented AMDGPU builtin call: ") +
787 getContext().BuiltinInfo.getName(builtinId));
788 return mlir::Value{};
789 }
790 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
791 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
792 cgm.errorNYI(expr->getSourceRange(),
793 std::string("unimplemented AMDGPU builtin call: ") +
794 getContext().BuiltinInfo.getName(builtinId));
795 return mlir::Value{};
796 }
797 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
798 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
799 cgm.errorNYI(expr->getSourceRange(),
800 std::string("unimplemented AMDGPU builtin call: ") +
801 getContext().BuiltinInfo.getName(builtinId));
802 return mlir::Value{};
803 }
804 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
805 cgm.errorNYI(expr->getSourceRange(),
806 std::string("unimplemented AMDGPU builtin call: ") +
807 getContext().BuiltinInfo.getName(builtinId));
808 return mlir::Value{};
809 }
810 case Builtin::BIlogbf:
811 case Builtin::BI__builtin_logbf: {
812 cgm.errorNYI(expr->getSourceRange(),
813 std::string("unimplemented AMDGPU builtin call: ") +
814 getContext().BuiltinInfo.getName(builtinId));
815 return mlir::Value{};
816 }
817 case Builtin::BIscalbnf:
818 case Builtin::BI__builtin_scalbnf:
819 case Builtin::BIscalbn:
820 case Builtin::BI__builtin_scalbn: {
821 cgm.errorNYI(expr->getSourceRange(),
822 std::string("unimplemented AMDGPU builtin call: ") +
823 getContext().BuiltinInfo.getName(builtinId));
824 return mlir::Value{};
825 }
826 default:
827 return std::nullopt;
828 }
829}
TargetBuiltins.h
Enumerates target-specific builtins in their own namespaces within namespace clang.
clang::CIRGen::CIRGenFunction::emitAMDGPUBuiltinExpr
std::optional< mlir::Value > emitAMDGPUBuiltinExpr(unsigned builtinID, const CallExpr *expr)
Emit a call to an AMDGPU builtin function.
Definition CIRGenBuiltinAMDGPU.cpp:23
clang::CIRGen::CIRGenFunction::getContext
clang::ASTContext & getContext() const
Definition CIRGenFunction.h:215
clang::CallExpr
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition Expr.h:2946
clang::ast_matchers::expr
const internal::VariadicDynCastAllOfMatcher< Stmt, Expr > expr
Matches expressions.
Definition ASTMatchersInternal.cpp:937
clang
The JSON file list parser is used to communicate input to InstallAPI.
Definition CalledOnceCheck.h:17
Generated on for clang by doxygen 1.14.0