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/IR/IntrinsicsAMDGPU.h"
18#include "llvm/Support/AMDGPUAddrSpace.h"
19#include "llvm/Support/ErrorHandling.h"
20
21using namespace clang;
22using namespace clang::CIRGen;
23using namespace cir;
24
25// Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the
26// result to match \p e's return type when needed.
27// If \p e is null, returns the raw AS-4 pointer.
28static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
29 const CallExpr *e = nullptr) {
30 CIRGenBuilderTy &builder = cgf.getBuilder();
31 mlir::Location loc =
32 e ? cgf.getLoc(e->getExprLoc()) : builder.getUnknownLoc();
33 // The intrinsic always returns a pointer in the constant AS.
34 mlir::Type retTy = cir::PointerType::get(
35 cir::VoidType::get(builder.getContext()),
36 cir::TargetAddressSpaceAttr::get(builder.getContext(),
37 llvm::AMDGPUAS::CONSTANT_ADDRESS));
38 mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr",
39 retTy, mlir::ValueRange{});
40 if (!e)
41 return call;
42 // Only cast when the caller-visible AS differs from the intrinsic's AS;
43 auto expectedPtrTy =
44 mlir::cast<cir::PointerType>(cgf.convertType(e->getType()));
45 auto callPtrTy = mlir::cast<cir::PointerType>(call.getType());
46 if (expectedPtrTy.getAddrSpace() == callPtrTy.getAddrSpace())
47 return call;
48 return builder.createAddrSpaceCast(loc, call, expectedPtrTy);
49}
50
51static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
52 CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName,
53 llvm::StringRef constrainedIntrinsicName) {
54 mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0));
55 mlir::Value src1 = cgf.emitScalarExpr(e->getArg(1));
56 mlir::Location loc = cgf.getLoc(e->getExprLoc());
57
58 CIRGenBuilderTy &builder = cgf.getBuilder();
59
60 CIRGenFunction::CIRGenFPOptionsRAII fpOptsRAII(cgf, e);
61
62 if (builder.getIsFPConstrained()) {
63 cgf.cgm.errorNYI(e->getSourceRange(),
64 "constrained FP intrinsic support is NYI.");
65 }
66
67 return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(),
68 mlir::ValueRange{src0, src1});
69}
70
71static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e,
72 const llvm::fltSemantics &fSem) {
73 CIRGenBuilderTy &builder = cgf.getBuilder();
74 mlir::Location loc = cgf.getLoc(e->getExprLoc());
75
76 mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0));
77 mlir::Type srcTy = src0.getType();
78 mlir::Type int32Ty = builder.getSInt32Ty();
79
80 mlir::Type frExpResMembers[] = {srcTy, int32Ty};
81 cir::RecordType frExpResTy = builder.getAnonRecordTy(
82 frExpResMembers, /*packed=*/false,
83 cir::RecordType::getAllDataKinds(frExpResMembers));
84
85 mlir::Value frExpResult = builder.emitIntrinsicCallOp(
86 loc, "frexp", frExpResTy, mlir::ValueRange{src0});
87
88 mlir::Value exp =
89 cir::ExtractMemberOp::create(builder, loc, int32Ty, frExpResult, 1);
90
91 mlir::Value negativeOne =
92 builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1));
93 mlir::Value expMinus1 = builder.createNSWAdd(loc, exp, negativeOne);
94
95 mlir::Value siToFp = cir::CastOp::create(
96 builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1);
97
98 mlir::Value fabs = cir::FAbsOp::create(builder, loc, src0);
99
100 llvm::APFloat infVal = llvm::APFloat::getInf(fSem);
101 mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal));
102
103 mlir::Value fabsNegInf =
104 builder.createCompare(loc, cir::CmpOpKind::one, fabs, inf);
105
106 mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs);
107
108 llvm::APFloat zeroValue = llvm::APFloat::getZero(fSem);
109 mlir::Value zero =
110 builder.getConstant(loc, cir::FPAttr::get(srcTy, zeroValue));
111
112 mlir::Value srcEqZero =
113 builder.createCompare(loc, cir::CmpOpKind::eq, src0, zero);
114
115 llvm::APFloat negInfVal = llvm::APFloat::getInf(fSem, true);
116 mlir::Value negInf =
117 builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal));
118
119 mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel);
120
121 return res;
122}
123
124static mlir::Value
125emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e,
126 llvm::StringRef intrinsicName,
127 bool isImageStore) {
128 auto &builder = cgf.getBuilder();
129
130 llvm::SmallVector<mlir::Value, 10> args;
131 for (unsigned i = 0, n = e->getNumArgs(); i < n; ++i)
132 args.push_back(cgf.emitScalarExpr(e->getArg(i)));
133
134 mlir::Type retTy = isImageStore ? cir::VoidType::get(builder.getContext())
135 : cgf.convertType(e->getType());
136
137 auto callOp = cir::LLVMIntrinsicCallOp::create(
138 builder, cgf.getLoc(e->getExprLoc()),
139 builder.getStringAttr(intrinsicName), retTy, args);
140
141 return callOp.getResult();
142}
143
144std::optional<mlir::Value>
145CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
146 const CallExpr *expr) {
147 switch (builtinId) {
148 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:
149 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:
150 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:
151 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:
152 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:
153 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:
154 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:
155 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:
156 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:
157 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:
158 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:
159 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:
160 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:
161 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:
162 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:
163 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:
164 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:
165 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {
166 cgm.errorNYI(expr->getSourceRange(),
167 std::string("unimplemented AMDGPU builtin call: ") +
168 getContext().BuiltinInfo.getName(builtinId));
169 return mlir::Value{};
170 }
171 case AMDGPU::BI__builtin_amdgcn_div_scale:
172 case AMDGPU::BI__builtin_amdgcn_div_scalef: {
173 Address flagOutPtr = emitPointerWithAlignment(expr->getArg(3));
174 llvm::StringRef intrinsicName = "amdgcn.div.scale";
175 mlir::Value x = emitScalarExpr(expr->getArg(0));
176 mlir::Value y = emitScalarExpr(expr->getArg(1));
177 mlir::Value z = emitScalarExpr(expr->getArg(2));
178
179 auto i1Ty = builder.getUIntNTy(1);
180 mlir::Type resMembers[] = {x.getType(), i1Ty};
181 cir::RecordType resTy =
182 builder.getAnonRecordTy(resMembers, /*packed=*/false,
183 cir::RecordType::getAllDataKinds(resMembers));
184
185 mlir::Value structResult =
186 cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()),
187 builder.getStringAttr(intrinsicName),
188 resTy, {x, y, z})
189 .getResult();
190
191 mlir::Value result = cir::ExtractMemberOp::create(
192 builder, getLoc(expr->getExprLoc()), x.getType(), structResult, 0);
193 mlir::Value flag = cir::ExtractMemberOp::create(
194 builder, getLoc(expr->getExprLoc()), i1Ty, structResult, 1);
195
196 mlir::Type flagType = flagOutPtr.getElementType();
197 mlir::Value flagToStore =
198 cir::CastOp::create(builder, getLoc(expr->getExprLoc()), flagType,
199 cir::CastKind::int_to_bool, flag);
200 builder.createStore(getLoc(expr->getExprLoc()), flagToStore, flagOutPtr);
201 return result;
202 }
203 case AMDGPU::BI__builtin_amdgcn_div_fmas:
204 case AMDGPU::BI__builtin_amdgcn_div_fmasf:
205 return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas")
206 .getValue();
207 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
208 case AMDGPU::BI__builtin_amdgcn_mov_dpp:
209 case AMDGPU::BI__builtin_amdgcn_update_dpp: {
210 mlir::Location loc = getLoc(expr->getExprLoc());
211 unsigned iceArguments = 0;
212 ASTContext::GetBuiltinTypeError error;
213 getContext().GetBuiltinType(builtinId, error, &iceArguments);
214 assert(error == ASTContext::GE_None && "Should not codegen an error");
215 assert(expr->getNumArgs() == 5 || expr->getNumArgs() == 6 ||
216 expr->getNumArgs() == 2);
217
218 mlir::Type dataTy = convertType(expr->getArg(0)->getType());
219 unsigned size = cgm.getDataLayout().getTypeSizeInBits(dataTy);
220 cir::IntType intTy = builder.getUIntNTy(std::max(size, 32u));
221
222 bool isMovDpp8 = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp8;
223 bool isMovDpp = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp;
224 bool isUpdateDpp = builtinId == AMDGPU::BI__builtin_amdgcn_update_dpp;
225 llvm::StringRef intrinsicName =
226 isMovDpp8 ? "amdgcn.mov.dpp8" : "amdgcn.update.dpp";
227 cir::FuncType intrinsicTy =
228 getIntrinsicType(isMovDpp8 ? llvm::Intrinsic::amdgcn_mov_dpp8
229 : llvm::Intrinsic::amdgcn_update_dpp,
230 {intTy});
231
232 auto coerceTo = [&](mlir::Value from, mlir::Type to) -> mlir::Value {
233 if (from.getType() == to)
234 return from;
235 if (mlir::isa<cir::IntType>(from.getType()) &&
236 mlir::isa<cir::IntType>(to))
237 return builder.createIntCast(from, to);
238 return builder.createBitcast(from, to);
239 };
240
241 llvm::SmallVector<mlir::Value, 6> args;
242 // __builtin_amdgcn_mov_dpp has no "old" operand at the source level, but
243 // the real intrinsic it lowers to requires one, so we synthesize a poison
244 // value for it since it is never meaningfully read.
245 if (isMovDpp)
246 args.push_back(builder.getConstant(loc, cir::PoisonAttr::get(intTy)));
247
248 // Number of builtin-level leading args that need zero-extend promotion when
249 // the data type is narrower than 32 bits.
250 unsigned numPromotedArgs = isUpdateDpp ? 2u : 1u;
251 for (unsigned i = 0; i != expr->getNumArgs(); ++i) {
252 mlir::Value v =
253 emitScalarOrConstFoldImmArg(iceArguments, i, expr->getArg(i));
254 if (i < numPromotedArgs && size < 32) {
255 mlir::Type sameWidthUTy = builder.getUIntNTy(size);
256 if (v.getType() != sameWidthUTy)
257 v = builder.createBitcast(v, sameWidthUTy);
258 v = builder.createIntCast(v, intTy);
259 }
260 args.push_back(coerceTo(v, intrinsicTy.getInput(i + unsigned(isMovDpp))));
261 }
262
263 mlir::Value result = builder.emitIntrinsicCallOp(loc, intrinsicName, intTy,
264 mlir::ValueRange(args));
265 if (size < 32 && !mlir::isa<cir::IntType>(dataTy))
266 result = builder.createIntCast(result, builder.getUIntNTy(size));
267 return coerceTo(result, dataTy);
268 }
269 case AMDGPU::BI__builtin_amdgcn_permlane16:
270 case AMDGPU::BI__builtin_amdgcn_permlanex16: {
271 llvm::StringRef intrinsicName =
272 builtinId == AMDGPU::BI__builtin_amdgcn_permlane16
273 ? "amdgcn.permlane16"
274 : "amdgcn.permlanex16";
275 return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName).getValue();
276 }
277 case AMDGPU::BI__builtin_amdgcn_permlane64:
278 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64")
279 .getValue();
280 case AMDGPU::BI__builtin_amdgcn_permlane_bcast:
281 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.bcast")
282 .getValue();
283 case AMDGPU::BI__builtin_amdgcn_permlane_up:
284 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.up")
285 .getValue();
286 case AMDGPU::BI__builtin_amdgcn_permlane_down:
287 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.down")
288 .getValue();
289 case AMDGPU::BI__builtin_amdgcn_permlane_xor:
290 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.xor")
291 .getValue();
292 case AMDGPU::BI__builtin_amdgcn_readlane:
293 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane")
294 .getValue();
295 case AMDGPU::BI__builtin_amdgcn_readfirstlane:
296 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane")
297 .getValue();
298 case AMDGPU::BI__builtin_amdgcn_wave_shuffle:
299 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.wave.shuffle")
300 .getValue();
301 case AMDGPU::BI__builtin_amdgcn_div_fixup:
302 case AMDGPU::BI__builtin_amdgcn_div_fixupf:
303 case AMDGPU::BI__builtin_amdgcn_div_fixuph: {
304 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
305 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
306 mlir::Value src2 = emitScalarExpr(expr->getArg(2));
307 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
308 "amdgcn.div.fixup", src0.getType(),
309 mlir::ValueRange{src0, src1, src2});
310 }
311 case AMDGPU::BI__builtin_amdgcn_trig_preop:
312 case AMDGPU::BI__builtin_amdgcn_trig_preopf:
313 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.trig.preop")
314 .getValue();
315 case AMDGPU::BI__builtin_amdgcn_rcp:
316 case AMDGPU::BI__builtin_amdgcn_rcpf:
317 case AMDGPU::BI__builtin_amdgcn_rcph:
318 case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
319 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue();
320 }
321 case AMDGPU::BI__builtin_amdgcn_sqrt:
322 case AMDGPU::BI__builtin_amdgcn_sqrtf:
323 case AMDGPU::BI__builtin_amdgcn_sqrth:
324 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: {
325 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt").getValue();
326 }
327 case AMDGPU::BI__builtin_amdgcn_rsq:
328 case AMDGPU::BI__builtin_amdgcn_rsqf:
329 case AMDGPU::BI__builtin_amdgcn_rsqh:
330 case AMDGPU::BI__builtin_amdgcn_rsq_bf16: {
331 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue();
332 }
333 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
334 case AMDGPU::BI__builtin_amdgcn_rsq_clampf: {
335 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp")
336 .getValue();
337 }
338 case AMDGPU::BI__builtin_amdgcn_sinf:
339 case AMDGPU::BI__builtin_amdgcn_sinh:
340 case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
341 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue();
342 }
343 case AMDGPU::BI__builtin_amdgcn_cosf:
344 case AMDGPU::BI__builtin_amdgcn_cosh:
345 case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
346 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue();
347 }
348 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
349 return emitAMDGPUDispatchPtr(*this, expr);
350 case AMDGPU::BI__builtin_amdgcn_logf:
351 case AMDGPU::BI__builtin_amdgcn_log_bf16: {
352 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.log").getValue();
353 }
354 case AMDGPU::BI__builtin_amdgcn_exp2f:
355 case AMDGPU::BI__builtin_amdgcn_exp2_bf16: {
356 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.exp2").getValue();
357 }
358 case AMDGPU::BI__builtin_amdgcn_log_clampf: {
359 cgm.errorNYI(expr->getSourceRange(),
360 std::string("unimplemented AMDGPU builtin call: ") +
361 getContext().BuiltinInfo.getName(builtinId));
362 return mlir::Value{};
363 }
364 case AMDGPU::BI__builtin_amdgcn_ldexp:
365 case AMDGPU::BI__builtin_amdgcn_ldexpf:
366 case AMDGPU::BI__builtin_amdgcn_ldexph: {
367 mlir::Value src0 = emitScalarExpr(expr->getArg(0));
368 mlir::Value src1 = emitScalarExpr(expr->getArg(1));
369 // For ldexph, the raw instruction has different behavior for out-of-bounds
370 // exponent values (implicit truncation instead of saturate to
371 // short_min/short_max), so truncate the exponent to i16 first.
372 mlir::Value exp =
373 builtinId == AMDGPU::BI__builtin_amdgcn_ldexph
374 ? cir::CastOp::create(builder, getLoc(expr->getExprLoc()),
375 builder.getSInt16Ty(),
376 cir::CastKind::integral, src1)
377 : src1;
378 return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
379 src0.getType(),
380 mlir::ValueRange{src0, exp});
381 }
382 case AMDGPU::BI__builtin_amdgcn_frexp_mant:
383 case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
384 case AMDGPU::BI__builtin_amdgcn_frexp_manth: {
385 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant")
386 .getValue();
387 }
388 case AMDGPU::BI__builtin_amdgcn_frexp_exp:
389 case AMDGPU::BI__builtin_amdgcn_frexp_expf:
390 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {
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_fract:
397 case AMDGPU::BI__builtin_amdgcn_fractf:
398 case AMDGPU::BI__builtin_amdgcn_fracth:
399 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.fract").getValue();
400 case AMDGPU::BI__builtin_amdgcn_lerp: {
401 cgm.errorNYI(expr->getSourceRange(),
402 std::string("unimplemented AMDGPU builtin call: ") +
403 getContext().BuiltinInfo.getName(builtinId));
404 return mlir::Value{};
405 }
406 case AMDGPU::BI__builtin_amdgcn_ubfe:
407 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.ubfe").getValue();
408 case AMDGPU::BI__builtin_amdgcn_sbfe:
409 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.sbfe").getValue();
410 case AMDGPU::BI__builtin_amdgcn_ballot_w32:
411 case AMDGPU::BI__builtin_amdgcn_ballot_w64:
412 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot",
413 convertType(expr->getType()))
414 .getValue();
415 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
416 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64:
417 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.inverse.ballot",
418 convertType(expr->getType()))
419 .getValue();
420 case AMDGPU::BI__builtin_amdgcn_tanhf:
421 case AMDGPU::BI__builtin_amdgcn_tanhh:
422 case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
423 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue();
424 }
425 case AMDGPU::BI__builtin_amdgcn_uicmp:
426 case AMDGPU::BI__builtin_amdgcn_uicmpl:
427 case AMDGPU::BI__builtin_amdgcn_sicmp:
428 case AMDGPU::BI__builtin_amdgcn_sicmpl:
429 case AMDGPU::BI__builtin_amdgcn_fcmp:
430 case AMDGPU::BI__builtin_amdgcn_fcmpf: {
431 mlir::Value lhs = emitScalarExpr(expr->getArg(0));
432 mlir::Value rhs = emitScalarExpr(expr->getArg(1));
433
434 uint64_t imm =
435 expr->getArg(2)->EvaluateKnownConstInt(getContext()).getZExtValue();
436
437 cir::CmpOpKind pred;
438 switch (imm) {
439 case 0x1: // FCMP_OEQ
440 case 0x20: // ICMP_EQ
441 pred = cir::CmpOpKind::eq;
442 break;
443 case 0xe: // FCMP_UNE
444 case 0x21: // ICMP_NE
445 pred = cir::CmpOpKind::ne;
446 break;
447 case 0x2: // FCMP_OGT
448 case 0x22: // ICMP_UGT
449 case 0x26: // ICMP_SGT
450 pred = cir::CmpOpKind::gt;
451 break;
452 case 0x3: // FCMP_OGE
453 case 0x23: // ICMP_UGE
454 case 0x27: // ICMP_SGE
455 pred = cir::CmpOpKind::ge;
456 break;
457 case 0x4: // FCMP_OLT
458 case 0x24: // ICMP_ULT
459 case 0x28: // ICMP_SLT
460 pred = cir::CmpOpKind::lt;
461 break;
462 case 0x5: // FCMP_OLE
463 case 0x25: // ICMP_ULE
464 case 0x29: // ICMP_SLE
465 pred = cir::CmpOpKind::le;
466 break;
467 case 0x6: // FCMP_ONE
468 pred = cir::CmpOpKind::one;
469 break;
470 case 0x8: // FCMP_UNO
471 pred = cir::CmpOpKind::uno;
472 break;
473 default:
474 cgm.errorNYI(expr->getSourceRange(),
475 "amdgcn compare with unsupported predicate");
476 return mlir::Value{};
477 }
478
479 mlir::Location loc = getLoc(expr->getExprLoc());
480 mlir::Value cmp = builder.createCompare(loc, pred, lhs, rhs);
481 return builder.emitIntrinsicCallOp(loc, "amdgcn.ballot",
482 convertType(expr->getType()), cmp);
483 }
484 case AMDGPU::BI__builtin_amdgcn_class:
485 case AMDGPU::BI__builtin_amdgcn_classf:
486 case AMDGPU::BI__builtin_amdgcn_classh:
487 return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.class",
488 convertType(expr->getType()))
489 .getValue();
490 case AMDGPU::BI__builtin_amdgcn_fmed3f:
491 case AMDGPU::BI__builtin_amdgcn_fmed3h:
492 return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.fmed3").getValue();
493 case AMDGPU::BI__builtin_amdgcn_ds_append:
494 case AMDGPU::BI__builtin_amdgcn_ds_consume: {
495 cgm.errorNYI(expr->getSourceRange(),
496 std::string("unimplemented AMDGPU builtin call: ") +
497 getContext().BuiltinInfo.getName(builtinId));
498 return mlir::Value{};
499 }
500 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:
501 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:
502 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:
503 return emitBuiltinWithOneOverloadedType<1>(
504 expr, "amdgcn.global.load.tr.b64", convertType(expr->getType()))
505 .getValue();
506 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:
507 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:
508 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:
509 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:
510 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:
511 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:
512 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:
513 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:
514 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:
515 return emitBuiltinWithOneOverloadedType<1>(
516 expr, "amdgcn.global.load.tr.b128", convertType(expr->getType()))
517 .getValue();
518 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:
519 return emitBuiltinWithOneOverloadedType<1>(
520 expr, "amdgcn.global.load.tr4.b64", convertType(expr->getType()))
521 .getValue();
522 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:
523 return emitBuiltinWithOneOverloadedType<1>(
524 expr, "amdgcn.global.load.tr6.b96", convertType(expr->getType()))
525 .getValue();
526 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:
527 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr4.b64",
528 convertType(expr->getType()))
529 .getValue();
530 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:
531 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr6.b96",
532 convertType(expr->getType()))
533 .getValue();
534 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:
535 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr8.b64",
536 convertType(expr->getType()))
537 .getValue();
538 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:
539 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:
540 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:
541 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr16.b128",
542 convertType(expr->getType()))
543 .getValue();
544 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:
545 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr4.b64",
546 convertType(expr->getType()))
547 .getValue();
548 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:
549 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr8.b64",
550 convertType(expr->getType()))
551 .getValue();
552 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:
553 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr6.b96",
554 convertType(expr->getType()))
555 .getValue();
556 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16:
557 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:
558 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:
559 return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr16.b64",
560 convertType(expr->getType()))
561 .getValue();
562 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:
563 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
564 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
565 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:
566 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:
567 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {
568 cgm.errorNYI(expr->getSourceRange(),
569 std::string("unimplemented AMDGPU builtin call: ") +
570 getContext().BuiltinInfo.getName(builtinId));
571 return mlir::Value{};
572 }
573 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:
574 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:
575 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {
576 cgm.errorNYI(expr->getSourceRange(),
577 std::string("unimplemented AMDGPU builtin call: ") +
578 getContext().BuiltinInfo.getName(builtinId));
579 return mlir::Value{};
580 }
581 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {
582 cgm.errorNYI(expr->getSourceRange(),
583 std::string("unimplemented AMDGPU builtin call: ") +
584 getContext().BuiltinInfo.getName(builtinId));
585 return mlir::Value{};
586 }
587 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:
588 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:
589 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:
590 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
591 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:
592 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {
593 cgm.errorNYI(expr->getSourceRange(),
594 std::string("unimplemented AMDGPU builtin call: ") +
595 getContext().BuiltinInfo.getName(builtinId));
596 return mlir::Value{};
597 }
598 case AMDGPU::BI__builtin_amdgcn_get_fpenv:
599 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {
600 cgm.errorNYI(expr->getSourceRange(),
601 std::string("unimplemented AMDGPU builtin call: ") +
602 getContext().BuiltinInfo.getName(builtinId));
603 return mlir::Value{};
604 }
605 case AMDGPU::BI__builtin_amdgcn_read_exec:
606 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
607 case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
608 // The exec mask is read as a ballot over an all-true predicate. The
609 // ballot is at least as wide as the wavefront, so that a wave64 target
610 // still reports both halves for the _lo/_hi forms.
611 mlir::Location loc = getLoc(expr->getExprLoc());
612 unsigned registerWidth =
613 builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_lo ? 32 : 64;
614 unsigned ballotWidth =
615 std::max(getTarget().getGridValue().GV_Warp_Size, registerWidth);
616 cir::IntType ballotTy = builder.getUIntNTy(ballotWidth);
617
618 mlir::Value truePred = builder.getBool(true, loc).getResult();
619 mlir::Value result =
620 builder.emitIntrinsicCallOp(loc, "amdgcn.ballot", ballotTy, truePred);
621
622 if (builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_hi)
623 result = builder.createShiftRight(loc, result, 32);
624 return builder.createIntCast(result, convertType(expr->getType()));
625 }
626 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
627 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
628 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:
629 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {
630 cgm.errorNYI(expr->getSourceRange(),
631 std::string("unimplemented AMDGPU builtin call: ") +
632 getContext().BuiltinInfo.getName(builtinId));
633 return mlir::Value{};
634 }
635 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:
636 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {
637 cgm.errorNYI(expr->getSourceRange(),
638 std::string("unimplemented AMDGPU builtin call: ") +
639 getContext().BuiltinInfo.getName(builtinId));
640 return mlir::Value{};
641 }
642 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:
643 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:
644 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:
645 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {
646 cgm.errorNYI(expr->getSourceRange(),
647 std::string("unimplemented AMDGPU builtin call: ") +
648 getContext().BuiltinInfo.getName(builtinId));
649 return mlir::Value{};
650 }
651 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
652 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
653 return emitAMDGCNImageOverloadedReturnType(*this, expr,
654 "amdgcn.image.load.1d", false);
655 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
656 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
657 return emitAMDGCNImageOverloadedReturnType(
658 *this, expr, "amdgcn.image.load.1darray", false);
659 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
660 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
661 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
662 return emitAMDGCNImageOverloadedReturnType(*this, expr,
663 "amdgcn.image.load.2d", false);
664 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
665 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
666 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
667 return emitAMDGCNImageOverloadedReturnType(
668 *this, expr, "amdgcn.image.load.2darray", false);
669 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
670 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
671 return emitAMDGCNImageOverloadedReturnType(*this, expr,
672 "amdgcn.image.load.3d", false);
673 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
674 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
675 return emitAMDGCNImageOverloadedReturnType(*this, expr,
676 "amdgcn.image.load.cube", false);
677 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
678 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
679 return emitAMDGCNImageOverloadedReturnType(
680 *this, expr, "amdgcn.image.load.mip.1d", false);
681 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
682 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
683 return emitAMDGCNImageOverloadedReturnType(
684 *this, expr, "amdgcn.image.load.mip.1darray", false);
685 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
686 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
687 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
688 return emitAMDGCNImageOverloadedReturnType(
689 *this, expr, "amdgcn.image.load.mip.2d", false);
690 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
691 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
692 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
693 return emitAMDGCNImageOverloadedReturnType(
694 *this, expr, "amdgcn.image.load.mip.2darray", false);
695 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
696 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
697 return emitAMDGCNImageOverloadedReturnType(
698 *this, expr, "amdgcn.image.load.mip.3d", false);
699 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
700 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
701 return emitAMDGCNImageOverloadedReturnType(
702 *this, expr, "amdgcn.image.load.mip.cube", false);
703 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
704 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
705 return emitAMDGCNImageOverloadedReturnType(*this, expr,
706 "amdgcn.image.store.1d", true);
707 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
708 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
709 return emitAMDGCNImageOverloadedReturnType(
710 *this, expr, "amdgcn.image.store.1darray", true);
711 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
712 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
713 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
714 return emitAMDGCNImageOverloadedReturnType(*this, expr,
715 "amdgcn.image.store.2d", true);
716 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
717 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
718 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
719 return emitAMDGCNImageOverloadedReturnType(
720 *this, expr, "amdgcn.image.store.2darray", true);
721 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
722 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
723 return emitAMDGCNImageOverloadedReturnType(*this, expr,
724 "amdgcn.image.store.3d", true);
725 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
726 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
727 return emitAMDGCNImageOverloadedReturnType(*this, expr,
728 "amdgcn.image.store.cube", true);
729 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
730 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
731 return emitAMDGCNImageOverloadedReturnType(
732 *this, expr, "amdgcn.image.store.mip.1d", true);
733 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
734 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
735 return emitAMDGCNImageOverloadedReturnType(
736 *this, expr, "amdgcn.image.store.mip.1darray", true);
737 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
738 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
739 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
740 return emitAMDGCNImageOverloadedReturnType(
741 *this, expr, "amdgcn.image.store.mip.2d", true);
742 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
743 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
744 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
745 return emitAMDGCNImageOverloadedReturnType(
746 *this, expr, "amdgcn.image.store.mip.2darray", true);
747 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
748 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
749 return emitAMDGCNImageOverloadedReturnType(
750 *this, expr, "amdgcn.image.store.mip.3d", true);
751 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
752 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:
753 return emitAMDGCNImageOverloadedReturnType(
754 *this, expr, "amdgcn.image.store.mip.cube", true);
755 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
756 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
757 return emitAMDGCNImageOverloadedReturnType(*this, expr,
758 "amdgcn.image.sample.1d", false);
759 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
760 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
761 return emitAMDGCNImageOverloadedReturnType(
762 *this, expr, "amdgcn.image.sample.1darray", false);
763 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
764 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
765 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
766 return emitAMDGCNImageOverloadedReturnType(*this, expr,
767 "amdgcn.image.sample.2d", false);
768 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
769 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
770 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
771 return emitAMDGCNImageOverloadedReturnType(
772 *this, expr, "amdgcn.image.sample.2darray", false);
773 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
774 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
775 return emitAMDGCNImageOverloadedReturnType(*this, expr,
776 "amdgcn.image.sample.3d", false);
777 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
778 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:
779 return emitAMDGCNImageOverloadedReturnType(
780 *this, expr, "amdgcn.image.sample.cube", false);
781 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:
782 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:
783 return emitAMDGCNImageOverloadedReturnType(
784 *this, expr, "amdgcn.image.sample.lz.1d", false);
785 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:
786 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:
787 return emitAMDGCNImageOverloadedReturnType(
788 *this, expr, "amdgcn.image.sample.l.1d", false);
789 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:
790 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:
791 return emitAMDGCNImageOverloadedReturnType(
792 *this, expr, "amdgcn.image.sample.d.1d", false);
793 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:
794 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:
795 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:
796 return emitAMDGCNImageOverloadedReturnType(
797 *this, expr, "amdgcn.image.sample.lz.2d", false);
798 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:
799 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:
800 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:
801 return emitAMDGCNImageOverloadedReturnType(
802 *this, expr, "amdgcn.image.sample.l.2d", false);
803 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:
804 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:
805 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:
806 return emitAMDGCNImageOverloadedReturnType(
807 *this, expr, "amdgcn.image.sample.d.2d", false);
808 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:
809 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:
810 return emitAMDGCNImageOverloadedReturnType(
811 *this, expr, "amdgcn.image.sample.lz.3d", false);
812 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:
813 case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:
814 return emitAMDGCNImageOverloadedReturnType(
815 *this, expr, "amdgcn.image.sample.l.3d", false);
816 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
817 case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
818 return emitAMDGCNImageOverloadedReturnType(
819 *this, expr, "amdgcn.image.sample.d.3d", false);
820 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:
821 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:
822 return emitAMDGCNImageOverloadedReturnType(
823 *this, expr, "amdgcn.image.sample.lz.cube", false);
824 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:
825 case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:
826 return emitAMDGCNImageOverloadedReturnType(
827 *this, expr, "amdgcn.image.sample.l.cube", false);
828 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:
829 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:
830 return emitAMDGCNImageOverloadedReturnType(
831 *this, expr, "amdgcn.image.sample.lz.1darray", false);
832 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:
833 case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:
834 return emitAMDGCNImageOverloadedReturnType(
835 *this, expr, "amdgcn.image.sample.l.1darray", false);
836 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:
837 case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:
838 return emitAMDGCNImageOverloadedReturnType(
839 *this, expr, "amdgcn.image.sample.d.1darray", false);
840 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:
841 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:
842 case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:
843 return emitAMDGCNImageOverloadedReturnType(
844 *this, expr, "amdgcn.image.sample.lz.2darray", false);
845 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:
846 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:
847 case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:
848 return emitAMDGCNImageOverloadedReturnType(
849 *this, expr, "amdgcn.image.sample.l.2darray", false);
850 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:
851 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
852 case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:
853 return emitAMDGCNImageOverloadedReturnType(
854 *this, expr, "amdgcn.image.sample.d.2darray", false);
855 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
856 case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32:
857 return emitAMDGCNImageOverloadedReturnType(
858 *this, expr, "amdgcn.image.gather4.lz.2d", false);
859 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
860 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {
861 cgm.errorNYI(expr->getSourceRange(),
862 std::string("unimplemented AMDGPU builtin call: ") +
863 getContext().BuiltinInfo.getName(builtinId));
864 return mlir::Value{};
865 }
866 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:
867 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:
868 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:
869 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:
870 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:
871 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:
872 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:
873 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:
874 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:
875 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:
876 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:
877 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:
878 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:
879 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:
880 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:
881 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:
882 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:
883 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:
884 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:
885 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:
886 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:
887 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:
888 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:
889 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:
890 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:
891 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:
892 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:
893 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:
894 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:
895 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:
896 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:
897 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:
898 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:
899 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:
900 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:
901 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:
902 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:
903 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: {
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_swmmac_f32_16x16x32_f16_w32:
910 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:
911 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:
912 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:
913 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:
914 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:
915 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:
916 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:
917 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:
918 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:
919 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:
920 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:
921 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:
922 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:
923 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:
924 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:
925 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:
926 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:
927 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:
928 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:
929 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:
930 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {
931 cgm.errorNYI(expr->getSourceRange(),
932 std::string("unimplemented AMDGPU builtin call: ") +
933 getContext().BuiltinInfo.getName(builtinId));
934 return mlir::Value{};
935 }
936 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:
937 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:
938 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:
939 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:
940 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:
941 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:
942 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:
943 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:
944 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:
945 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:
946 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:
947 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:
948 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:
949 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:
950 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:
951 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:
952 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:
953 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:
954 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:
955 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:
956 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:
957 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:
958 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:
959 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:
960 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:
961 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:
962 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:
963 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:
964 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: {
965 cgm.errorNYI(expr->getSourceRange(),
966 std::string("unimplemented AMDGPU builtin call: ") +
967 getContext().BuiltinInfo.getName(builtinId));
968 return mlir::Value{};
969 }
970 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:
971 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:
972 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:
973 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:
974 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:
975 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:
976 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:
977 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:
978 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:
979 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:
980 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:
981 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:
982 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:
983 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {
984 cgm.errorNYI(expr->getSourceRange(),
985 std::string("unimplemented AMDGPU builtin call: ") +
986 getContext().BuiltinInfo.getName(builtinId));
987 return mlir::Value{};
988 }
989 // amdgcn workgroup size
990 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:
991 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:
992 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: {
993 cgm.errorNYI(expr->getSourceRange(),
994 std::string("unimplemented AMDGPU builtin call: ") +
995 getContext().BuiltinInfo.getName(builtinId));
996 return mlir::Value{};
997 }
998 case AMDGPU::BI__builtin_amdgcn_grid_size_x:
999 case AMDGPU::BI__builtin_amdgcn_grid_size_y:
1000 case AMDGPU::BI__builtin_amdgcn_grid_size_z: {
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_r600_recipsqrt_ieee:
1007 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: {
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_alignbit: {
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_fence: {
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_atomic_inc32:
1026 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:
1027 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:
1028 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:
1029 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:
1030 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:
1031 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:
1032 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:
1033 case AMDGPU::BI__builtin_amdgcn_ds_faddf:
1034 case AMDGPU::BI__builtin_amdgcn_ds_fminf:
1035 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:
1036 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:
1037 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:
1038 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:
1039 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:
1040 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:
1041 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:
1042 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:
1043 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:
1044 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:
1045 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:
1046 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:
1047 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {
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_s_sendmsg_rtn:
1054 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {
1055 cgm.errorNYI(expr->getSourceRange(),
1056 std::string("unimplemented AMDGPU builtin call: ") +
1057 getContext().BuiltinInfo.getName(builtinId));
1058 return mlir::Value{};
1059 }
1060 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:
1061 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {
1062 cgm.errorNYI(expr->getSourceRange(),
1063 std::string("unimplemented AMDGPU builtin call: ") +
1064 getContext().BuiltinInfo.getName(builtinId));
1065 return mlir::Value{};
1066 }
1067 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
1068 case AMDGPU::BI__builtin_amdgcn_bitop3_b16:
1069 return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.bitop3")
1070 .getValue();
1071 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {
1072 cgm.errorNYI(expr->getSourceRange(),
1073 std::string("unimplemented AMDGPU builtin call: ") +
1074 getContext().BuiltinInfo.getName(builtinId));
1075 return mlir::Value{};
1076 }
1077 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:
1078 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:
1079 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:
1080 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:
1081 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:
1082 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: {
1083 cgm.errorNYI(expr->getSourceRange(),
1084 std::string("unimplemented AMDGPU builtin call: ") +
1085 getContext().BuiltinInfo.getName(builtinId));
1086 return mlir::Value{};
1087 }
1088 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:
1089 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:
1090 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:
1091 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:
1092 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:
1093 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {
1094 cgm.errorNYI(expr->getSourceRange(),
1095 std::string("unimplemented AMDGPU builtin call: ") +
1096 getContext().BuiltinInfo.getName(builtinId));
1097 return mlir::Value{};
1098 }
1099 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: {
1100 cgm.errorNYI(expr->getSourceRange(),
1101 std::string("unimplemented AMDGPU builtin call: ") +
1102 getContext().BuiltinInfo.getName(builtinId));
1103 return mlir::Value{};
1104 }
1105 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:
1106 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64:
1107 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: {
1108 cgm.errorNYI(expr->getSourceRange(),
1109 std::string("unimplemented AMDGPU builtin call: ") +
1110 getContext().BuiltinInfo.getName(builtinId));
1111 return mlir::Value{};
1112 }
1113 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:
1114 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: {
1115 cgm.errorNYI(expr->getSourceRange(),
1116 std::string("unimplemented AMDGPU builtin call: ") +
1117 getContext().BuiltinInfo.getName(builtinId));
1118 return mlir::Value{};
1119 }
1120 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:
1121 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: {
1122 cgm.errorNYI(expr->getSourceRange(),
1123 std::string("unimplemented AMDGPU builtin call: ") +
1124 getContext().BuiltinInfo.getName(builtinId));
1125 return mlir::Value{};
1126 }
1127 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
1128 return emitBuiltinWithOneOverloadedType<2>(
1129 expr, "amdgcn.s.prefetch.data",
1130 cir::VoidType::get(builder.getContext()))
1131 .getValue();
1132 case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst:
1133 return emitBuiltinWithOneOverloadedType<2>(
1134 expr, "amdgcn.s.prefetch.inst",
1135 cir::VoidType::get(builder.getContext()))
1136 .getValue();
1137 case Builtin::BIlogbf:
1138 case Builtin::BI__builtin_logbf:
1139 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle());
1140 case Builtin::BIlogb:
1141 case Builtin::BI__builtin_logb:
1142 return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble());
1143 case Builtin::BIscalbnf:
1144 case Builtin::BI__builtin_scalbnf:
1145 case Builtin::BIscalbn:
1146 case Builtin::BI__builtin_scalbn: {
1147 return emitBinaryExpMaybeConstrainedFPBuiltin(
1148 *this, expr, "ldexp", "experimental.constrained.ldexp");
1149 }
1150 default:
1151 return std::nullopt;
1152 }
1153}
emitAMDGPUDispatchPtr
static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, const CallExpr *e=nullptr)
Definition CIRGenBuiltinAMDGPU.cpp:28
emitLogbBuiltin
static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, const llvm::fltSemantics &fSem)
Definition CIRGenBuiltinAMDGPU.cpp:71
emitAMDGCNImageOverloadedReturnType
static mlir::Value emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, bool isImageStore)
Definition CIRGenBuiltinAMDGPU.cpp:125
emitBinaryExpMaybeConstrainedFPBuiltin
static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, llvm::StringRef constrainedIntrinsicName)
Definition CIRGenBuiltinAMDGPU.cpp:51
TargetBuiltins.h
Enumerates target-specific builtins in their own namespaces within namespace clang.
cir::CIRBaseBuilderTy::createNSWAdd
mlir::Value createNSWAdd(mlir::Location loc, mlir::Value lhs, mlir::Value rhs)
Definition CIRBaseBuilder.h:831
cir::CIRBaseBuilderTy::getConstant
cir::ConstantOp getConstant(mlir::Location loc, mlir::TypedAttr attr)
Definition CIRBaseBuilder.h:83
cir::CIRBaseBuilderTy::createCompare
cir::CmpOp createCompare(mlir::Location loc, cir::CmpOpKind kind, mlir::Value lhs, mlir::Value rhs)
Definition CIRBaseBuilder.h:901
cir::CIRBaseBuilderTy::createSelect
mlir::Value createSelect(mlir::Location loc, mlir::Value condition, mlir::Value trueValue, mlir::Value falseValue)
Definition CIRBaseBuilder.h:769
cir::CIRBaseBuilderTy::getIsFPConstrained
bool getIsFPConstrained() const
Query for the use of constrained floating point math.
Definition CIRBaseBuilder.h:231
cir::CIRBaseBuilderTy::createAddrSpaceCast
mlir::Value createAddrSpaceCast(mlir::Location loc, mlir::Value src, mlir::Type newTy)
Definition CIRBaseBuilder.h:718
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::ASTContext::GetBuiltinType
QualType GetBuiltinType(unsigned ID, GetBuiltinTypeError &Error, unsigned *IntegerConstantArgs=nullptr) const
Return the type for the specified builtin.
Definition ASTContext.cpp:13087
clang::ASTContext::GE_None
@ GE_None
No error.
Definition ASTContext.h:2746
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:847
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:398
clang::CIRGen::CIRGenFunction::convertType
mlir::Type convertType(clang::QualType t)
Definition CIRGenFunction.cpp:123
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:125
clang::CIRGen::CIRGenFunction::getTarget
const TargetInfo & getTarget() const
Definition CIRGenFunction.h:404
clang::CIRGen::CIRGenFunction::getLoc
mlir::Location getLoc(clang::SourceLocation srcLoc)
Helpers to convert Clang's SourceLocation to a MLIR Location.
Definition CIRGenFunction.cpp:127
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:1956
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:145
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:1884
clang::CIRGen::CIRGenFunction::getBuilder
CIRGenBuilderTy & getBuilder()
Definition CIRGenFunction.h:316
clang::CIRGen::CIRGenFunction::getIntrinsicType
cir::FuncType getIntrinsicType(llvm::Intrinsic::ID id, llvm::ArrayRef< mlir::Type > overloadTys={})
Return the CIR signature of the LLVM intrinsic id, resolving its overloaded types to overloadTys.
Definition CIRGenBuiltin.cpp:1237
clang::CIRGen::CIRGenFunction::getContext
clang::ASTContext & getContext() const
Definition CIRGenFunction.h:314
clang::CIRGen::CIRGenFunction::emitScalarOrConstFoldImmArg
mlir::Value emitScalarOrConstFoldImmArg(unsigned iceArguments, unsigned idx, const Expr *argExpr)
Definition CIRGenBuiltin.cpp:3434
clang::CIRGen::CIRGenModule::errorNYI
DiagnosticBuilder errorNYI(SourceLocation, llvm::StringRef)
Helpers to emit "not yet implemented" error diagnostics.
Definition CIRGenModule.cpp:4168
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