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, 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_mov_dpp8:
207 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
208 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
209 cgm.errorNYI(expr->getSourceRange(),
210 std::string("unimplemented AMDGPU builtin call: ") +
211 getContext().BuiltinInfo.getName(builtinId));
212 return mlir::Value{};
213 }
214 case AMDGPU::BI__builtin_amdgcn_permlane16:
215 case AMDGPU::BI__builtin_amdgcn_permlanex16: {
216 llvm::StringRef intrinsicName =
217 builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
218 ? "amdgcn.permlane16"
219 : "amdgcn.permlanex16";
220 return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName).getValue();
221 }
222 case AMDGPU::BI__builtin_amdgcn_permlane64:
223 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64")
224 .getValue();
225 case AMDGPU::BI__builtin_amdgcn_permlane_bcast:
226 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.bcast")
227 .getValue();
228 case AMDGPU::BI__builtin_amdgcn_permlane_up:
229 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.up")
230 .getValue();
231 case AMDGPU::BI__builtin_amdgcn_permlane_down:
232 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.down")
233 .getValue();
234 case AMDGPU::BI__builtin_amdgcn_permlane_xor:
235 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.xor")
236 .getValue();
237 case AMDGPU::BI__builtin_amdgcn_readlane:
238 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
239 .getValue();
240 case AMDGPU::BI__builtin_amdgcn_readfirstlane:
241 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
242 .getValue();
243 case AMDGPU::BI__builtin_amdgcn_wave_shuffle: {
244 cgm.errorNYI(expr->getSourceRange(),
245 std::string("unimplemented AMDGPU builtin call: ") +
246 getContext().BuiltinInfo.getName(builtinId));
247 return mlir::Value{};
248 }
249 case AMDGPU::BI__builtin_amdgcn_div_fixup:
250 case AMDGPU::BI__builtin_amdgcn_div_fixupf:
251 case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
252 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
253 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
254 mlir::Value src2 = emitScalarExpr(expr->getArg(2));
255 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
256 "amdgcn.div.fixup", src0.getType(),
257 mlir::ValueRange{src0, src1, src2});
258 }
259 case AMDGPU::BI__builtin_amdgcn_trig_preop:
260 case AMDGPU::BI__builtin_amdgcn_trig_preopf:
261 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.trig.preop")
262 .getValue();
263 case AMDGPU::BI__builtin_amdgcn_rcp:
264 case AMDGPU::BI__builtin_amdgcn_rcpf:
265 case AMDGPU::BI__builtin_amdgcn_rcph:
266 case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
267 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue();
268 }
269 case AMDGPU::BI__builtin_amdgcn_sqrt:
270 case AMDGPU::BI__builtin_amdgcn_sqrtf:
271 case AMDGPU::BI__builtin_amdgcn_sqrth:
272 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
273 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt").getValue();
274 }
275 case AMDGPU::BI__builtin_amdgcn_rsq:
276 case AMDGPU::BI__builtin_amdgcn_rsqf:
277 case AMDGPU::BI__builtin_amdgcn_rsqh:
278 case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
279 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue();
280 }
281 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
282 case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
283 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp")
284 .getValue();
285 }
286 case AMDGPU::BI__builtin_amdgcn_sinf:
287 case AMDGPU::BI__builtin_amdgcn_sinh:
288 case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
289 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue();
290 }
291 case AMDGPU::BI__builtin_amdgcn_cosf:
292 case AMDGPU::BI__builtin_amdgcn_cosh:
293 case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
294 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue();
295 }
296 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
297 return emitAMDGPUDispatchPtr(*this, expr);
298 case AMDGPU::BI__builtin_amdgcn_logf:
299 case AMDGPU::BI__builtin_amdgcn_log_bf16: {
300 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.log").getValue();
301 }
302 case AMDGPU::BI__builtin_amdgcn_exp2f:
303 case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
304 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.exp2").getValue();
305 }
306 case AMDGPU::BI__builtin_amdgcn_log_clampf: {
307 cgm.errorNYI(expr->getSourceRange(),
308 std::string("unimplemented AMDGPU builtin call: ") +
309 getContext().BuiltinInfo.getName(builtinId));
310 return mlir::Value{};
311 }
312 case AMDGPU::BI__builtin_amdgcn_ldexp:
313 case AMDGPU::BI__builtin_amdgcn_ldexpf:
314 case AMDGPU::BI__builtin_amdgcn_ldexph: {
315 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
316 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
317 // For ldexph, the raw instruction has different behavior for out-of-bounds
318 // exponent values (implicit truncation instead of saturate to
319 // short_min/short_max), so truncate the exponent to i16 first.
320 mlir::Value exp =
321 builtinId == AMDGPU::BI__builtin_amdgcn_ldexph
322 ? cir::CastOp::create(builder, getLoc(expr->getExprLoc()),
323 builder.getSInt16Ty(),
324 cir::CastKind::integral, src1)
325 : src1;
326 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
327 src0.getType(),
328 mlir::ValueRange{src0, exp});
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 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant")
334 .getValue();
335 }
336 case AMDGPU::BI__builtin_amdgcn_frexp_exp:
337 case AMDGPU::BI__builtin_amdgcn_frexp_expf:
338 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
339 cgm.errorNYI(expr->getSourceRange(),
340 std::string("unimplemented AMDGPU builtin call: ") +
341 getContext().BuiltinInfo.getName(builtinId));
342 return mlir::Value{};
343 }
344 case AMDGPU::BI__builtin_amdgcn_fract:
345 case AMDGPU::BI__builtin_amdgcn_fractf:
346 case AMDGPU::BI__builtin_amdgcn_fracth:
347 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.fract").getValue();
348 case AMDGPU::BI__builtin_amdgcn_lerp: {
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_ubfe:
355 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.ubfe").getValue();
356 case AMDGPU::BI__builtin_amdgcn_sbfe:
357 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.sbfe").getValue();
358 case AMDGPU::BI__builtin_amdgcn_ballot_w32:
359 case AMDGPU::BI__builtin_amdgcn_ballot_w64:
360 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot",
361 convertType(expr->getType()))
362 .getValue();
363 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
364 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64:
365 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.inverse.ballot",
366 convertType(expr->getType()))
367 .getValue();
368 case AMDGPU::BI__builtin_amdgcn_tanhf:
369 case AMDGPU::BI__builtin_amdgcn_tanhh:
370 case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
371 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue();
372 }
373 case AMDGPU::BI__builtin_amdgcn_uicmp:
374 case AMDGPU::BI__builtin_amdgcn_uicmpl:
375 case AMDGPU::BI__builtin_amdgcn_sicmp:
376 case AMDGPU::BI__builtin_amdgcn_sicmpl:
377 case AMDGPU::BI__builtin_amdgcn_fcmp:
378 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
379 mlir::Value lhs = emitScalarExpr(expr->getArg(0));
380 mlir::Value rhs = emitScalarExpr(expr->getArg(1));
381
382 uint64_t imm =
383 expr->getArg(2)->EvaluateKnownConstInt(getContext()).getZExtValue();
384
385 cir::CmpOpKind pred;
386 switch (imm) {
387 case 0x1: // FCMP_OEQ
388 case 0x20: // ICMP_EQ
389 pred = cir::CmpOpKind::eq;
390 break;
391 case 0xe: // FCMP_UNE
392 case 0x21: // ICMP_NE
393 pred = cir::CmpOpKind::ne;
394 break;
395 case 0x2: // FCMP_OGT
396 case 0x22: // ICMP_UGT
397 case 0x26: // ICMP_SGT
398 pred = cir::CmpOpKind::gt;
399 break;
400 case 0x3: // FCMP_OGE
401 case 0x23: // ICMP_UGE
402 case 0x27: // ICMP_SGE
403 pred = cir::CmpOpKind::ge;
404 break;
405 case 0x4: // FCMP_OLT
406 case 0x24: // ICMP_ULT
407 case 0x28: // ICMP_SLT
408 pred = cir::CmpOpKind::lt;
409 break;
410 case 0x5: // FCMP_OLE
411 case 0x25: // ICMP_ULE
412 case 0x29: // ICMP_SLE
413 pred = cir::CmpOpKind::le;
414 break;
415 case 0x6: // FCMP_ONE
416 pred = cir::CmpOpKind::one;
417 break;
418 case 0x8: // FCMP_UNO
419 pred = cir::CmpOpKind::uno;
420 break;
421 default:
422 cgm.errorNYI(expr->getSourceRange(),
423 "amdgcn compare with unsupported predicate");
424 return mlir::Value{};
425 }
426
427 mlir::Location loc = getLoc(expr->getExprLoc());
428 mlir::Value cmp = builder.createCompare(loc, pred, lhs, rhs);
429 return builder.emitIntrinsicCallOp(loc, "amdgcn.ballot",
430 convertType(expr->getType()), cmp);
431 }
432 case AMDGPU::BI__builtin_amdgcn_class:
433 case AMDGPU::BI__builtin_amdgcn_classf:
434 case AMDGPU::BI__builtin_amdgcn_classh:
435 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.class",
436 convertType(expr->getType()))
437 .getValue();
438 case AMDGPU::BI__builtin_amdgcn_fmed3f:
439 case AMDGPU::BI__builtin_amdgcn_fmed3h:
440 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.fmed3").getValue();
441 case AMDGPU::BI__builtin_amdgcn_ds_append:
442 case AMDGPU::BI__builtin_amdgcn_ds_consume: {
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_global_load_tr_b64_i32:
449 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
450 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
451 return emitBuiltinWithOneOverloadedType<1>(
452 expr, "amdgcn.global.load.tr.b64", convertType(expr->getType()))
453 .getValue();
454 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
455 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
456 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
457 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
458 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
459 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
460 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
461 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
462 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:
463 return emitBuiltinWithOneOverloadedType<1>(
464 expr, "amdgcn.global.load.tr.b128", convertType(expr->getType()))
465 .getValue();
466 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
467 return emitBuiltinWithOneOverloadedType<1>(
468 expr, "amdgcn.global.load.tr4.b64", convertType(expr->getType()))
469 .getValue();
470 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
471 return emitBuiltinWithOneOverloadedType<1>(
472 expr, "amdgcn.global.load.tr6.b96", convertType(expr->getType()))
473 .getValue();
474 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
475 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr4.b64",
476 convertType(expr->getType()))
477 .getValue();
478 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
479 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr6.b96",
480 convertType(expr->getType()))
481 .getValue();
482 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
483 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr8.b64",
484 convertType(expr->getType()))
485 .getValue();
486 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
487 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
488 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:
489 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr16.b128",
490 convertType(expr->getType()))
491 .getValue();
492 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
493 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr4.b64",
494 convertType(expr->getType()))
495 .getValue();
496 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
497 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr8.b64",
498 convertType(expr->getType()))
499 .getValue();
500 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
501 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr6.b96",
502 convertType(expr->getType()))
503 .getValue();
504 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16:
505 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
506 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
507 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr16.b64",
508 convertType(expr->getType()))
509 .getValue();
510 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
511 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
512 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
513 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
514 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
515 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
516 cgm.errorNYI(expr->getSourceRange(),
517 std::string("unimplemented AMDGPU builtin call: ") +
518 getContext().BuiltinInfo.getName(builtinId));
519 return mlir::Value{};
520 }
521 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
522 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
523 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
524 cgm.errorNYI(expr->getSourceRange(),
525 std::string("unimplemented AMDGPU builtin call: ") +
526 getContext().BuiltinInfo.getName(builtinId));
527 return mlir::Value{};
528 }
529 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
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_cooperative_atomic_load_32x4B:
536 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
537 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
538 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
539 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
540 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
541 cgm.errorNYI(expr->getSourceRange(),
542 std::string("unimplemented AMDGPU builtin call: ") +
543 getContext().BuiltinInfo.getName(builtinId));
544 return mlir::Value{};
545 }
546 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
547 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
548 cgm.errorNYI(expr->getSourceRange(),
549 std::string("unimplemented AMDGPU builtin call: ") +
550 getContext().BuiltinInfo.getName(builtinId));
551 return mlir::Value{};
552 }
553 case AMDGPU::BI__builtin_amdgcn_read_exec:
554 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
555 case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
556 // The exec mask is read as a ballot over an all-true predicate. The
557 // ballot is at least as wide as the wavefront, so that a wave64 target
558 // still reports both halves for the _lo/_hi forms.
559 mlir::Location loc = getLoc(expr->getExprLoc());
560 unsigned registerWidth =
561 builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_lo ? 32 : 64;
562 unsigned ballotWidth =
563 std::max(getTarget().getGridValue().GV_Warp_Size, registerWidth);
564 cir::IntType ballotTy = builder.getUIntNTy(ballotWidth);
565
566 mlir::Value truePred = builder.getBool(true, loc).getResult();
567 mlir::Value result =
568 builder.emitIntrinsicCallOp(loc, "amdgcn.ballot", ballotTy, truePred);
569
570 if (builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_hi)
571 result = builder.createShiftRight(loc, result, 32);
572 return builder.createIntCast(result, convertType(expr->getType()));
573 }
574 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
575 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
576 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
577 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
578 cgm.errorNYI(expr->getSourceRange(),
579 std::string("unimplemented AMDGPU builtin call: ") +
580 getContext().BuiltinInfo.getName(builtinId));
581 return mlir::Value{};
582 }
583 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:
584 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
585 cgm.errorNYI(expr->getSourceRange(),
586 std::string("unimplemented AMDGPU builtin call: ") +
587 getContext().BuiltinInfo.getName(builtinId));
588 return mlir::Value{};
589 }
590 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:
591 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
592 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
593 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
594 cgm.errorNYI(expr->getSourceRange(),
595 std::string("unimplemented AMDGPU builtin call: ") +
596 getContext().BuiltinInfo.getName(builtinId));
597 return mlir::Value{};
598 }
599 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
600 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
601 return emitAMDGCNImageOverloadedReturnType(*this, expr,
602 "amdgcn.image.load.1d", false);
603 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
604 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
605 return emitAMDGCNImageOverloadedReturnType(
606 *this, expr, "amdgcn.image.load.1darray", false);
607 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
608 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
609 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
610 return emitAMDGCNImageOverloadedReturnType(*this, expr,
611 "amdgcn.image.load.2d", false);
612 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
613 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
614 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
615 return emitAMDGCNImageOverloadedReturnType(
616 *this, expr, "amdgcn.image.load.2darray", false);
617 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
618 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
619 return emitAMDGCNImageOverloadedReturnType(*this, expr,
620 "amdgcn.image.load.3d", false);
621 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
622 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
623 return emitAMDGCNImageOverloadedReturnType(*this, expr,
624 "amdgcn.image.load.cube", false);
625 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
626 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
627 return emitAMDGCNImageOverloadedReturnType(
628 *this, expr, "amdgcn.image.load.mip.1d", false);
629 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
630 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
631 return emitAMDGCNImageOverloadedReturnType(
632 *this, expr, "amdgcn.image.load.mip.1darray", false);
633 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
634 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
635 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
636 return emitAMDGCNImageOverloadedReturnType(
637 *this, expr, "amdgcn.image.load.mip.2d", false);
638 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
639 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
640 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
641 return emitAMDGCNImageOverloadedReturnType(
642 *this, expr, "amdgcn.image.load.mip.2darray", false);
643 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
644 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
645 return emitAMDGCNImageOverloadedReturnType(
646 *this, expr, "amdgcn.image.load.mip.3d", false);
647 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
648 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
649 return emitAMDGCNImageOverloadedReturnType(
650 *this, expr, "amdgcn.image.load.mip.cube", false);
651 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
652 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
653 return emitAMDGCNImageOverloadedReturnType(*this, expr,
654 "amdgcn.image.store.1d", true);
655 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
656 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
657 return emitAMDGCNImageOverloadedReturnType(
658 *this, expr, "amdgcn.image.store.1darray", true);
659 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
660 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
661 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
662 return emitAMDGCNImageOverloadedReturnType(*this, expr,
663 "amdgcn.image.store.2d", true);
664 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
665 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
666 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
667 return emitAMDGCNImageOverloadedReturnType(
668 *this, expr, "amdgcn.image.store.2darray", true);
669 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
670 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
671 return emitAMDGCNImageOverloadedReturnType(*this, expr,
672 "amdgcn.image.store.3d", true);
673 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
674 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
675 return emitAMDGCNImageOverloadedReturnType(*this, expr,
676 "amdgcn.image.store.cube", true);
677 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
678 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
679 return emitAMDGCNImageOverloadedReturnType(
680 *this, expr, "amdgcn.image.store.mip.1d", true);
681 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
682 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
683 return emitAMDGCNImageOverloadedReturnType(
684 *this, expr, "amdgcn.image.store.mip.1darray", true);
685 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
686 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
687 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
688 return emitAMDGCNImageOverloadedReturnType(
689 *this, expr, "amdgcn.image.store.mip.2d", true);
690 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
691 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
692 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
693 return emitAMDGCNImageOverloadedReturnType(
694 *this, expr, "amdgcn.image.store.mip.2darray", true);
695 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
696 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
697 return emitAMDGCNImageOverloadedReturnType(
698 *this, expr, "amdgcn.image.store.mip.3d", true);
699 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
700 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
701 return emitAMDGCNImageOverloadedReturnType(
702 *this, expr, "amdgcn.image.store.mip.cube", true);
703 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
704 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
705 return emitAMDGCNImageOverloadedReturnType(*this, expr,
706 "amdgcn.image.sample.1d", false);
707 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
708 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
709 return emitAMDGCNImageOverloadedReturnType(
710 *this, expr, "amdgcn.image.sample.1darray", false);
711 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
712 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
713 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
714 return emitAMDGCNImageOverloadedReturnType(*this, expr,
715 "amdgcn.image.sample.2d", false);
716 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
717 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
718 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
719 return emitAMDGCNImageOverloadedReturnType(
720 *this, expr, "amdgcn.image.sample.2darray", false);
721 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
722 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
723 return emitAMDGCNImageOverloadedReturnType(*this, expr,
724 "amdgcn.image.sample.3d", false);
725 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
726 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
727 return emitAMDGCNImageOverloadedReturnType(
728 *this, expr, "amdgcn.image.sample.cube", false);
729 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
730 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
731 return emitAMDGCNImageOverloadedReturnType(
732 *this, expr, "amdgcn.image.sample.lz.1d", false);
733 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
734 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
735 return emitAMDGCNImageOverloadedReturnType(
736 *this, expr, "amdgcn.image.sample.l.1d", false);
737 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
738 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
739 return emitAMDGCNImageOverloadedReturnType(
740 *this, expr, "amdgcn.image.sample.d.1d", false);
741 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
742 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
743 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
744 return emitAMDGCNImageOverloadedReturnType(
745 *this, expr, "amdgcn.image.sample.lz.2d", false);
746 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
747 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
748 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
749 return emitAMDGCNImageOverloadedReturnType(
750 *this, expr, "amdgcn.image.sample.l.2d", false);
751 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
752 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
753 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
754 return emitAMDGCNImageOverloadedReturnType(
755 *this, expr, "amdgcn.image.sample.d.2d", false);
756 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
757 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
758 return emitAMDGCNImageOverloadedReturnType(
759 *this, expr, "amdgcn.image.sample.lz.3d", false);
760 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
761 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
762 return emitAMDGCNImageOverloadedReturnType(
763 *this, expr, "amdgcn.image.sample.l.3d", false);
764 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
765 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
766 return emitAMDGCNImageOverloadedReturnType(
767 *this, expr, "amdgcn.image.sample.d.3d", false);
768 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
769 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
770 return emitAMDGCNImageOverloadedReturnType(
771 *this, expr, "amdgcn.image.sample.lz.cube", false);
772 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
773 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
774 return emitAMDGCNImageOverloadedReturnType(
775 *this, expr, "amdgcn.image.sample.l.cube", false);
776 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
777 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
778 return emitAMDGCNImageOverloadedReturnType(
779 *this, expr, "amdgcn.image.sample.lz.1darray", false);
780 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
781 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
782 return emitAMDGCNImageOverloadedReturnType(
783 *this, expr, "amdgcn.image.sample.l.1darray", false);
784 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
785 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
786 return emitAMDGCNImageOverloadedReturnType(
787 *this, expr, "amdgcn.image.sample.d.1darray", false);
788 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
789 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
790 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
791 return emitAMDGCNImageOverloadedReturnType(
792 *this, expr, "amdgcn.image.sample.lz.2darray", false);
793 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
794 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
795 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
796 return emitAMDGCNImageOverloadedReturnType(
797 *this, expr, "amdgcn.image.sample.l.2darray", false);
798 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
799 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
800 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
801 return emitAMDGCNImageOverloadedReturnType(
802 *this, expr, "amdgcn.image.sample.d.2darray", false);
803 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
804 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32:
805 return emitAMDGCNImageOverloadedReturnType(
806 *this, expr, "amdgcn.image.gather4.lz.2d", false);
807 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
808 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
809 cgm.errorNYI(expr->getSourceRange(),
810 std::string("unimplemented AMDGPU builtin call: ") +
811 getContext().BuiltinInfo.getName(builtinId));
812 return mlir::Value{};
813 }
814 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
815 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
816 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
817 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
818 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
819 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
820 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
821 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
822 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
823 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
824 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
825 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
826 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
827 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
828 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
829 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
830 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
831 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
832 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
833 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:
834 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:
835 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:
836 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:
837 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:
838 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:
839 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:
840 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:
841 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:
842 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:
843 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:
844 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:
845 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:
846 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:
847 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:
848 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:
849 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:
850 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:
851 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: {
852 cgm.errorNYI(expr->getSourceRange(),
853 std::string("unimplemented AMDGPU builtin call: ") +
854 getContext().BuiltinInfo.getName(builtinId));
855 return mlir::Value{};
856 }
857 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:
858 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:
859 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:
860 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:
861 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:
862 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:
863 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:
864 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:
865 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:
866 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:
867 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:
868 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:
869 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:
870 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:
871 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:
872 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:
873 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:
874 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:
875 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:
876 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:
877 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:
878 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {
879 cgm.errorNYI(expr->getSourceRange(),
880 std::string("unimplemented AMDGPU builtin call: ") +
881 getContext().BuiltinInfo.getName(builtinId));
882 return mlir::Value{};
883 }
884 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
885 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
886 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
887 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
888 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
889 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
890 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:
891 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:
892 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:
893 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:
894 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:
895 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:
896 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:
897 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:
898 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:
899 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:
900 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:
901 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:
902 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:
903 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:
904 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:
905 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:
906 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
907 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
908 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
909 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
910 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
911 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
912 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: {
913 cgm.errorNYI(expr->getSourceRange(),
914 std::string("unimplemented AMDGPU builtin call: ") +
915 getContext().BuiltinInfo.getName(builtinId));
916 return mlir::Value{};
917 }
918 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
919 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
920 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
921 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:
922 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:
923 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:
924 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:
925 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:
926 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:
927 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:
928 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:
929 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:
930 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:
931 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
932 cgm.errorNYI(expr->getSourceRange(),
933 std::string("unimplemented AMDGPU builtin call: ") +
934 getContext().BuiltinInfo.getName(builtinId));
935 return mlir::Value{};
936 }
937 // amdgcn workgroup size
938 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
939 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
940 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
941 cgm.errorNYI(expr->getSourceRange(),
942 std::string("unimplemented AMDGPU builtin call: ") +
943 getContext().BuiltinInfo.getName(builtinId));
944 return mlir::Value{};
945 }
946 case AMDGPU::BI__builtin_amdgcn_grid_size_x:
947 case AMDGPU::BI__builtin_amdgcn_grid_size_y:
948 case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
949 cgm.errorNYI(expr->getSourceRange(),
950 std::string("unimplemented AMDGPU builtin call: ") +
951 getContext().BuiltinInfo.getName(builtinId));
952 return mlir::Value{};
953 }
954 case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
955 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
956 cgm.errorNYI(expr->getSourceRange(),
957 std::string("unimplemented AMDGPU builtin call: ") +
958 getContext().BuiltinInfo.getName(builtinId));
959 return mlir::Value{};
960 }
961 case AMDGPU::BI__builtin_amdgcn_alignbit: {
962 cgm.errorNYI(expr->getSourceRange(),
963 std::string("unimplemented AMDGPU builtin call: ") +
964 getContext().BuiltinInfo.getName(builtinId));
965 return mlir::Value{};
966 }
967 case AMDGPU::BI__builtin_amdgcn_fence: {
968 cgm.errorNYI(expr->getSourceRange(),
969 std::string("unimplemented AMDGPU builtin call: ") +
970 getContext().BuiltinInfo.getName(builtinId));
971 return mlir::Value{};
972 }
973 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:
974 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
975 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
976 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
977 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
978 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
979 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
980 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:
981 case AMDGPU::BI__builtin_amdgcn_ds_faddf:
982 case AMDGPU::BI__builtin_amdgcn_ds_fminf:
983 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
984 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
985 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
986 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
987 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
988 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
989 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
990 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
991 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
992 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
993 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
994 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
995 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
996 cgm.errorNYI(expr->getSourceRange(),
997 std::string("unimplemented AMDGPU builtin call: ") +
998 getContext().BuiltinInfo.getName(builtinId));
999 return mlir::Value{};
1000 }
1001 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:
1002 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
1003 cgm.errorNYI(expr->getSourceRange(),
1004 std::string("unimplemented AMDGPU builtin call: ") +
1005 getContext().BuiltinInfo.getName(builtinId));
1006 return mlir::Value{};
1007 }
1008 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:
1009 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {
1010 cgm.errorNYI(expr->getSourceRange(),
1011 std::string("unimplemented AMDGPU builtin call: ") +
1012 getContext().BuiltinInfo.getName(builtinId));
1013 return mlir::Value{};
1014 }
1015 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
1016 case AMDGPU::BI__builtin_amdgcn_bitop3_b16:
1017 return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.bitop3")
1018 .getValue();
1019 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
1020 cgm.errorNYI(expr->getSourceRange(),
1021 std::string("unimplemented AMDGPU builtin call: ") +
1022 getContext().BuiltinInfo.getName(builtinId));
1023 return mlir::Value{};
1024 }
1025 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:
1026 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
1027 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
1028 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
1029 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
1030 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: {
1031 cgm.errorNYI(expr->getSourceRange(),
1032 std::string("unimplemented AMDGPU builtin call: ") +
1033 getContext().BuiltinInfo.getName(builtinId));
1034 return mlir::Value{};
1035 }
1036 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
1037 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
1038 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
1039 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
1040 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
1041 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
1042 cgm.errorNYI(expr->getSourceRange(),
1043 std::string("unimplemented AMDGPU builtin call: ") +
1044 getContext().BuiltinInfo.getName(builtinId));
1045 return mlir::Value{};
1046 }
1047 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
1048 cgm.errorNYI(expr->getSourceRange(),
1049 std::string("unimplemented AMDGPU builtin call: ") +
1050 getContext().BuiltinInfo.getName(builtinId));
1051 return mlir::Value{};
1052 }
1053 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
1054 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64:
1055 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
1056 cgm.errorNYI(expr->getSourceRange(),
1057 std::string("unimplemented AMDGPU builtin call: ") +
1058 getContext().BuiltinInfo.getName(builtinId));
1059 return mlir::Value{};
1060 }
1061 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
1062 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
1063 cgm.errorNYI(expr->getSourceRange(),
1064 std::string("unimplemented AMDGPU builtin call: ") +
1065 getContext().BuiltinInfo.getName(builtinId));
1066 return mlir::Value{};
1067 }
1068 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
1069 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
1070 cgm.errorNYI(expr->getSourceRange(),
1071 std::string("unimplemented AMDGPU builtin call: ") +
1072 getContext().BuiltinInfo.getName(builtinId));
1073 return mlir::Value{};
1074 }
1075 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: {
1076 cgm.errorNYI(expr->getSourceRange(),
1077 std::string("unimplemented AMDGPU builtin call: ") +
1078 getContext().BuiltinInfo.getName(builtinId));
1079 return mlir::Value{};
1080 }
1081 case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: {
1082 cgm.errorNYI(expr->getSourceRange(),
1083 std::string("unimplemented AMDGPU builtin call: ") +
1084 getContext().BuiltinInfo.getName(builtinId));
1085 return mlir::Value{};
1086 }
1087 case Builtin::BIlogbf:
1088 case Builtin::BI__builtin_logbf:
1089 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
1090 case Builtin::BIlogb:
1091 case Builtin::BI__builtin_logb:
1092 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble());
1093 case Builtin::BIscalbnf:
1094 case Builtin::BI__builtin_scalbnf:
1095 case Builtin::BIscalbn:
1096 case Builtin::BI__builtin_scalbn: {
1097 return emitBinaryExpMaybeConstrainedFPBuiltin(
1098 *this, expr, "ldexp", "experimental.constrained.ldexp");
1099 }
1100 default:
1101 return std::nullopt;
1102 }
1103}
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:820
cir::CIRBaseBuilderTy::createCompare
cir::CmpOp createCompare(mlir::Location loc, cir::CmpOpKind kind, mlir::Value lhs, mlir::Value rhs)
Definition CIRBaseBuilder.h:895
cir::CIRBaseBuilderTy::createSelect
mlir::Value createSelect(mlir::Location loc, mlir::Value condition, mlir::Value trueValue, mlir::Value falseValue)
Definition CIRBaseBuilder.h:767
cir::CIRBaseBuilderTy::getIsFPConstrained
bool getIsFPConstrained() const
Query for the use of constrained floating point math.
Definition CIRBaseBuilder.h:229
cir::CIRBaseBuilderTy::createAddrSpaceCast
mlir::Value createAddrSpaceCast(mlir::Location loc, mlir::Value src, mlir::Type newTy)
Definition CIRBaseBuilder.h:716
cir::RecordType
C++ view class that accepts both !cir.struct and !cir.union types.
Definition CIRTypes.h:149
cir::RecordType::getAllDataKinds
static llvm::SmallVector< RecordMemberKind > getAllDataKinds(llvm::ArrayRef< mlir::Type > members)
One Data kind per member.
Definition CIRTypes.cpp:162
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:826
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:395
clang::CIRGen::CIRGenFunction::convertType
mlir::Type convertType(clang::QualType t)
Definition CIRGenFunction.cpp:140
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:122
clang::CIRGen::CIRGenFunction::getTarget
const TargetInfo & getTarget() const
Definition CIRGenFunction.h:351
clang::CIRGen::CIRGenFunction::getLoc
mlir::Location getLoc(clang::SourceLocation srcLoc)
Helpers to convert Clang's SourceLocation to a MLIR Location.
Definition CIRGenFunction.cpp:144
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:1837
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:1857
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:4128
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:2987
clang::CallExpr::getArg
Expr * getArg(unsigned Arg)
getArg - Return the specified argument.
Definition Expr.h:3191
clang::CallExpr::getNumArgs
unsigned getNumArgs() const
getNumArgs - Return the number of actual arguments to this call.
Definition Expr.h:3178
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