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