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/AMDGPUAddrSpace.h"
18#include "llvm/Support/ErrorHandling.h"
19
20using namespace clang;
21using namespace clang::CIRGen;
22using namespace cir;
23
24// Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the
25// result to match \p e's return type when needed.
26// If \p e is null, returns the raw AS-4 pointer.
27static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
28 const CallExpr *e = nullptr) {
29 CIRGenBuilderTy &builder = cgf.getBuilder();
30 mlir::Location loc =
31 e ? cgf.getLoc(e->getExprLoc()) : builder.getUnknownLoc();
32 // The intrinsic always returns a pointer in the constant AS.
33 mlir::Type retTy = cir::PointerType::get(
34 cir::VoidType::get(builder.getContext()),
35 cir::TargetAddressSpaceAttr::get(builder.getContext(),
36 llvm::AMDGPUAS::CONSTANT_ADDRESS));
37 mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr",
38 retTy, mlir::ValueRange{});
39 if (!e)
40 return call;
41 // Only cast when the caller-visible AS differs from the intrinsic's AS;
42 auto expectedPtrTy =
43 mlir::cast<cir::PointerType>(cgf.convertType(e->getType()));
44 auto callPtrTy = mlir::cast<cir::PointerType>(call.getType());
45 if (expectedPtrTy.getAddrSpace() == callPtrTy.getAddrSpace())
46 return call;
47 return builder.createAddrSpaceCast(loc, call, expectedPtrTy);
48}
49
50static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
51 CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName,
52 llvm::StringRef constrainedIntrinsicName) {
53 mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0));
54 mlir::Value src1 = cgf.emitScalarExpr(e->getArg(1));
55 mlir::Location loc = cgf.getLoc(e->getExprLoc());
56
57 CIRGenBuilderTy &builder = cgf.getBuilder();
58
59 CIRGenFunction::CIRGenFPOptionsRAII fpOptsRAII(cgf, e);
60
61 if (builder.getIsFPConstrained()) {
62 cgf.cgm.errorNYI(e->getSourceRange(),
63 "constrained FP intrinsic support is NYI.");
64 }
65
66 return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(),
67 mlir::ValueRange{src0, src1});
68}
69
70static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e,
71 const llvm::fltSemantics &fSem) {
72 CIRGenBuilderTy &builder = cgf.getBuilder();
73 mlir::Location loc = cgf.getLoc(e->getExprLoc());
74
75 mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0));
76 mlir::Type srcTy = src0.getType();
77 mlir::Type int32Ty = builder.getSInt32Ty();
78
79 cir::RecordType frExpResTy =
80 builder.getAnonRecordTy({srcTy, int32Ty}, false, false);
81
82 mlir::Value frExpResult = builder.emitIntrinsicCallOp(
83 loc, "frexp", frExpResTy, mlir::ValueRange{src0});
84
85 mlir::Value exp =
86 cir::ExtractMemberOp::create(builder, loc, int32Ty, frExpResult, 1);
87
88 mlir::Value negativeOne =
89 builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1));
90 mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne);
91
92 mlir::Value siToFp = cir::CastOp::create(
93 builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1);
94
95 mlir::Value fabs = cir::FAbsOp::create(builder, loc, srcTy, src0);
96
97 llvm::APFloat infVal = llvm::APFloat::getInf(fSem);
98 mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal));
99
100 mlir::Value fabsNegInf =
101 builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf);
102
103 mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs);
104
105 llvm::APFloat zeroValue = llvm::APFloat::getZero(fSem);
106 mlir::Value zero =
107 builder.getConstant(loc, cir::FPAttr::get(srcTy, zeroValue));
108
109 mlir::Value srcEqZero =
110 builder.createCompare(loc, cir::CmpOpKind::eq, src0, zero);
111
112 llvm::APFloat negInfVal = llvm::APFloat::getInf(fSem, true);
113 mlir::Value negInf =
114 builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal));
115
116 mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel);
117
118 return res;
119}
120
121static mlir::Value
122emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e,
123 llvm::StringRef intrinsicName,
124 bool isImageStore) {
125 auto &builder = cgf.getBuilder();
126
127 llvm::SmallVector<mlir::Value, 10> args;
128 for (unsigned i = 0, n = e->getNumArgs(); i < n; ++i)
129 args.push_back(cgf.emitScalarExpr(e->getArg(i)));
130
131 mlir::Type retTy = isImageStore ? cir::VoidType::get(builder.getContext())
132 : cgf.convertType(e->getType());
133
134 auto callOp = cir::LLVMIntrinsicCallOp::create(
135 builder, cgf.getLoc(e->getExprLoc()),
136 builder.getStringAttr(intrinsicName), retTy, args);
137
138 return callOp.getResult();
139}
140
141std::optional<mlir::Value>
142CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
143 const CallExpr *expr) {
144 switch (builtinId) {
145 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
146 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
147 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
148 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
149 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
150 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
151 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
152 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
153 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
154 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
155 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
156 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
157 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
158 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
159 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
160 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
161 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
162 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
163 cgm.errorNYI(expr->getSourceRange(),
164 std::string("unimplemented AMDGPU builtin call: ") +
165 getContext().BuiltinInfo.getName(builtinId));
166 return mlir::Value{};
167 }
168 case AMDGPU::BI__builtin_amdgcn_div_scale:
169 case AMDGPU::BI__builtin_amdgcn_div_scalef: {
170 Address flagOutPtr = emitPointerWithAlignment(expr->getArg(3));
171 llvm::StringRef intrinsicName = "amdgcn.div.scale";
172 mlir::Value x = emitScalarExpr(expr->getArg(0));
173 mlir::Value y = emitScalarExpr(expr->getArg(1));
174 mlir::Value z = emitScalarExpr(expr->getArg(2));
175
176 auto i1Ty = builder.getUIntNTy(1);
177 cir::RecordType resTy = builder.getAnonRecordTy(
178 {x.getType(), i1Ty}, /*packed=*/false, /*padded=*/false);
179
180 mlir::Value structResult =
181 cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()),
182 builder.getStringAttr(intrinsicName),
183 resTy, {x, y, z})
184 .getResult();
185
186 mlir::Value result = cir::ExtractMemberOp::create(
187 builder, getLoc(expr->getExprLoc()), x.getType(), structResult, 0);
188 mlir::Value flag = cir::ExtractMemberOp::create(
189 builder, getLoc(expr->getExprLoc()), i1Ty, structResult, 1);
190
191 mlir::Type flagType = flagOutPtr.getElementType();
192 mlir::Value flagToStore =
193 cir::CastOp::create(builder, getLoc(expr->getExprLoc()), flagType,
194 cir::CastKind::int_to_bool, flag);
195 builder.createStore(getLoc(expr->getExprLoc()), flagToStore, flagOutPtr);
196 return result;
197 }
198 case AMDGPU::BI__builtin_amdgcn_div_fmas:
199 case AMDGPU::BI__builtin_amdgcn_div_fmasf:
200 return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas")
201 .getValue();
202 case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
203 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ds.swizzle")
204 .getValue();
205 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
206 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
207 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
208 cgm.errorNYI(expr->getSourceRange(),
209 std::string("unimplemented AMDGPU builtin call: ") +
210 getContext().BuiltinInfo.getName(builtinId));
211 return mlir::Value{};
212 }
213 case AMDGPU::BI__builtin_amdgcn_permlane16:
214 case AMDGPU::BI__builtin_amdgcn_permlanex16:
215 case AMDGPU::BI__builtin_amdgcn_permlane64: {
216 cgm.errorNYI(expr->getSourceRange(),
217 std::string("unimplemented AMDGPU builtin call: ") +
218 getContext().BuiltinInfo.getName(builtinId));
219 return mlir::Value{};
220 }
221 case AMDGPU::BI__builtin_amdgcn_readlane:
222 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
223 .getValue();
224 case AMDGPU::BI__builtin_amdgcn_readfirstlane:
225 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
226 .getValue();
227 case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
228 cgm.errorNYI(expr->getSourceRange(),
229 std::string("unimplemented AMDGPU builtin call: ") +
230 getContext().BuiltinInfo.getName(builtinId));
231 return mlir::Value{};
232 }
233 case AMDGPU::BI__builtin_amdgcn_div_fixup:
234 case AMDGPU::BI__builtin_amdgcn_div_fixupf:
235 case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
236 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
237 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
238 mlir::Value src2 = emitScalarExpr(expr->getArg(2));
239 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
240 "amdgcn.div.fixup", src0.getType(),
241 mlir::ValueRange{src0, src1, src2});
242 }
243 case AMDGPU::BI__builtin_amdgcn_trig_preop:
244 case AMDGPU::BI__builtin_amdgcn_trig_preopf: {
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_rcp:
251 case AMDGPU::BI__builtin_amdgcn_rcpf:
252 case AMDGPU::BI__builtin_amdgcn_rcph:
253 case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
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_sqrt:
260 case AMDGPU::BI__builtin_amdgcn_sqrtf:
261 case AMDGPU::BI__builtin_amdgcn_sqrth:
262 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
263 cgm.errorNYI(expr->getSourceRange(),
264 std::string("unimplemented AMDGPU builtin call: ") +
265 getContext().BuiltinInfo.getName(builtinId));
266 return mlir::Value{};
267 }
268 case AMDGPU::BI__builtin_amdgcn_rsq:
269 case AMDGPU::BI__builtin_amdgcn_rsqf:
270 case AMDGPU::BI__builtin_amdgcn_rsqh:
271 case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
272 cgm.errorNYI(expr->getSourceRange(),
273 std::string("unimplemented AMDGPU builtin call: ") +
274 getContext().BuiltinInfo.getName(builtinId));
275 return mlir::Value{};
276 }
277 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
278 case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
279 cgm.errorNYI(expr->getSourceRange(),
280 std::string("unimplemented AMDGPU builtin call: ") +
281 getContext().BuiltinInfo.getName(builtinId));
282 return mlir::Value{};
283 }
284 case AMDGPU::BI__builtin_amdgcn_sinf:
285 case AMDGPU::BI__builtin_amdgcn_sinh:
286 case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
287 cgm.errorNYI(expr->getSourceRange(),
288 std::string("unimplemented AMDGPU builtin call: ") +
289 getContext().BuiltinInfo.getName(builtinId));
290 return mlir::Value{};
291 }
292 case AMDGPU::BI__builtin_amdgcn_cosf:
293 case AMDGPU::BI__builtin_amdgcn_cosh:
294 case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
295 cgm.errorNYI(expr->getSourceRange(),
296 std::string("unimplemented AMDGPU builtin call: ") +
297 getContext().BuiltinInfo.getName(builtinId));
298 return mlir::Value{};
299 }
300 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
301 return emitAMDGPUDispatchPtr(*this, expr);
302 case AMDGPU::BI__builtin_amdgcn_logf:
303 case AMDGPU::BI__builtin_amdgcn_log_bf16: {
304 cgm.errorNYI(expr->getSourceRange(),
305 std::string("unimplemented AMDGPU builtin call: ") +
306 getContext().BuiltinInfo.getName(builtinId));
307 return mlir::Value{};
308 }
309 case AMDGPU::BI__builtin_amdgcn_exp2f:
310 case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
311 cgm.errorNYI(expr->getSourceRange(),
312 std::string("unimplemented AMDGPU builtin call: ") +
313 getContext().BuiltinInfo.getName(builtinId));
314 return mlir::Value{};
315 }
316 case AMDGPU::BI__builtin_amdgcn_log_clampf: {
317 cgm.errorNYI(expr->getSourceRange(),
318 std::string("unimplemented AMDGPU builtin call: ") +
319 getContext().BuiltinInfo.getName(builtinId));
320 return mlir::Value{};
321 }
322 case AMDGPU::BI__builtin_amdgcn_ldexp:
323 case AMDGPU::BI__builtin_amdgcn_ldexpf:
324 case AMDGPU::BI__builtin_amdgcn_ldexph: {
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_frexp_mant:
331 case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
332 case AMDGPU::BI__builtin_amdgcn_frexp_manth: {
333 cgm.errorNYI(expr->getSourceRange(),
334 std::string("unimplemented AMDGPU builtin call: ") +
335 getContext().BuiltinInfo.getName(builtinId));
336 return mlir::Value{};
337 }
338 case AMDGPU::BI__builtin_amdgcn_frexp_exp:
339 case AMDGPU::BI__builtin_amdgcn_frexp_expf:
340 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
341 cgm.errorNYI(expr->getSourceRange(),
342 std::string("unimplemented AMDGPU builtin call: ") +
343 getContext().BuiltinInfo.getName(builtinId));
344 return mlir::Value{};
345 }
346 case AMDGPU::BI__builtin_amdgcn_fract:
347 case AMDGPU::BI__builtin_amdgcn_fractf:
348 case AMDGPU::BI__builtin_amdgcn_fracth: {
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_lerp: {
355 cgm.errorNYI(expr->getSourceRange(),
356 std::string("unimplemented AMDGPU builtin call: ") +
357 getContext().BuiltinInfo.getName(builtinId));
358 return mlir::Value{};
359 }
360 case AMDGPU::BI__builtin_amdgcn_ubfe: {
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_sbfe: {
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_ballot_w32:
373 case AMDGPU::BI__builtin_amdgcn_ballot_w64: {
374 cgm.errorNYI(expr->getSourceRange(),
375 std::string("unimplemented AMDGPU builtin call: ") +
376 getContext().BuiltinInfo.getName(builtinId));
377 return mlir::Value{};
378 }
379 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
380 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
381 cgm.errorNYI(expr->getSourceRange(),
382 std::string("unimplemented AMDGPU builtin call: ") +
383 getContext().BuiltinInfo.getName(builtinId));
384 return mlir::Value{};
385 }
386 case AMDGPU::BI__builtin_amdgcn_tanhf:
387 case AMDGPU::BI__builtin_amdgcn_tanhh:
388 case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
389 cgm.errorNYI(expr->getSourceRange(),
390 std::string("unimplemented AMDGPU builtin call: ") +
391 getContext().BuiltinInfo.getName(builtinId));
392 return mlir::Value{};
393 }
394 case AMDGPU::BI__builtin_amdgcn_uicmp:
395 case AMDGPU::BI__builtin_amdgcn_uicmpl:
396 case AMDGPU::BI__builtin_amdgcn_sicmp:
397 case AMDGPU::BI__builtin_amdgcn_sicmpl: {
398 cgm.errorNYI(expr->getSourceRange(),
399 std::string("unimplemented AMDGPU builtin call: ") +
400 getContext().BuiltinInfo.getName(builtinId));
401 return mlir::Value{};
402 }
403 case AMDGPU::BI__builtin_amdgcn_fcmp:
404 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
405 cgm.errorNYI(expr->getSourceRange(),
406 std::string("unimplemented AMDGPU builtin call: ") +
407 getContext().BuiltinInfo.getName(builtinId));
408 return mlir::Value{};
409 }
410 case AMDGPU::BI__builtin_amdgcn_class:
411 case AMDGPU::BI__builtin_amdgcn_classf:
412 case AMDGPU::BI__builtin_amdgcn_classh: {
413 cgm.errorNYI(expr->getSourceRange(),
414 std::string("unimplemented AMDGPU builtin call: ") +
415 getContext().BuiltinInfo.getName(builtinId));
416 return mlir::Value{};
417 }
418 case AMDGPU::BI__builtin_amdgcn_fmed3f:
419 case AMDGPU::BI__builtin_amdgcn_fmed3h: {
420 cgm.errorNYI(expr->getSourceRange(),
421 std::string("unimplemented AMDGPU builtin call: ") +
422 getContext().BuiltinInfo.getName(builtinId));
423 return mlir::Value{};
424 }
425 case AMDGPU::BI__builtin_amdgcn_ds_append:
426 case AMDGPU::BI__builtin_amdgcn_ds_consume: {
427 cgm.errorNYI(expr->getSourceRange(),
428 std::string("unimplemented AMDGPU builtin call: ") +
429 getContext().BuiltinInfo.getName(builtinId));
430 return mlir::Value{};
431 }
432 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
433 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
434 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
435 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
436 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
437 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
438 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
439 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
440 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
441 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
442 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
443 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
444 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
445 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: {
446 cgm.errorNYI(expr->getSourceRange(),
447 std::string("unimplemented AMDGPU builtin call: ") +
448 getContext().BuiltinInfo.getName(builtinId));
449 return mlir::Value{};
450 }
451 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
452 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
453 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
454 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
455 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
456 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: {
457 cgm.errorNYI(expr->getSourceRange(),
458 std::string("unimplemented AMDGPU builtin call: ") +
459 getContext().BuiltinInfo.getName(builtinId));
460 return mlir::Value{};
461 }
462 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
463 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
464 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
465 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
466 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
467 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {
468 cgm.errorNYI(expr->getSourceRange(),
469 std::string("unimplemented AMDGPU builtin call: ") +
470 getContext().BuiltinInfo.getName(builtinId));
471 return mlir::Value{};
472 }
473 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
474 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
475 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
476 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
477 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
478 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
479 cgm.errorNYI(expr->getSourceRange(),
480 std::string("unimplemented AMDGPU builtin call: ") +
481 getContext().BuiltinInfo.getName(builtinId));
482 return mlir::Value{};
483 }
484 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
485 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
486 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
487 cgm.errorNYI(expr->getSourceRange(),
488 std::string("unimplemented AMDGPU builtin call: ") +
489 getContext().BuiltinInfo.getName(builtinId));
490 return mlir::Value{};
491 }
492 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
493 cgm.errorNYI(expr->getSourceRange(),
494 std::string("unimplemented AMDGPU builtin call: ") +
495 getContext().BuiltinInfo.getName(builtinId));
496 return mlir::Value{};
497 }
498 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
499 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
500 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
501 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
502 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
503 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
504 cgm.errorNYI(expr->getSourceRange(),
505 std::string("unimplemented AMDGPU builtin call: ") +
506 getContext().BuiltinInfo.getName(builtinId));
507 return mlir::Value{};
508 }
509 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
510 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
511 cgm.errorNYI(expr->getSourceRange(),
512 std::string("unimplemented AMDGPU builtin call: ") +
513 getContext().BuiltinInfo.getName(builtinId));
514 return mlir::Value{};
515 }
516 case AMDGPU::BI__builtin_amdgcn_read_exec:
517 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
518 case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
519 cgm.errorNYI(expr->getSourceRange(),
520 std::string("unimplemented AMDGPU builtin call: ") +
521 getContext().BuiltinInfo.getName(builtinId));
522 return mlir::Value{};
523 }
524 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
525 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
526 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
527 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
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_image_bvh8_intersect_ray:
534 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
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_ds_bvh_stack_rtn:
541 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
542 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
543 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
544 cgm.errorNYI(expr->getSourceRange(),
545 std::string("unimplemented AMDGPU builtin call: ") +
546 getContext().BuiltinInfo.getName(builtinId));
547 return mlir::Value{};
548 }
549 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
550 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
551 return emitAMDGCNImageOverloadedReturnType(*this, expr,
552 "amdgcn.image.load.1d", false);
553 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
554 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
555 return emitAMDGCNImageOverloadedReturnType(
556 *this, expr, "amdgcn.image.load.1darray", false);
557 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
558 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
559 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
560 return emitAMDGCNImageOverloadedReturnType(*this, expr,
561 "amdgcn.image.load.2d", false);
562 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
563 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
564 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
565 return emitAMDGCNImageOverloadedReturnType(
566 *this, expr, "amdgcn.image.load.2darray", false);
567 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
568 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
569 return emitAMDGCNImageOverloadedReturnType(*this, expr,
570 "amdgcn.image.load.3d", false);
571 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
572 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
573 return emitAMDGCNImageOverloadedReturnType(*this, expr,
574 "amdgcn.image.load.cube", false);
575 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
576 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
577 return emitAMDGCNImageOverloadedReturnType(
578 *this, expr, "amdgcn.image.load.mip.1d", false);
579 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
580 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
581 return emitAMDGCNImageOverloadedReturnType(
582 *this, expr, "amdgcn.image.load.mip.1darray", false);
583 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
584 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
585 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
586 return emitAMDGCNImageOverloadedReturnType(
587 *this, expr, "amdgcn.image.load.mip.2d", false);
588 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
589 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
590 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
591 return emitAMDGCNImageOverloadedReturnType(
592 *this, expr, "amdgcn.image.load.mip.2darray", false);
593 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
594 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
595 return emitAMDGCNImageOverloadedReturnType(
596 *this, expr, "amdgcn.image.load.mip.3d", false);
597 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
598 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
599 return emitAMDGCNImageOverloadedReturnType(
600 *this, expr, "amdgcn.image.load.mip.cube", false);
601 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
602 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
603 return emitAMDGCNImageOverloadedReturnType(*this, expr,
604 "amdgcn.image.store.1d", true);
605 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
606 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
607 return emitAMDGCNImageOverloadedReturnType(
608 *this, expr, "amdgcn.image.store.1darray", true);
609 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
610 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
611 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
612 return emitAMDGCNImageOverloadedReturnType(*this, expr,
613 "amdgcn.image.store.2d", true);
614 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
615 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
616 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
617 return emitAMDGCNImageOverloadedReturnType(
618 *this, expr, "amdgcn.image.store.2darray", true);
619 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
620 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
621 return emitAMDGCNImageOverloadedReturnType(*this, expr,
622 "amdgcn.image.store.3d", true);
623 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
624 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
625 return emitAMDGCNImageOverloadedReturnType(*this, expr,
626 "amdgcn.image.store.cube", true);
627 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
628 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
629 return emitAMDGCNImageOverloadedReturnType(
630 *this, expr, "amdgcn.image.store.mip.1d", true);
631 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
632 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
633 return emitAMDGCNImageOverloadedReturnType(
634 *this, expr, "amdgcn.image.store.mip.1darray", true);
635 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
636 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
637 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
638 return emitAMDGCNImageOverloadedReturnType(
639 *this, expr, "amdgcn.image.store.mip.2d", true);
640 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
641 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
642 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
643 return emitAMDGCNImageOverloadedReturnType(
644 *this, expr, "amdgcn.image.store.mip.2darray", true);
645 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
646 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
647 return emitAMDGCNImageOverloadedReturnType(
648 *this, expr, "amdgcn.image.store.mip.3d", true);
649 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
650 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
651 return emitAMDGCNImageOverloadedReturnType(
652 *this, expr, "amdgcn.image.store.mip.cube", true);
653 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
654 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
655 return emitAMDGCNImageOverloadedReturnType(*this, expr,
656 "amdgcn.image.sample.1d", false);
657 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
658 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
659 return emitAMDGCNImageOverloadedReturnType(
660 *this, expr, "amdgcn.image.sample.1darray", false);
661 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
662 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
663 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
664 return emitAMDGCNImageOverloadedReturnType(*this, expr,
665 "amdgcn.image.sample.2d", false);
666 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
667 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
668 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
669 return emitAMDGCNImageOverloadedReturnType(
670 *this, expr, "amdgcn.image.sample.2darray", false);
671 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
672 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
673 return emitAMDGCNImageOverloadedReturnType(*this, expr,
674 "amdgcn.image.sample.3d", false);
675 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
676 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
677 return emitAMDGCNImageOverloadedReturnType(
678 *this, expr, "amdgcn.image.sample.cube", false);
679 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
680 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
681 return emitAMDGCNImageOverloadedReturnType(
682 *this, expr, "amdgcn.image.sample.lz.1d", false);
683 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
684 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
685 return emitAMDGCNImageOverloadedReturnType(
686 *this, expr, "amdgcn.image.sample.l.1d", false);
687 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
688 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
689 return emitAMDGCNImageOverloadedReturnType(
690 *this, expr, "amdgcn.image.sample.d.1d", false);
691 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
692 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
693 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
694 return emitAMDGCNImageOverloadedReturnType(
695 *this, expr, "amdgcn.image.sample.lz.2d", false);
696 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
697 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
698 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
699 return emitAMDGCNImageOverloadedReturnType(
700 *this, expr, "amdgcn.image.sample.l.2d", false);
701 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
702 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
703 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
704 return emitAMDGCNImageOverloadedReturnType(
705 *this, expr, "amdgcn.image.sample.d.2d", false);
706 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
707 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
708 return emitAMDGCNImageOverloadedReturnType(
709 *this, expr, "amdgcn.image.sample.lz.3d", false);
710 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
711 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
712 return emitAMDGCNImageOverloadedReturnType(
713 *this, expr, "amdgcn.image.sample.l.3d", false);
714 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
715 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
716 return emitAMDGCNImageOverloadedReturnType(
717 *this, expr, "amdgcn.image.sample.d.3d", false);
718 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
719 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
720 return emitAMDGCNImageOverloadedReturnType(
721 *this, expr, "amdgcn.image.sample.lz.cube", false);
722 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
723 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
724 return emitAMDGCNImageOverloadedReturnType(
725 *this, expr, "amdgcn.image.sample.l.cube", false);
726 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
727 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
728 return emitAMDGCNImageOverloadedReturnType(
729 *this, expr, "amdgcn.image.sample.lz.1darray", false);
730 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
731 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
732 return emitAMDGCNImageOverloadedReturnType(
733 *this, expr, "amdgcn.image.sample.l.1darray", false);
734 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
735 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
736 return emitAMDGCNImageOverloadedReturnType(
737 *this, expr, "amdgcn.image.sample.d.1darray", false);
738 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
739 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
740 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
741 return emitAMDGCNImageOverloadedReturnType(
742 *this, expr, "amdgcn.image.sample.lz.2darray", false);
743 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
744 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
745 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
746 return emitAMDGCNImageOverloadedReturnType(
747 *this, expr, "amdgcn.image.sample.l.2darray", false);
748 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
749 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
750 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
751 return emitAMDGCNImageOverloadedReturnType(
752 *this, expr, "amdgcn.image.sample.d.2darray", false);
753 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
754 return emitAMDGCNImageOverloadedReturnType(
755 *this, expr, "amdgcn.image.gather4.lz.2d", false);
756 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
757 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
758 cgm.errorNYI(expr->getSourceRange(),
759 std::string("unimplemented AMDGPU builtin call: ") +
760 getContext().BuiltinInfo.getName(builtinId));
761 return mlir::Value{};
762 }
763 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
764 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
765 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
766 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
767 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
768 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
769 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
770 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
771 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
772 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
773 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
774 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
775 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
776 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
777 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
778 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
779 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
780 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
781 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
782 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:
783 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:
784 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:
785 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:
786 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:
787 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:
788 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:
789 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:
790 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:
791 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:
792 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:
793 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:
794 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:
795 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:
796 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:
797 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:
798 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:
799 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:
800 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: {
801 cgm.errorNYI(expr->getSourceRange(),
802 std::string("unimplemented AMDGPU builtin call: ") +
803 getContext().BuiltinInfo.getName(builtinId));
804 return mlir::Value{};
805 }
806 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:
807 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:
808 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:
809 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:
810 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:
811 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:
812 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:
813 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:
814 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:
815 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:
816 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:
817 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:
818 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:
819 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:
820 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:
821 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:
822 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:
823 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:
824 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:
825 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:
826 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:
827 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {
828 cgm.errorNYI(expr->getSourceRange(),
829 std::string("unimplemented AMDGPU builtin call: ") +
830 getContext().BuiltinInfo.getName(builtinId));
831 return mlir::Value{};
832 }
833 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
834 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
835 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
836 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
837 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
838 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
839 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:
840 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:
841 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:
842 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:
843 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:
844 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:
845 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:
846 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:
847 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:
848 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:
849 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:
850 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:
851 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:
852 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:
853 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:
854 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:
855 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
856 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
857 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
858 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
859 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
860 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
861 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: {
862 cgm.errorNYI(expr->getSourceRange(),
863 std::string("unimplemented AMDGPU builtin call: ") +
864 getContext().BuiltinInfo.getName(builtinId));
865 return mlir::Value{};
866 }
867 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
868 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
869 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
870 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:
871 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:
872 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:
873 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:
874 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:
875 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:
876 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:
877 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:
878 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:
879 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:
880 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
881 cgm.errorNYI(expr->getSourceRange(),
882 std::string("unimplemented AMDGPU builtin call: ") +
883 getContext().BuiltinInfo.getName(builtinId));
884 return mlir::Value{};
885 }
886 // amdgcn workgroup size
887 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
888 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
889 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
890 cgm.errorNYI(expr->getSourceRange(),
891 std::string("unimplemented AMDGPU builtin call: ") +
892 getContext().BuiltinInfo.getName(builtinId));
893 return mlir::Value{};
894 }
895 case AMDGPU::BI__builtin_amdgcn_grid_size_x:
896 case AMDGPU::BI__builtin_amdgcn_grid_size_y:
897 case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
898 cgm.errorNYI(expr->getSourceRange(),
899 std::string("unimplemented AMDGPU builtin call: ") +
900 getContext().BuiltinInfo.getName(builtinId));
901 return mlir::Value{};
902 }
903 case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
904 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
905 cgm.errorNYI(expr->getSourceRange(),
906 std::string("unimplemented AMDGPU builtin call: ") +
907 getContext().BuiltinInfo.getName(builtinId));
908 return mlir::Value{};
909 }
910 case AMDGPU::BI__builtin_amdgcn_alignbit: {
911 cgm.errorNYI(expr->getSourceRange(),
912 std::string("unimplemented AMDGPU builtin call: ") +
913 getContext().BuiltinInfo.getName(builtinId));
914 return mlir::Value{};
915 }
916 case AMDGPU::BI__builtin_amdgcn_fence: {
917 cgm.errorNYI(expr->getSourceRange(),
918 std::string("unimplemented AMDGPU builtin call: ") +
919 getContext().BuiltinInfo.getName(builtinId));
920 return mlir::Value{};
921 }
922 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
923 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
924 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
925 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
926 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
927 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
928 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
929 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:
930 case AMDGPU::BI__builtin_amdgcn_ds_faddf:
931 case AMDGPU::BI__builtin_amdgcn_ds_fminf:
932 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
933 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
934 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
935 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
936 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
937 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
938 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
939 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
940 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
941 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
942 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
943 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
944 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
945 cgm.errorNYI(expr->getSourceRange(),
946 std::string("unimplemented AMDGPU builtin call: ") +
947 getContext().BuiltinInfo.getName(builtinId));
948 return mlir::Value{};
949 }
950 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
951 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
952 cgm.errorNYI(expr->getSourceRange(),
953 std::string("unimplemented AMDGPU builtin call: ") +
954 getContext().BuiltinInfo.getName(builtinId));
955 return mlir::Value{};
956 }
957 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:
958 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {
959 cgm.errorNYI(expr->getSourceRange(),
960 std::string("unimplemented AMDGPU builtin call: ") +
961 getContext().BuiltinInfo.getName(builtinId));
962 return mlir::Value{};
963 }
964 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
965 case AMDGPU::BI__builtin_amdgcn_bitop3_b16: {
966 cgm.errorNYI(expr->getSourceRange(),
967 std::string("unimplemented AMDGPU builtin call: ") +
968 getContext().BuiltinInfo.getName(builtinId));
969 return mlir::Value{};
970 }
971 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
972 cgm.errorNYI(expr->getSourceRange(),
973 std::string("unimplemented AMDGPU builtin call: ") +
974 getContext().BuiltinInfo.getName(builtinId));
975 return mlir::Value{};
976 }
977 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:
978 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
979 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
980 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
981 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
982 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: {
983 cgm.errorNYI(expr->getSourceRange(),
984 std::string("unimplemented AMDGPU builtin call: ") +
985 getContext().BuiltinInfo.getName(builtinId));
986 return mlir::Value{};
987 }
988 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
989 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
990 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
991 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
992 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
993 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
994 cgm.errorNYI(expr->getSourceRange(),
995 std::string("unimplemented AMDGPU builtin call: ") +
996 getContext().BuiltinInfo.getName(builtinId));
997 return mlir::Value{};
998 }
999 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
1000 cgm.errorNYI(expr->getSourceRange(),
1001 std::string("unimplemented AMDGPU builtin call: ") +
1002 getContext().BuiltinInfo.getName(builtinId));
1003 return mlir::Value{};
1004 }
1005 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
1006 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
1007 cgm.errorNYI(expr->getSourceRange(),
1008 std::string("unimplemented AMDGPU builtin call: ") +
1009 getContext().BuiltinInfo.getName(builtinId));
1010 return mlir::Value{};
1011 }
1012 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
1013 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
1014 cgm.errorNYI(expr->getSourceRange(),
1015 std::string("unimplemented AMDGPU builtin call: ") +
1016 getContext().BuiltinInfo.getName(builtinId));
1017 return mlir::Value{};
1018 }
1019 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
1020 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
1021 cgm.errorNYI(expr->getSourceRange(),
1022 std::string("unimplemented AMDGPU builtin call: ") +
1023 getContext().BuiltinInfo.getName(builtinId));
1024 return mlir::Value{};
1025 }
1026 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
1027 cgm.errorNYI(expr->getSourceRange(),
1028 std::string("unimplemented AMDGPU builtin call: ") +
1029 getContext().BuiltinInfo.getName(builtinId));
1030 return mlir::Value{};
1031 }
1032 case Builtin::BIlogbf:
1033 case Builtin::BI__builtin_logbf:
1034 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
1035 case Builtin::BIlogb:
1036 case Builtin::BI__builtin_logb:
1037 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble());
1038 case Builtin::BIscalbnf:
1039 case Builtin::BI__builtin_scalbnf:
1040 case Builtin::BIscalbn:
1041 case Builtin::BI__builtin_scalbn: {
1042 return emitBinaryExpMaybeConstrainedFPBuiltin(
1043 *this, expr, "ldexp", "experimental.constrained.ldexp");
1044 }
1045 default:
1046 return std::nullopt;
1047 }
1048}
emitAMDGPUDispatchPtr
static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, const CallExpr *e=nullptr)
Definition CIRGenBuiltinAMDGPU.cpp:27
emitLogbBuiltin
static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, const llvm::fltSemantics &fSem)
Definition CIRGenBuiltinAMDGPU.cpp:70
emitAMDGCNImageOverloadedReturnType
static mlir::Value emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, bool isImageStore)
Definition CIRGenBuiltinAMDGPU.cpp:122
emitBinaryExpMaybeConstrainedFPBuiltin
static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, llvm::StringRef constrainedIntrinsicName)
Definition CIRGenBuiltinAMDGPU.cpp:50
TargetBuiltins.h
Enumerates target-specific builtins in their own namespaces within namespace clang.
cir::CIRBaseBuilderTy::getConstant
cir::ConstantOp getConstant(mlir::Location loc, mlir::TypedAttr attr)
Definition CIRBaseBuilder.h:76
cir::CIRBaseBuilderTy::createAdd
mlir::Value createAdd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, OverflowBehavior ob=OverflowBehavior::None)
Definition CIRBaseBuilder.h:692
cir::CIRBaseBuilderTy::createCompare
cir::CmpOp createCompare(mlir::Location loc, cir::CmpOpKind kind, mlir::Value lhs, mlir::Value rhs)
Definition CIRBaseBuilder.h:771
cir::CIRBaseBuilderTy::createSelect
mlir::Value createSelect(mlir::Location loc, mlir::Value condition, mlir::Value trueValue, mlir::Value falseValue)
Definition CIRBaseBuilder.h:639
cir::CIRBaseBuilderTy::createAddrSpaceCast
mlir::Value createAddrSpaceCast(mlir::Location loc, mlir::Value src, mlir::Type newTy)
Definition CIRBaseBuilder.h:588
cir::RecordType
C++ view class that accepts both !cir.struct and !cir.union types.
Definition CIRTypes.h:93
clang::CIRGen::Address::getElementType
mlir::Type getElementType() const
Definition Address.h:125
clang::CIRGen::CIRGenBuilderTy::emitIntrinsicCallOp
mlir::Value emitIntrinsicCallOp(mlir::Location loc, const llvm::StringRef str, const mlir::Type &resTy, Operands &&...op)
Definition CIRGenBuilder.h:855
clang::CIRGen::CIRGenBuilderTy::getAnonRecordTy
cir::StructType getAnonRecordTy(llvm::ArrayRef< mlir::Type > members, bool packed=false, bool padded=false)
Get a CIR anonymous struct type.
Definition CIRGenBuilder.h:436
clang::CIRGen::CIRGenBuilderTy::getIsFPConstrained
bool getIsFPConstrained() const
Query for the use of constrained floating point math.
Definition CIRGenBuilder.h:134
clang::CIRGen::CIRGenFunction::convertType
mlir::Type convertType(clang::QualType t)
Definition CIRGenFunction.cpp:104
clang::CIRGen::CIRGenFunction::emitPointerWithAlignment
Address emitPointerWithAlignment(const clang::Expr *expr, LValueBaseInfo *baseInfo=nullptr)
Given an expression with a pointer type, emit the value and compute our best estimate of the alignmen...
Definition CIRGenExpr.cpp:117
clang::CIRGen::CIRGenFunction::getLoc
mlir::Location getLoc(clang::SourceLocation srcLoc)
Helpers to convert Clang's SourceLocation to a MLIR Location.
Definition CIRGenFunction.cpp:108
clang::CIRGen::CIRGenFunction::emitBuiltinWithOneOverloadedType
RValue emitBuiltinWithOneOverloadedType(const CallExpr *e, llvm::StringRef intrinName, mlir::Type resultType={})
Emit a simple LLVM intrinsic that takes N scalar arguments.
Definition CIRGenFunction.h:1722
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:142
clang::CIRGen::CIRGenFunction::emitScalarExpr
mlir::Value emitScalarExpr(const clang::Expr *e, bool ignoreResultAssign=false)
Emit the computation of the specified expression of scalar type.
Definition CIRGenExprScalar.cpp:1768
clang::CIRGen::CIRGenFunction::getBuilder
CIRGenBuilderTy & getBuilder()
Definition CIRGenFunction.h:222
clang::CIRGen::CIRGenFunction::getContext
clang::ASTContext & getContext() const
Definition CIRGenFunction.h:220
clang::CIRGen::CIRGenModule::errorNYI
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
Definition CIRGenModule.cpp:3844
clang::CIRGen::RValue::getValue
mlir::Value getValue() const
Return the value of this scalar value.
Definition CIRGenValue.h:57
clang::CallExpr
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition Expr.h:2949
clang::CallExpr::getArg
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition Expr.h:3153
clang::CallExpr::getNumArgs
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
Definition Expr.h:3140
clang::Expr::getExprLoc
SourceLocation getExprLoc() const LLVM_READONLY
getExprLoc - Return the preferred location for the arrow when diagnosing a problem with a generic exp...
Definition Expr.cpp:283
clang::Expr::getType
QualType getType() const
Definition Expr.h:144
clang::Stmt::getSourceRange
SourceRange getSourceRange() const LLVM_READONLY
SourceLocation tokens are not useful in isolation - they are low level value objects created/interpre...
Definition Stmt.cpp:343
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
exp
#define exp(__x)
Definition tgmath.h:431
fabs
#define fabs(__x)
Definition tgmath.h:549
Generated on for clang by doxygen 1.14.0