clang 24.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 mlir::Type frExpResMembers[] = {srcTy, int32Ty};
80 cir::RecordType frExpResTy = builder.getAnonRecordTy(
81 frExpResMembers, /*packed=*/false,
82 cir::RecordType::getAllDataKinds(frExpResMembers));
83
84 mlir::Value frExpResult = builder.emitIntrinsicCallOp(
85 loc, "frexp", frExpResTy, mlir::ValueRange{src0});
86
87 mlir::Value exp =
88 cir::ExtractMemberOp::create(builder, loc, int32Ty, frExpResult, 1);
89
90 mlir::Value negativeOne =
91 builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1));
92 mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne);
93
94 mlir::Value siToFp = cir::CastOp::create(
95 builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1);
96
97 mlir::Value fabs = cir::FAbsOp::create(builder, loc, srcTy, src0);
98
99 llvm::APFloat infVal = llvm::APFloat::getInf(fSem);
100 mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal));
101
102 mlir::Value fabsNegInf =
103 builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf);
104
105 mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs);
106
107 llvm::APFloat zeroValue = llvm::APFloat::getZero(fSem);
108 mlir::Value zero =
109 builder.getConstant(loc, cir::FPAttr::get(srcTy, zeroValue));
110
111 mlir::Value srcEqZero =
112 builder.createCompare(loc, cir::CmpOpKind::eq, src0, zero);
113
114 llvm::APFloat negInfVal = llvm::APFloat::getInf(fSem, true);
115 mlir::Value negInf =
116 builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal));
117
118 mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel);
119
120 return res;
121}
122
123static mlir::Value
124emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e,
125 llvm::StringRef intrinsicName,
126 bool isImageStore) {
127 auto &builder = cgf.getBuilder();
128
129 llvm::SmallVector<mlir::Value, 10> args;
130 for (unsigned i = 0, n = e->getNumArgs(); i < n; ++i)
131 args.push_back(cgf.emitScalarExpr(e->getArg(i)));
132
133 mlir::Type retTy = isImageStore ? cir::VoidType::get(builder.getContext())
134 : cgf.convertType(e->getType());
135
136 auto callOp = cir::LLVMIntrinsicCallOp::create(
137 builder, cgf.getLoc(e->getExprLoc()),
138 builder.getStringAttr(intrinsicName), retTy, args);
139
140 return callOp.getResult();
141}
142
143std::optional<mlir::Value>
144CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
145 const CallExpr *expr) {
146 switch (builtinId) {
147 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
148 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
149 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
150 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
151 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
152 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
153 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
154 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
155 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
156 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
157 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
158 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
159 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
160 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
161 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
162 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
163 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
164 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
165 cgm.errorNYI(expr->getSourceRange(),
166 std::string("unimplemented AMDGPU builtin call: ") +
167 getContext().BuiltinInfo.getName(builtinId));
168 return mlir::Value{};
169 }
170 case AMDGPU::BI__builtin_amdgcn_div_scale:
171 case AMDGPU::BI__builtin_amdgcn_div_scalef: {
172 Address flagOutPtr = emitPointerWithAlignment(expr->getArg(3));
173 llvm::StringRef intrinsicName = "amdgcn.div.scale";
174 mlir::Value x = emitScalarExpr(expr->getArg(0));
175 mlir::Value y = emitScalarExpr(expr->getArg(1));
176 mlir::Value z = emitScalarExpr(expr->getArg(2));
177
178 auto i1Ty = builder.getUIntNTy(1);
179 mlir::Type resMembers[] = {x.getType(), i1Ty};
180 cir::RecordType resTy =
181 builder.getAnonRecordTy(resMembers, /*packed=*/false,
182 cir::RecordType::getAllDataKinds(resMembers));
183
184 mlir::Value structResult =
185 cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()),
186 builder.getStringAttr(intrinsicName),
187 resTy, {x, y, z})
188 .getResult();
189
190 mlir::Value result = cir::ExtractMemberOp::create(
191 builder, getLoc(expr->getExprLoc()), x.getType(), structResult, 0);
192 mlir::Value flag = cir::ExtractMemberOp::create(
193 builder, getLoc(expr->getExprLoc()), i1Ty, structResult, 1);
194
195 mlir::Type flagType = flagOutPtr.getElementType();
196 mlir::Value flagToStore =
197 cir::CastOp::create(builder, getLoc(expr->getExprLoc()), flagType,
198 cir::CastKind::int_to_bool, flag);
199 builder.createStore(getLoc(expr->getExprLoc()), flagToStore, flagOutPtr);
200 return result;
201 }
202 case AMDGPU::BI__builtin_amdgcn_div_fmas:
203 case AMDGPU::BI__builtin_amdgcn_div_fmasf:
204 return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas")
205 .getValue();
206 case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
207 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.ds.swizzle")
208 .getValue();
209 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
210 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
211 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
212 cgm.errorNYI(expr->getSourceRange(),
213 std::string("unimplemented AMDGPU builtin call: ") +
214 getContext().BuiltinInfo.getName(builtinId));
215 return mlir::Value{};
216 }
217 case AMDGPU::BI__builtin_amdgcn_permlane16:
218 case AMDGPU::BI__builtin_amdgcn_permlanex16: {
219 llvm::StringRef intrinsicName =
220 builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
221 ? "amdgcn.permlane16"
222 : "amdgcn.permlanex16";
223 return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName).getValue();
224 }
225 case AMDGPU::BI__builtin_amdgcn_permlane64:
226 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64")
227 .getValue();
228 case AMDGPU::BI__builtin_amdgcn_readlane:
229 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
230 .getValue();
231 case AMDGPU::BI__builtin_amdgcn_readfirstlane:
232 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
233 .getValue();
234 case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
235 cgm.errorNYI(expr->getSourceRange(),
236 std::string("unimplemented AMDGPU builtin call: ") +
237 getContext().BuiltinInfo.getName(builtinId));
238 return mlir::Value{};
239 }
240 case AMDGPU::BI__builtin_amdgcn_div_fixup:
241 case AMDGPU::BI__builtin_amdgcn_div_fixupf:
242 case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
243 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
244 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
245 mlir::Value src2 = emitScalarExpr(expr->getArg(2));
246 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
247 "amdgcn.div.fixup", src0.getType(),
248 mlir::ValueRange{src0, src1, src2});
249 }
250 case AMDGPU::BI__builtin_amdgcn_trig_preop:
251 case AMDGPU::BI__builtin_amdgcn_trig_preopf:
252 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.trig.preop")
253 .getValue();
254 case AMDGPU::BI__builtin_amdgcn_rcp:
255 case AMDGPU::BI__builtin_amdgcn_rcpf:
256 case AMDGPU::BI__builtin_amdgcn_rcph:
257 case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
258 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue();
259 }
260 case AMDGPU::BI__builtin_amdgcn_sqrt:
261 case AMDGPU::BI__builtin_amdgcn_sqrtf:
262 case AMDGPU::BI__builtin_amdgcn_sqrth:
263 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
264 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt").getValue();
265 }
266 case AMDGPU::BI__builtin_amdgcn_rsq:
267 case AMDGPU::BI__builtin_amdgcn_rsqf:
268 case AMDGPU::BI__builtin_amdgcn_rsqh:
269 case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
270 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue();
271 }
272 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
273 case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
274 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp")
275 .getValue();
276 }
277 case AMDGPU::BI__builtin_amdgcn_sinf:
278 case AMDGPU::BI__builtin_amdgcn_sinh:
279 case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
280 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue();
281 }
282 case AMDGPU::BI__builtin_amdgcn_cosf:
283 case AMDGPU::BI__builtin_amdgcn_cosh:
284 case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
285 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue();
286 }
287 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
288 return emitAMDGPUDispatchPtr(*this, expr);
289 case AMDGPU::BI__builtin_amdgcn_logf:
290 case AMDGPU::BI__builtin_amdgcn_log_bf16: {
291 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.log").getValue();
292 }
293 case AMDGPU::BI__builtin_amdgcn_exp2f:
294 case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
295 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.exp2").getValue();
296 }
297 case AMDGPU::BI__builtin_amdgcn_log_clampf: {
298 cgm.errorNYI(expr->getSourceRange(),
299 std::string("unimplemented AMDGPU builtin call: ") +
300 getContext().BuiltinInfo.getName(builtinId));
301 return mlir::Value{};
302 }
303 case AMDGPU::BI__builtin_amdgcn_ldexp:
304 case AMDGPU::BI__builtin_amdgcn_ldexpf:
305 case AMDGPU::BI__builtin_amdgcn_ldexph: {
306 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
307 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
308 // For ldexph, the raw instruction has different behavior for out-of-bounds
309 // exponent values (implicit truncation instead of saturate to
310 // short_min/short_max), so truncate the exponent to i16 first.
311 mlir::Value exp =
312 builtinId == AMDGPU::BI__builtin_amdgcn_ldexph
313 ? cir::CastOp::create(builder, getLoc(expr->getExprLoc()),
314 builder.getSInt16Ty(),
315 cir::CastKind::integral, src1)
316 : src1;
317 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
318 src0.getType(),
319 mlir::ValueRange{src0, exp});
320 }
321 case AMDGPU::BI__builtin_amdgcn_frexp_mant:
322 case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
323 case AMDGPU::BI__builtin_amdgcn_frexp_manth: {
324 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant")
325 .getValue();
326 }
327 case AMDGPU::BI__builtin_amdgcn_frexp_exp:
328 case AMDGPU::BI__builtin_amdgcn_frexp_expf:
329 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
330 cgm.errorNYI(expr->getSourceRange(),
331 std::string("unimplemented AMDGPU builtin call: ") +
332 getContext().BuiltinInfo.getName(builtinId));
333 return mlir::Value{};
334 }
335 case AMDGPU::BI__builtin_amdgcn_fract:
336 case AMDGPU::BI__builtin_amdgcn_fractf:
337 case AMDGPU::BI__builtin_amdgcn_fracth: {
338 cgm.errorNYI(expr->getSourceRange(),
339 std::string("unimplemented AMDGPU builtin call: ") +
340 getContext().BuiltinInfo.getName(builtinId));
341 return mlir::Value{};
342 }
343 case AMDGPU::BI__builtin_amdgcn_lerp: {
344 cgm.errorNYI(expr->getSourceRange(),
345 std::string("unimplemented AMDGPU builtin call: ") +
346 getContext().BuiltinInfo.getName(builtinId));
347 return mlir::Value{};
348 }
349 case AMDGPU::BI__builtin_amdgcn_ubfe: {
350 cgm.errorNYI(expr->getSourceRange(),
351 std::string("unimplemented AMDGPU builtin call: ") +
352 getContext().BuiltinInfo.getName(builtinId));
353 return mlir::Value{};
354 }
355 case AMDGPU::BI__builtin_amdgcn_sbfe: {
356 cgm.errorNYI(expr->getSourceRange(),
357 std::string("unimplemented AMDGPU builtin call: ") +
358 getContext().BuiltinInfo.getName(builtinId));
359 return mlir::Value{};
360 }
361 case AMDGPU::BI__builtin_amdgcn_ballot_w32:
362 case AMDGPU::BI__builtin_amdgcn_ballot_w64: {
363 cgm.errorNYI(expr->getSourceRange(),
364 std::string("unimplemented AMDGPU builtin call: ") +
365 getContext().BuiltinInfo.getName(builtinId));
366 return mlir::Value{};
367 }
368 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
369 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
370 cgm.errorNYI(expr->getSourceRange(),
371 std::string("unimplemented AMDGPU builtin call: ") +
372 getContext().BuiltinInfo.getName(builtinId));
373 return mlir::Value{};
374 }
375 case AMDGPU::BI__builtin_amdgcn_tanhf:
376 case AMDGPU::BI__builtin_amdgcn_tanhh:
377 case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
378 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue();
379 }
380 case AMDGPU::BI__builtin_amdgcn_uicmp:
381 case AMDGPU::BI__builtin_amdgcn_uicmpl:
382 case AMDGPU::BI__builtin_amdgcn_sicmp:
383 case AMDGPU::BI__builtin_amdgcn_sicmpl: {
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_fcmp:
390 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
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_class:
397 case AMDGPU::BI__builtin_amdgcn_classf:
398 case AMDGPU::BI__builtin_amdgcn_classh: {
399 cgm.errorNYI(expr->getSourceRange(),
400 std::string("unimplemented AMDGPU builtin call: ") +
401 getContext().BuiltinInfo.getName(builtinId));
402 return mlir::Value{};
403 }
404 case AMDGPU::BI__builtin_amdgcn_fmed3f:
405 case AMDGPU::BI__builtin_amdgcn_fmed3h: {
406 cgm.errorNYI(expr->getSourceRange(),
407 std::string("unimplemented AMDGPU builtin call: ") +
408 getContext().BuiltinInfo.getName(builtinId));
409 return mlir::Value{};
410 }
411 case AMDGPU::BI__builtin_amdgcn_ds_append:
412 case AMDGPU::BI__builtin_amdgcn_ds_consume: {
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_global_load_tr_b64_i32:
419 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
420 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
421 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
422 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
423 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
424 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
425 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
426 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
427 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
428 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
429 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
430 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
431 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: {
432 cgm.errorNYI(expr->getSourceRange(),
433 std::string("unimplemented AMDGPU builtin call: ") +
434 getContext().BuiltinInfo.getName(builtinId));
435 return mlir::Value{};
436 }
437 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
438 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
439 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
440 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
441 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
442 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: {
443 cgm.errorNYI(expr->getSourceRange(),
444 std::string("unimplemented AMDGPU builtin call: ") +
445 getContext().BuiltinInfo.getName(builtinId));
446 return mlir::Value{};
447 }
448 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
449 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
450 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
451 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
452 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
453 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {
454 cgm.errorNYI(expr->getSourceRange(),
455 std::string("unimplemented AMDGPU builtin call: ") +
456 getContext().BuiltinInfo.getName(builtinId));
457 return mlir::Value{};
458 }
459 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
460 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
461 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
462 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
463 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
464 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
465 cgm.errorNYI(expr->getSourceRange(),
466 std::string("unimplemented AMDGPU builtin call: ") +
467 getContext().BuiltinInfo.getName(builtinId));
468 return mlir::Value{};
469 }
470 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
471 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
472 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
473 cgm.errorNYI(expr->getSourceRange(),
474 std::string("unimplemented AMDGPU builtin call: ") +
475 getContext().BuiltinInfo.getName(builtinId));
476 return mlir::Value{};
477 }
478 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
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_cooperative_atomic_load_32x4B:
485 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
486 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
487 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
488 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
489 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
490 cgm.errorNYI(expr->getSourceRange(),
491 std::string("unimplemented AMDGPU builtin call: ") +
492 getContext().BuiltinInfo.getName(builtinId));
493 return mlir::Value{};
494 }
495 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
496 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
497 cgm.errorNYI(expr->getSourceRange(),
498 std::string("unimplemented AMDGPU builtin call: ") +
499 getContext().BuiltinInfo.getName(builtinId));
500 return mlir::Value{};
501 }
502 case AMDGPU::BI__builtin_amdgcn_read_exec:
503 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
504 case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
505 cgm.errorNYI(expr->getSourceRange(),
506 std::string("unimplemented AMDGPU builtin call: ") +
507 getContext().BuiltinInfo.getName(builtinId));
508 return mlir::Value{};
509 }
510 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
511 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
512 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
513 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
514 cgm.errorNYI(expr->getSourceRange(),
515 std::string("unimplemented AMDGPU builtin call: ") +
516 getContext().BuiltinInfo.getName(builtinId));
517 return mlir::Value{};
518 }
519 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:
520 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
521 cgm.errorNYI(expr->getSourceRange(),
522 std::string("unimplemented AMDGPU builtin call: ") +
523 getContext().BuiltinInfo.getName(builtinId));
524 return mlir::Value{};
525 }
526 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:
527 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
528 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
529 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
530 cgm.errorNYI(expr->getSourceRange(),
531 std::string("unimplemented AMDGPU builtin call: ") +
532 getContext().BuiltinInfo.getName(builtinId));
533 return mlir::Value{};
534 }
535 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
536 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
537 return emitAMDGCNImageOverloadedReturnType(*this, expr,
538 "amdgcn.image.load.1d", false);
539 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
540 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
541 return emitAMDGCNImageOverloadedReturnType(
542 *this, expr, "amdgcn.image.load.1darray", false);
543 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
544 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
545 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
546 return emitAMDGCNImageOverloadedReturnType(*this, expr,
547 "amdgcn.image.load.2d", false);
548 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
549 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
550 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
551 return emitAMDGCNImageOverloadedReturnType(
552 *this, expr, "amdgcn.image.load.2darray", false);
553 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
554 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
555 return emitAMDGCNImageOverloadedReturnType(*this, expr,
556 "amdgcn.image.load.3d", false);
557 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
558 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
559 return emitAMDGCNImageOverloadedReturnType(*this, expr,
560 "amdgcn.image.load.cube", false);
561 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
562 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
563 return emitAMDGCNImageOverloadedReturnType(
564 *this, expr, "amdgcn.image.load.mip.1d", false);
565 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
566 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
567 return emitAMDGCNImageOverloadedReturnType(
568 *this, expr, "amdgcn.image.load.mip.1darray", false);
569 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
570 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
571 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
572 return emitAMDGCNImageOverloadedReturnType(
573 *this, expr, "amdgcn.image.load.mip.2d", false);
574 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
575 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
576 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
577 return emitAMDGCNImageOverloadedReturnType(
578 *this, expr, "amdgcn.image.load.mip.2darray", false);
579 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
580 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
581 return emitAMDGCNImageOverloadedReturnType(
582 *this, expr, "amdgcn.image.load.mip.3d", false);
583 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
584 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
585 return emitAMDGCNImageOverloadedReturnType(
586 *this, expr, "amdgcn.image.load.mip.cube", false);
587 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
588 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
589 return emitAMDGCNImageOverloadedReturnType(*this, expr,
590 "amdgcn.image.store.1d", true);
591 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
592 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
593 return emitAMDGCNImageOverloadedReturnType(
594 *this, expr, "amdgcn.image.store.1darray", true);
595 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
596 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
597 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
598 return emitAMDGCNImageOverloadedReturnType(*this, expr,
599 "amdgcn.image.store.2d", true);
600 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
601 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
602 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
603 return emitAMDGCNImageOverloadedReturnType(
604 *this, expr, "amdgcn.image.store.2darray", true);
605 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
606 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
607 return emitAMDGCNImageOverloadedReturnType(*this, expr,
608 "amdgcn.image.store.3d", true);
609 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
610 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
611 return emitAMDGCNImageOverloadedReturnType(*this, expr,
612 "amdgcn.image.store.cube", true);
613 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
614 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
615 return emitAMDGCNImageOverloadedReturnType(
616 *this, expr, "amdgcn.image.store.mip.1d", true);
617 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
618 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
619 return emitAMDGCNImageOverloadedReturnType(
620 *this, expr, "amdgcn.image.store.mip.1darray", true);
621 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
622 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
623 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
624 return emitAMDGCNImageOverloadedReturnType(
625 *this, expr, "amdgcn.image.store.mip.2d", true);
626 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
627 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
628 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
629 return emitAMDGCNImageOverloadedReturnType(
630 *this, expr, "amdgcn.image.store.mip.2darray", true);
631 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
632 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
633 return emitAMDGCNImageOverloadedReturnType(
634 *this, expr, "amdgcn.image.store.mip.3d", true);
635 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
636 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
637 return emitAMDGCNImageOverloadedReturnType(
638 *this, expr, "amdgcn.image.store.mip.cube", true);
639 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
640 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
641 return emitAMDGCNImageOverloadedReturnType(*this, expr,
642 "amdgcn.image.sample.1d", false);
643 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
644 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
645 return emitAMDGCNImageOverloadedReturnType(
646 *this, expr, "amdgcn.image.sample.1darray", false);
647 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
648 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
649 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
650 return emitAMDGCNImageOverloadedReturnType(*this, expr,
651 "amdgcn.image.sample.2d", false);
652 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
653 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
654 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
655 return emitAMDGCNImageOverloadedReturnType(
656 *this, expr, "amdgcn.image.sample.2darray", false);
657 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
658 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
659 return emitAMDGCNImageOverloadedReturnType(*this, expr,
660 "amdgcn.image.sample.3d", false);
661 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
662 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
663 return emitAMDGCNImageOverloadedReturnType(
664 *this, expr, "amdgcn.image.sample.cube", false);
665 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
666 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
667 return emitAMDGCNImageOverloadedReturnType(
668 *this, expr, "amdgcn.image.sample.lz.1d", false);
669 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
670 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
671 return emitAMDGCNImageOverloadedReturnType(
672 *this, expr, "amdgcn.image.sample.l.1d", false);
673 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
674 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
675 return emitAMDGCNImageOverloadedReturnType(
676 *this, expr, "amdgcn.image.sample.d.1d", false);
677 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
678 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
679 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
680 return emitAMDGCNImageOverloadedReturnType(
681 *this, expr, "amdgcn.image.sample.lz.2d", false);
682 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
683 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
684 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
685 return emitAMDGCNImageOverloadedReturnType(
686 *this, expr, "amdgcn.image.sample.l.2d", false);
687 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
688 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
689 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
690 return emitAMDGCNImageOverloadedReturnType(
691 *this, expr, "amdgcn.image.sample.d.2d", false);
692 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
693 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
694 return emitAMDGCNImageOverloadedReturnType(
695 *this, expr, "amdgcn.image.sample.lz.3d", false);
696 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
697 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
698 return emitAMDGCNImageOverloadedReturnType(
699 *this, expr, "amdgcn.image.sample.l.3d", false);
700 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
701 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
702 return emitAMDGCNImageOverloadedReturnType(
703 *this, expr, "amdgcn.image.sample.d.3d", false);
704 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
705 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
706 return emitAMDGCNImageOverloadedReturnType(
707 *this, expr, "amdgcn.image.sample.lz.cube", false);
708 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
709 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
710 return emitAMDGCNImageOverloadedReturnType(
711 *this, expr, "amdgcn.image.sample.l.cube", false);
712 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
713 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
714 return emitAMDGCNImageOverloadedReturnType(
715 *this, expr, "amdgcn.image.sample.lz.1darray", false);
716 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
717 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
718 return emitAMDGCNImageOverloadedReturnType(
719 *this, expr, "amdgcn.image.sample.l.1darray", false);
720 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
721 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
722 return emitAMDGCNImageOverloadedReturnType(
723 *this, expr, "amdgcn.image.sample.d.1darray", false);
724 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
725 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
726 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
727 return emitAMDGCNImageOverloadedReturnType(
728 *this, expr, "amdgcn.image.sample.lz.2darray", false);
729 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
730 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
731 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
732 return emitAMDGCNImageOverloadedReturnType(
733 *this, expr, "amdgcn.image.sample.l.2darray", false);
734 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
735 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
736 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
737 return emitAMDGCNImageOverloadedReturnType(
738 *this, expr, "amdgcn.image.sample.d.2darray", false);
739 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
740 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32:
741 return emitAMDGCNImageOverloadedReturnType(
742 *this, expr, "amdgcn.image.gather4.lz.2d", false);
743 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
744 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
745 cgm.errorNYI(expr->getSourceRange(),
746 std::string("unimplemented AMDGPU builtin call: ") +
747 getContext().BuiltinInfo.getName(builtinId));
748 return mlir::Value{};
749 }
750 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
751 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
752 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
753 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
754 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
755 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
756 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
757 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
758 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
759 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
760 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
761 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
762 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
763 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
764 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
765 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
766 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
767 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
768 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
769 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:
770 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:
771 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:
772 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:
773 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:
774 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:
775 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:
776 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:
777 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:
778 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:
779 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:
780 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:
781 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:
782 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:
783 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:
784 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:
785 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:
786 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:
787 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: {
788 cgm.errorNYI(expr->getSourceRange(),
789 std::string("unimplemented AMDGPU builtin call: ") +
790 getContext().BuiltinInfo.getName(builtinId));
791 return mlir::Value{};
792 }
793 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:
794 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:
795 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:
796 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:
797 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:
798 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:
799 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:
800 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:
801 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:
802 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:
803 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:
804 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:
805 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:
806 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:
807 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:
808 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:
809 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:
810 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:
811 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:
812 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:
813 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:
814 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {
815 cgm.errorNYI(expr->getSourceRange(),
816 std::string("unimplemented AMDGPU builtin call: ") +
817 getContext().BuiltinInfo.getName(builtinId));
818 return mlir::Value{};
819 }
820 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
821 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
822 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
823 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
824 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
825 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
826 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:
827 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:
828 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:
829 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:
830 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:
831 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:
832 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:
833 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:
834 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:
835 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:
836 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:
837 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:
838 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:
839 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:
840 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:
841 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:
842 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
843 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
844 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
845 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
846 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
847 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
848 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: {
849 cgm.errorNYI(expr->getSourceRange(),
850 std::string("unimplemented AMDGPU builtin call: ") +
851 getContext().BuiltinInfo.getName(builtinId));
852 return mlir::Value{};
853 }
854 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
855 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
856 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
857 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:
858 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:
859 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:
860 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:
861 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:
862 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:
863 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:
864 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:
865 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:
866 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:
867 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
868 cgm.errorNYI(expr->getSourceRange(),
869 std::string("unimplemented AMDGPU builtin call: ") +
870 getContext().BuiltinInfo.getName(builtinId));
871 return mlir::Value{};
872 }
873 // amdgcn workgroup size
874 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
875 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
876 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
877 cgm.errorNYI(expr->getSourceRange(),
878 std::string("unimplemented AMDGPU builtin call: ") +
879 getContext().BuiltinInfo.getName(builtinId));
880 return mlir::Value{};
881 }
882 case AMDGPU::BI__builtin_amdgcn_grid_size_x:
883 case AMDGPU::BI__builtin_amdgcn_grid_size_y:
884 case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
885 cgm.errorNYI(expr->getSourceRange(),
886 std::string("unimplemented AMDGPU builtin call: ") +
887 getContext().BuiltinInfo.getName(builtinId));
888 return mlir::Value{};
889 }
890 case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
891 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
892 cgm.errorNYI(expr->getSourceRange(),
893 std::string("unimplemented AMDGPU builtin call: ") +
894 getContext().BuiltinInfo.getName(builtinId));
895 return mlir::Value{};
896 }
897 case AMDGPU::BI__builtin_amdgcn_alignbit: {
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_amdgcn_fence: {
904 cgm.errorNYI(expr->getSourceRange(),
905 std::string("unimplemented AMDGPU builtin call: ") +
906 getContext().BuiltinInfo.getName(builtinId));
907 return mlir::Value{};
908 }
909 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
910 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
911 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
912 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
913 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
914 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
915 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
916 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:
917 case AMDGPU::BI__builtin_amdgcn_ds_faddf:
918 case AMDGPU::BI__builtin_amdgcn_ds_fminf:
919 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
920 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
921 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
922 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
923 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
924 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
925 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
926 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
927 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
928 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
929 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
930 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
931 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
932 cgm.errorNYI(expr->getSourceRange(),
933 std::string("unimplemented AMDGPU builtin call: ") +
934 getContext().BuiltinInfo.getName(builtinId));
935 return mlir::Value{};
936 }
937 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
938 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
939 cgm.errorNYI(expr->getSourceRange(),
940 std::string("unimplemented AMDGPU builtin call: ") +
941 getContext().BuiltinInfo.getName(builtinId));
942 return mlir::Value{};
943 }
944 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:
945 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {
946 cgm.errorNYI(expr->getSourceRange(),
947 std::string("unimplemented AMDGPU builtin call: ") +
948 getContext().BuiltinInfo.getName(builtinId));
949 return mlir::Value{};
950 }
951 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
952 case AMDGPU::BI__builtin_amdgcn_bitop3_b16: {
953 cgm.errorNYI(expr->getSourceRange(),
954 std::string("unimplemented AMDGPU builtin call: ") +
955 getContext().BuiltinInfo.getName(builtinId));
956 return mlir::Value{};
957 }
958 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
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_raw_buffer_store_b8:
965 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
966 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
967 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
968 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
969 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: {
970 cgm.errorNYI(expr->getSourceRange(),
971 std::string("unimplemented AMDGPU builtin call: ") +
972 getContext().BuiltinInfo.getName(builtinId));
973 return mlir::Value{};
974 }
975 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
976 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
977 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
978 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
979 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
980 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
981 cgm.errorNYI(expr->getSourceRange(),
982 std::string("unimplemented AMDGPU builtin call: ") +
983 getContext().BuiltinInfo.getName(builtinId));
984 return mlir::Value{};
985 }
986 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
987 cgm.errorNYI(expr->getSourceRange(),
988 std::string("unimplemented AMDGPU builtin call: ") +
989 getContext().BuiltinInfo.getName(builtinId));
990 return mlir::Value{};
991 }
992 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
993 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
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_fmin_f32:
1000 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
1001 cgm.errorNYI(expr->getSourceRange(),
1002 std::string("unimplemented AMDGPU builtin call: ") +
1003 getContext().BuiltinInfo.getName(builtinId));
1004 return mlir::Value{};
1005 }
1006 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
1007 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
1008 cgm.errorNYI(expr->getSourceRange(),
1009 std::string("unimplemented AMDGPU builtin call: ") +
1010 getContext().BuiltinInfo.getName(builtinId));
1011 return mlir::Value{};
1012 }
1013 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
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_s_prefetch_inst: {
1020 cgm.errorNYI(expr->getSourceRange(),
1021 std::string("unimplemented AMDGPU builtin call: ") +
1022 getContext().BuiltinInfo.getName(builtinId));
1023 return mlir::Value{};
1024 }
1025 case Builtin::BIlogbf:
1026 case Builtin::BI__builtin_logbf:
1027 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
1028 case Builtin::BIlogb:
1029 case Builtin::BI__builtin_logb:
1030 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble());
1031 case Builtin::BIscalbnf:
1032 case Builtin::BI__builtin_scalbnf:
1033 case Builtin::BIscalbn:
1034 case Builtin::BI__builtin_scalbn: {
1035 return emitBinaryExpMaybeConstrainedFPBuiltin(
1036 *this, expr, "ldexp", "experimental.constrained.ldexp");
1037 }
1038 default:
1039 return std::nullopt;
1040 }
1041}
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:124
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:83
cir::CIRBaseBuilderTy::createAdd
mlir::Value createAdd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, OverflowBehavior ob=OverflowBehavior::None)
Definition CIRBaseBuilder.h:815
cir::CIRBaseBuilderTy::createCompare
cir::CmpOp createCompare(mlir::Location loc, cir::CmpOpKind kind, mlir::Value lhs, mlir::Value rhs)
Definition CIRBaseBuilder.h:890
cir::CIRBaseBuilderTy::createSelect
mlir::Value createSelect(mlir::Location loc, mlir::Value condition, mlir::Value trueValue, mlir::Value falseValue)
Definition CIRBaseBuilder.h:762
cir::CIRBaseBuilderTy::getIsFPConstrained
bool getIsFPConstrained() const
Query for the use of constrained floating point math.
Definition CIRBaseBuilder.h:224
cir::CIRBaseBuilderTy::createAddrSpaceCast
mlir::Value createAddrSpaceCast(mlir::Location loc, mlir::Value src, mlir::Type newTy)
Definition CIRBaseBuilder.h:711
cir::RecordType
C++ view class that accepts both !cir.struct and !cir.union types.
Definition CIRTypes.h:143
cir::RecordType::getAllDataKinds
static llvm::SmallVector< RecordMemberKind > getAllDataKinds(llvm::ArrayRef< mlir::Type > members)
One Data kind per member.
Definition CIRTypes.cpp:168
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:818
clang::CIRGen::CIRGenBuilderTy::getAnonRecordTy
cir::StructType getAnonRecordTy(llvm::ArrayRef< mlir::Type > members, bool packed, llvm::ArrayRef< cir::RecordMemberKind > memberKinds)
Get a CIR anonymous struct type.
Definition CIRGenBuilder.h:393
clang::CIRGen::CIRGenFunction::convertType
mlir::Type convertType(clang::QualType t)
Definition CIRGenFunction.cpp:139
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:143
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:1828
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:144
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:1789
clang::CIRGen::CIRGenFunction::getBuilder
CIRGenBuilderTy & getBuilder()
Definition CIRGenFunction.h:263
clang::CIRGen::CIRGenFunction::getContext
clang::ASTContext & getContext() const
Definition CIRGenFunction.h:261
clang::CIRGen::CIRGenModule::errorNYI
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
Definition CIRGenModule.cpp:4110
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:2963
clang::CallExpr::getArg
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition Expr.h:3167
clang::CallExpr::getNumArgs
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
Definition Expr.h:3154
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:145
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
Top level wrappers for InstallAPI frontend operations.
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