| //===---- CIRGenBuiltinAMDGPU.cpp - Emit CIR for AMDGPU builtins ----------===// | |
| // | |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | |
| // See https://llvm.org/LICENSE.txt for license information. | |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | |
| // | |
| //===----------------------------------------------------------------------===// | |
| // | |
| // This contains code to emit AMDGPU Builtin calls. | |
| // | |
| //===----------------------------------------------------------------------===// | |
| #include "CIRGenFunction.h" | |
| #include "mlir/IR/Value.h" | |
| #include "clang/Basic/TargetBuiltins.h" | |
| #include "llvm/Support/AMDGPUAddrSpace.h" | |
| #include "llvm/Support/ErrorHandling.h" | |
| using namespace clang; | |
| using namespace clang::CIRGen; | |
| using namespace cir; | |
| // Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the | |
| // result to match \p e's return type when needed. | |
| // If \p e is null, returns the raw AS-4 pointer. | |
| static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf, | |
| const CallExpr *e = nullptr) { | |
| CIRGenBuilderTy &builder = cgf.getBuilder(); | |
| mlir::Location loc = | |
| e ? cgf.getLoc(e->getExprLoc()) : builder.getUnknownLoc(); | |
| // The intrinsic always returns a pointer in the constant AS. | |
| mlir::Type retTy = cir::PointerType::get( | |
| cir::VoidType::get(builder.getContext()), | |
| cir::TargetAddressSpaceAttr::get(builder.getContext(), | |
| llvm::AMDGPUAS::CONSTANT_ADDRESS)); | |
| mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr", | |
| retTy, mlir::ValueRange{}); | |
| if (!e) | |
| return call; | |
| // Only cast when the caller-visible AS differs from the intrinsic's AS; | |
| auto expectedPtrTy = | |
| mlir::cast<cir::PointerType>(cgf.convertType(e->getType())); | |
| auto callPtrTy = mlir::cast<cir::PointerType>(call.getType()); | |
| if (expectedPtrTy.getAddrSpace() == callPtrTy.getAddrSpace()) | |
| return call; | |
| return builder.createAddrSpaceCast(loc, call, expectedPtrTy); | |
| } | |
| static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin( | |
| CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName, | |
| llvm::StringRef constrainedIntrinsicName) { | |
| mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0)); | |
| mlir::Value src1 = cgf.emitScalarExpr(e->getArg(1)); | |
| mlir::Location loc = cgf.getLoc(e->getExprLoc()); | |
| CIRGenBuilderTy &builder = cgf.getBuilder(); | |
| CIRGenFunction::CIRGenFPOptionsRAII fpOptsRAII(cgf, e); | |
| if (builder.getIsFPConstrained()) { | |
| cgf.cgm.errorNYI(e->getSourceRange(), | |
| "constrained FP intrinsic support is NYI."); | |
| } | |
| return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(), | |
| mlir::ValueRange{src0, src1}); | |
| } | |
| static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, | |
| const llvm::fltSemantics &fSem) { | |
| CIRGenBuilderTy &builder = cgf.getBuilder(); | |
| mlir::Location loc = cgf.getLoc(e->getExprLoc()); | |
| mlir::Value src0 = cgf.emitScalarExpr(e->getArg(0)); | |
| mlir::Type srcTy = src0.getType(); | |
| mlir::Type int32Ty = builder.getSInt32Ty(); | |
| mlir::Type frExpResMembers[] = {srcTy, int32Ty}; | |
| cir::RecordType frExpResTy = builder.getAnonRecordTy( | |
| frExpResMembers, /*packed=*/false, | |
| cir::RecordType::getAllDataKinds(frExpResMembers)); | |
| mlir::Value frExpResult = builder.emitIntrinsicCallOp( | |
| loc, "frexp", frExpResTy, mlir::ValueRange{src0}); | |
| mlir::Value exp = | |
| cir::ExtractMemberOp::create(builder, loc, int32Ty, frExpResult, 1); | |
| mlir::Value negativeOne = | |
| builder.getConstant(loc, cir::IntAttr::get(int32Ty, -1)); | |
| mlir::Value expMinus1 = builder.createAdd(loc, exp, negativeOne); | |
| mlir::Value siToFp = cir::CastOp::create( | |
| builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1); | |
| mlir::Value fabs = cir::FAbsOp::create(builder, loc, src0); | |
| llvm::APFloat infVal = llvm::APFloat::getInf(fSem); | |
| mlir::Value inf = builder.getConstant(loc, cir::FPAttr::get(srcTy, infVal)); | |
| mlir::Value fabsNegInf = | |
| builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf); | |
| mlir::Value sel = builder.createSelect(loc, fabsNegInf, siToFp, fabs); | |
| llvm::APFloat zeroValue = llvm::APFloat::getZero(fSem); | |
| mlir::Value zero = | |
| builder.getConstant(loc, cir::FPAttr::get(srcTy, zeroValue)); | |
| mlir::Value srcEqZero = | |
| builder.createCompare(loc, cir::CmpOpKind::eq, src0, zero); | |
| llvm::APFloat negInfVal = llvm::APFloat::getInf(fSem, true); | |
| mlir::Value negInf = | |
| builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal)); | |
| mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel); | |
| return res; | |
| } | |
| static mlir::Value | |
| emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e, | |
| llvm::StringRef intrinsicName, | |
| bool isImageStore) { | |
| auto &builder = cgf.getBuilder(); | |
| llvm::SmallVector<mlir::Value, 10> args; | |
| for (unsigned i = 0, n = e->getNumArgs(); i < n; ++i) | |
| args.push_back(cgf.emitScalarExpr(e->getArg(i))); | |
| mlir::Type retTy = isImageStore ? cir::VoidType::get(builder.getContext()) | |
| : cgf.convertType(e->getType()); | |
| auto callOp = cir::LLVMIntrinsicCallOp::create( | |
| builder, cgf.getLoc(e->getExprLoc()), | |
| builder.getStringAttr(intrinsicName), retTy, args); | |
| return callOp.getResult(); | |
| } | |
| std::optional<mlir::Value> | |
| CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, | |
| const CallExpr *expr) { | |
| switch (builtinId) { | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64: | |
| case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_div_scale: | |
| case AMDGPU::BI__builtin_amdgcn_div_scalef: { | |
| Address flagOutPtr = emitPointerWithAlignment(expr->getArg(3)); | |
| llvm::StringRef intrinsicName = "amdgcn.div.scale"; | |
| mlir::Value x = emitScalarExpr(expr->getArg(0)); | |
| mlir::Value y = emitScalarExpr(expr->getArg(1)); | |
| mlir::Value z = emitScalarExpr(expr->getArg(2)); | |
| auto i1Ty = builder.getUIntNTy(1); | |
| mlir::Type resMembers[] = {x.getType(), i1Ty}; | |
| cir::RecordType resTy = | |
| builder.getAnonRecordTy(resMembers, /*packed=*/false, | |
| cir::RecordType::getAllDataKinds(resMembers)); | |
| mlir::Value structResult = | |
| cir::LLVMIntrinsicCallOp::create(builder, getLoc(expr->getExprLoc()), | |
| builder.getStringAttr(intrinsicName), | |
| resTy, {x, y, z}) | |
| .getResult(); | |
| mlir::Value result = cir::ExtractMemberOp::create( | |
| builder, getLoc(expr->getExprLoc()), x.getType(), structResult, 0); | |
| mlir::Value flag = cir::ExtractMemberOp::create( | |
| builder, getLoc(expr->getExprLoc()), i1Ty, structResult, 1); | |
| mlir::Type flagType = flagOutPtr.getElementType(); | |
| mlir::Value flagToStore = | |
| cir::CastOp::create(builder, getLoc(expr->getExprLoc()), flagType, | |
| cir::CastKind::int_to_bool, flag); | |
| builder.createStore(getLoc(expr->getExprLoc()), flagToStore, flagOutPtr); | |
| return result; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_div_fmas: | |
| case AMDGPU::BI__builtin_amdgcn_div_fmasf: | |
| return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.div.fmas") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_mov_dpp8: | |
| case AMDGPU::BI__builtin_amdgcn_mov_dpp: | |
| case AMDGPU::BI__builtin_amdgcn_update_dpp: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_permlane16: | |
| case AMDGPU::BI__builtin_amdgcn_permlanex16: { | |
| llvm::StringRef intrinsicName = | |
| builtinId == AMDGPU::BI__builtin_amdgcn_permlane16 | |
| ? "amdgcn.permlane16" | |
| : "amdgcn.permlanex16"; | |
| return emitBuiltinWithOneOverloadedType<6>(expr, intrinsicName).getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_permlane64: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.permlane64") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_permlane_bcast: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.bcast") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_permlane_up: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.up") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_permlane_down: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.down") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_permlane_xor: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.permlane.xor") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_readlane: | |
| return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.readlane") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_readfirstlane: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.readfirstlane") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_wave_shuffle: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_div_fixup: | |
| case AMDGPU::BI__builtin_amdgcn_div_fixupf: | |
| case AMDGPU::BI__builtin_amdgcn_div_fixuph: { | |
| mlir::Value src0 = emitScalarExpr(expr->getArg(0)); | |
| mlir::Value src1 = emitScalarExpr(expr->getArg(1)); | |
| mlir::Value src2 = emitScalarExpr(expr->getArg(2)); | |
| return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), | |
| "amdgcn.div.fixup", src0.getType(), | |
| mlir::ValueRange{src0, src1, src2}); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_trig_preop: | |
| case AMDGPU::BI__builtin_amdgcn_trig_preopf: | |
| return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.trig.preop") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_rcp: | |
| case AMDGPU::BI__builtin_amdgcn_rcpf: | |
| case AMDGPU::BI__builtin_amdgcn_rcph: | |
| case AMDGPU::BI__builtin_amdgcn_rcp_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_sqrt: | |
| case AMDGPU::BI__builtin_amdgcn_sqrtf: | |
| case AMDGPU::BI__builtin_amdgcn_sqrth: | |
| case AMDGPU::BI__builtin_amdgcn_sqrt_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sqrt").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_rsq: | |
| case AMDGPU::BI__builtin_amdgcn_rsqf: | |
| case AMDGPU::BI__builtin_amdgcn_rsqh: | |
| case AMDGPU::BI__builtin_amdgcn_rsq_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_rsq_clamp: | |
| case AMDGPU::BI__builtin_amdgcn_rsq_clampf: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp") | |
| .getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_sinf: | |
| case AMDGPU::BI__builtin_amdgcn_sinh: | |
| case AMDGPU::BI__builtin_amdgcn_sin_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_cosf: | |
| case AMDGPU::BI__builtin_amdgcn_cosh: | |
| case AMDGPU::BI__builtin_amdgcn_cos_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: | |
| return emitAMDGPUDispatchPtr(*this, expr); | |
| case AMDGPU::BI__builtin_amdgcn_logf: | |
| case AMDGPU::BI__builtin_amdgcn_log_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.log").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_exp2f: | |
| case AMDGPU::BI__builtin_amdgcn_exp2_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.exp2").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_log_clampf: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_ldexp: | |
| case AMDGPU::BI__builtin_amdgcn_ldexpf: | |
| case AMDGPU::BI__builtin_amdgcn_ldexph: { | |
| mlir::Value src0 = emitScalarExpr(expr->getArg(0)); | |
| mlir::Value src1 = emitScalarExpr(expr->getArg(1)); | |
| // For ldexph, the raw instruction has different behavior for out-of-bounds | |
| // exponent values (implicit truncation instead of saturate to | |
| // short_min/short_max), so truncate the exponent to i16 first. | |
| mlir::Value exp = | |
| builtinId == AMDGPU::BI__builtin_amdgcn_ldexph | |
| ? cir::CastOp::create(builder, getLoc(expr->getExprLoc()), | |
| builder.getSInt16Ty(), | |
| cir::CastKind::integral, src1) | |
| : src1; | |
| return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp", | |
| src0.getType(), | |
| mlir::ValueRange{src0, exp}); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_frexp_mant: | |
| case AMDGPU::BI__builtin_amdgcn_frexp_mantf: | |
| case AMDGPU::BI__builtin_amdgcn_frexp_manth: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant") | |
| .getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_frexp_exp: | |
| case AMDGPU::BI__builtin_amdgcn_frexp_expf: | |
| case AMDGPU::BI__builtin_amdgcn_frexp_exph: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_fract: | |
| case AMDGPU::BI__builtin_amdgcn_fractf: | |
| case AMDGPU::BI__builtin_amdgcn_fracth: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.fract").getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_lerp: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_ubfe: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.ubfe").getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_sbfe: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.sbfe").getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ballot_w32: | |
| case AMDGPU::BI__builtin_amdgcn_ballot_w64: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32: | |
| case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.inverse.ballot", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_tanhf: | |
| case AMDGPU::BI__builtin_amdgcn_tanhh: | |
| case AMDGPU::BI__builtin_amdgcn_tanh_bf16: { | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue(); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_uicmp: | |
| case AMDGPU::BI__builtin_amdgcn_uicmpl: | |
| case AMDGPU::BI__builtin_amdgcn_sicmp: | |
| case AMDGPU::BI__builtin_amdgcn_sicmpl: | |
| case AMDGPU::BI__builtin_amdgcn_fcmp: | |
| case AMDGPU::BI__builtin_amdgcn_fcmpf: { | |
| mlir::Value lhs = emitScalarExpr(expr->getArg(0)); | |
| mlir::Value rhs = emitScalarExpr(expr->getArg(1)); | |
| uint64_t imm = | |
| expr->getArg(2)->EvaluateKnownConstInt(getContext()).getZExtValue(); | |
| cir::CmpOpKind pred; | |
| switch (imm) { | |
| case 0x1: // FCMP_OEQ | |
| case 0x20: // ICMP_EQ | |
| pred = cir::CmpOpKind::eq; | |
| break; | |
| case 0xe: // FCMP_UNE | |
| case 0x21: // ICMP_NE | |
| pred = cir::CmpOpKind::ne; | |
| break; | |
| case 0x2: // FCMP_OGT | |
| case 0x22: // ICMP_UGT | |
| case 0x26: // ICMP_SGT | |
| pred = cir::CmpOpKind::gt; | |
| break; | |
| case 0x3: // FCMP_OGE | |
| case 0x23: // ICMP_UGE | |
| case 0x27: // ICMP_SGE | |
| pred = cir::CmpOpKind::ge; | |
| break; | |
| case 0x4: // FCMP_OLT | |
| case 0x24: // ICMP_ULT | |
| case 0x28: // ICMP_SLT | |
| pred = cir::CmpOpKind::lt; | |
| break; | |
| case 0x5: // FCMP_OLE | |
| case 0x25: // ICMP_ULE | |
| case 0x29: // ICMP_SLE | |
| pred = cir::CmpOpKind::le; | |
| break; | |
| case 0x6: // FCMP_ONE | |
| pred = cir::CmpOpKind::one; | |
| break; | |
| case 0x8: // FCMP_UNO | |
| pred = cir::CmpOpKind::uno; | |
| break; | |
| default: | |
| cgm.errorNYI(expr->getSourceRange(), | |
| "amdgcn compare with unsupported predicate"); | |
| return mlir::Value{}; | |
| } | |
| mlir::Location loc = getLoc(expr->getExprLoc()); | |
| mlir::Value cmp = builder.createCompare(loc, pred, lhs, rhs); | |
| return builder.emitIntrinsicCallOp(loc, "amdgcn.ballot", | |
| convertType(expr->getType()), cmp); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_class: | |
| case AMDGPU::BI__builtin_amdgcn_classf: | |
| case AMDGPU::BI__builtin_amdgcn_classh: | |
| return emitBuiltinWithOneOverloadedType<2>(expr, "amdgcn.class", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_fmed3f: | |
| case AMDGPU::BI__builtin_amdgcn_fmed3h: | |
| return emitBuiltinWithOneOverloadedType<3>(expr, "amdgcn.fmed3").getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_append: | |
| case AMDGPU::BI__builtin_amdgcn_ds_consume: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>( | |
| expr, "amdgcn.global.load.tr.b64", convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16: | |
| return emitBuiltinWithOneOverloadedType<1>( | |
| expr, "amdgcn.global.load.tr.b128", convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>( | |
| expr, "amdgcn.global.load.tr4.b64", convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: | |
| return emitBuiltinWithOneOverloadedType<1>( | |
| expr, "amdgcn.global.load.tr6.b96", convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr4.b64", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr6.b96", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr8.b64", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.load.tr16.b128", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr4.b64", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr8.b64", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr6.b96", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16: | |
| return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ds.read.tr16.b64", | |
| convertType(expr->getType())) | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64: | |
| case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128: | |
| case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32: | |
| case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64: | |
| case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_cluster_load_b32: | |
| case AMDGPU::BI__builtin_amdgcn_cluster_load_b64: | |
| case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_load_to_lds: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B: | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B: | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B: | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B: | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B: | |
| case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_get_fpenv: | |
| case AMDGPU::BI__builtin_amdgcn_set_fpenv: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_read_exec: | |
| case AMDGPU::BI__builtin_amdgcn_read_exec_lo: | |
| case AMDGPU::BI__builtin_amdgcn_read_exec_hi: { | |
| // The exec mask is read as a ballot over an all-true predicate. The | |
| // ballot is at least as wide as the wavefront, so that a wave64 target | |
| // still reports both halves for the _lo/_hi forms. | |
| mlir::Location loc = getLoc(expr->getExprLoc()); | |
| unsigned registerWidth = | |
| builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_lo ? 32 : 64; | |
| unsigned ballotWidth = | |
| std::max(getTarget().getGridValue().GV_Warp_Size, registerWidth); | |
| cir::IntType ballotTy = builder.getUIntNTy(ballotWidth); | |
| mlir::Value truePred = builder.getBool(true, loc).getResult(); | |
| mlir::Value result = | |
| builder.emitIntrinsicCallOp(loc, "amdgcn.ballot", ballotTy, truePred); | |
| if (builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_hi) | |
| result = builder.createShiftRight(loc, result, 32); | |
| return builder.createIntCast(result, convertType(expr->getType())); | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray: | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h: | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l: | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray: | |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn: | |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn: | |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn: | |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.load.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.load.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.load.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.load.cube", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.load.mip.cube", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.store.1d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.1darray", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.store.2d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.2darray", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.store.3d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.store.cube", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.1d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.1darray", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.2d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.2darray", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.3d", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32: | |
| case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.store.mip.cube", true); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.sample.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.sample.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType(*this, expr, | |
| "amdgcn.image.sample.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.cube", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.d.1d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.d.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.d.3d", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.cube", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.cube", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.d.1darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.lz.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.l.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.sample.d.2darray", false); | |
| case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: | |
| case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32: | |
| return emitAMDGCNImageOverloadedReturnType( | |
| *this, expr, "amdgcn.image.gather4.lz.2d", false); | |
| case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4: | |
| case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4: | |
| case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8: | |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| // amdgcn workgroup size | |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_x: | |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_y: | |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_grid_size_x: | |
| case AMDGPU::BI__builtin_amdgcn_grid_size_y: | |
| case AMDGPU::BI__builtin_amdgcn_grid_size_z: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_r600_recipsqrt_ieee: | |
| case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_alignbit: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_fence: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_atomic_inc32: | |
| case AMDGPU::BI__builtin_amdgcn_atomic_inc64: | |
| case AMDGPU::BI__builtin_amdgcn_atomic_dec32: | |
| case AMDGPU::BI__builtin_amdgcn_atomic_dec64: | |
| case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64: | |
| case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: | |
| case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16: | |
| case AMDGPU::BI__builtin_amdgcn_ds_faddf: | |
| case AMDGPU::BI__builtin_amdgcn_ds_fminf: | |
| case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64: | |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64: | |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn: | |
| case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_permlane16_swap: | |
| case AMDGPU::BI__builtin_amdgcn_permlane32_swap: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_bitop3_b32: | |
| case AMDGPU::BI__builtin_amdgcn_bitop3_b16: | |
| return emitBuiltinWithOneOverloadedType<4>(expr, "amdgcn.bitop3") | |
| .getValue(); | |
| case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96: | |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32: | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f64: | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32: | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32: | |
| case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case AMDGPU::BI__builtin_amdgcn_s_prefetch_inst: { | |
| cgm.errorNYI(expr->getSourceRange(), | |
| std::string("unimplemented AMDGPU builtin call: ") + | |
| getContext().BuiltinInfo.getName(builtinId)); | |
| return mlir::Value{}; | |
| } | |
| case Builtin::BIlogbf: | |
| case Builtin::BI__builtin_logbf: | |
| return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEsingle()); | |
| case Builtin::BIlogb: | |
| case Builtin::BI__builtin_logb: | |
| return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble()); | |
| case Builtin::BIscalbnf: | |
| case Builtin::BI__builtin_scalbnf: | |
| case Builtin::BIscalbn: | |
| case Builtin::BI__builtin_scalbn: { | |
| return emitBinaryExpMaybeConstrainedFPBuiltin( | |
| *this, expr, "ldexp", "experimental.constrained.ldexp"); | |
| } | |
| default: | |
| return std::nullopt; | |
| } | |
| } |