| //===------- AMDCPU.cpp - Emit LLVM Code for 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 Builtin calls as LLVM code. |
| // |
| //===----------------------------------------------------------------------===// |
| |
| #include "CGBuiltin.h" |
| #include "clang/Basic/TargetBuiltins.h" |
| #include "llvm/Analysis/ValueTracking.h" |
| #include "llvm/IR/IntrinsicsAMDGPU.h" |
| #include "llvm/IR/IntrinsicsR600.h" |
| #include "llvm/IR/MemoryModelRelaxationAnnotations.h" |
| #include "llvm/Support/AMDGPUAddrSpace.h" |
| |
| using namespace clang; |
| using namespace CodeGen; |
| using namespace llvm; |
| |
| namespace { |
| |
| // Has second type mangled argument. |
| static Value * |
| emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E, |
| Intrinsic::ID IntrinsicID, |
| Intrinsic::ID ConstrainedIntrinsicID) { |
| llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1)); |
| |
| CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, E); |
| if (CGF.Builder.getIsFPConstrained()) { |
| Function *F = CGF.CGM.getIntrinsic(ConstrainedIntrinsicID, |
| {Src0->getType(), Src1->getType()}); |
| return CGF.Builder.CreateConstrainedFPCall(F, {Src0, Src1}); |
| } |
| |
| Function *F = |
| CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()}); |
| return CGF.Builder.CreateCall(F, {Src0, Src1}); |
| } |
| |
| // If \p E is not null pointer, insert address space cast to match return |
| // type of \p E if necessary. |
| Value *EmitAMDGPUDispatchPtr(CodeGenFunction &CGF, |
| const CallExpr *E = nullptr) { |
| auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_dispatch_ptr); |
| auto *Call = CGF.Builder.CreateCall(F); |
| Call->addRetAttr( |
| Attribute::getWithDereferenceableBytes(Call->getContext(), 64)); |
| Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(4))); |
| if (!E) |
| return Call; |
| QualType BuiltinRetType = E->getType(); |
| auto *RetTy = cast<llvm::PointerType>(CGF.ConvertType(BuiltinRetType)); |
| if (RetTy == Call->getType()) |
| return Call; |
| return CGF.Builder.CreateAddrSpaceCast(Call, RetTy); |
| } |
| |
| Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) { |
| auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_implicitarg_ptr); |
| auto *Call = CGF.Builder.CreateCall(F); |
| Call->addRetAttr( |
| Attribute::getWithDereferenceableBytes(Call->getContext(), 256)); |
| Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(8))); |
| return Call; |
| } |
| |
| // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively. |
| /// Emit code based on Code Object ABI version. |
| /// COV_4 : Emit code to use dispatch ptr |
| /// COV_5+ : Emit code to use implicitarg ptr |
| /// COV_NONE : Emit code to load a global variable "__oclc_ABI_version" |
| /// and use its value for COV_4 or COV_5+ approach. It is used for |
| /// compiling device libraries in an ABI-agnostic way. |
| Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) { |
| llvm::LoadInst *LD; |
| |
| auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion; |
| |
| if (Cov == CodeObjectVersionKind::COV_None) { |
| StringRef Name = "__oclc_ABI_version"; |
| auto *ABIVersionC = CGF.CGM.getModule().getNamedGlobal(Name); |
| if (!ABIVersionC) |
| ABIVersionC = new llvm::GlobalVariable( |
| CGF.CGM.getModule(), CGF.Int32Ty, false, |
| llvm::GlobalValue::ExternalLinkage, nullptr, Name, nullptr, |
| llvm::GlobalVariable::NotThreadLocal, |
| CGF.CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant)); |
| |
| // This load will be eliminated by the IPSCCP because it is constant |
| // weak_odr without externally_initialized. Either changing it to weak or |
| // adding externally_initialized will keep the load. |
| Value *ABIVersion = CGF.Builder.CreateAlignedLoad(CGF.Int32Ty, ABIVersionC, |
| CGF.CGM.getIntAlign()); |
| |
| Value *IsCOV5 = CGF.Builder.CreateICmpSGE( |
| ABIVersion, |
| llvm::ConstantInt::get(CGF.Int32Ty, CodeObjectVersionKind::COV_5)); |
| |
| // Indexing the implicit kernarg segment. |
| Value *ImplicitGEP = CGF.Builder.CreateConstGEP1_32( |
| CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2); |
| |
| // Indexing the HSA kernel_dispatch_packet struct. |
| Value *DispatchGEP = CGF.Builder.CreateConstGEP1_32( |
| CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2); |
| |
| auto Result = CGF.Builder.CreateSelect(IsCOV5, ImplicitGEP, DispatchGEP); |
| LD = CGF.Builder.CreateLoad( |
| Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2))); |
| } else { |
| Value *GEP = nullptr; |
| if (Cov >= CodeObjectVersionKind::COV_5) { |
| // Indexing the implicit kernarg segment. |
| GEP = CGF.Builder.CreateConstGEP1_32( |
| CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2); |
| } else { |
| // Indexing the HSA kernel_dispatch_packet struct. |
| GEP = CGF.Builder.CreateConstGEP1_32( |
| CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2); |
| } |
| LD = CGF.Builder.CreateLoad( |
| Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2))); |
| } |
| |
| llvm::MDBuilder MDHelper(CGF.getLLVMContext()); |
| llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1), |
| APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1)); |
| LD->setMetadata(llvm::LLVMContext::MD_range, RNode); |
| LD->setMetadata(llvm::LLVMContext::MD_noundef, |
| llvm::MDNode::get(CGF.getLLVMContext(), {})); |
| LD->setMetadata(llvm::LLVMContext::MD_invariant_load, |
| llvm::MDNode::get(CGF.getLLVMContext(), {})); |
| return LD; |
| } |
| |
| // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively. |
| Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) { |
| const unsigned XOffset = 12; |
| auto *DP = EmitAMDGPUDispatchPtr(CGF); |
| // Indexing the HSA kernel_dispatch_packet struct. |
| auto *Offset = llvm::ConstantInt::get(CGF.Int32Ty, XOffset + Index * 4); |
| auto *GEP = CGF.Builder.CreateGEP(CGF.Int8Ty, DP, Offset); |
| auto *LD = CGF.Builder.CreateLoad( |
| Address(GEP, CGF.Int32Ty, CharUnits::fromQuantity(4))); |
| |
| llvm::MDBuilder MDB(CGF.getLLVMContext()); |
| |
| // Known non-zero. |
| LD->setMetadata(llvm::LLVMContext::MD_range, |
| MDB.createRange(APInt(32, 1), APInt::getZero(32))); |
| LD->setMetadata(llvm::LLVMContext::MD_invariant_load, |
| llvm::MDNode::get(CGF.getLLVMContext(), {})); |
| return LD; |
| } |
| } // namespace |
| |
| // Generates the IR for __builtin_read_exec_*. |
| // Lowers the builtin to amdgcn_ballot intrinsic. |
| static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E, |
| llvm::Type *RegisterType, |
| llvm::Type *ValueType, bool isExecHi) { |
| CodeGen::CGBuilderTy &Builder = CGF.Builder; |
| CodeGen::CodeGenModule &CGM = CGF.CGM; |
| |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {RegisterType}); |
| llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)}); |
| |
| if (isExecHi) { |
| Value *Rt2 = Builder.CreateLShr(Call, 32); |
| Rt2 = Builder.CreateTrunc(Rt2, CGF.Int32Ty); |
| return Rt2; |
| } |
| |
| return Call; |
| } |
| |
| // Emit an intrinsic that has 1 float or double operand, and 1 integer. |
| static Value *emitFPIntBuiltin(CodeGenFunction &CGF, |
| const CallExpr *E, |
| unsigned IntrinsicID) { |
| llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1)); |
| |
| Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); |
| return CGF.Builder.CreateCall(F, {Src0, Src1}); |
| } |
| |
| // For processing memory ordering and memory scope arguments of various |
| // amdgcn builtins. |
| // \p Order takes a C++11 comptabile memory-ordering specifier and converts |
| // it into LLVM's memory ordering specifier using atomic C ABI, and writes |
| // to \p AO. \p Scope takes a const char * and converts it into AMDGCN |
| // specific SyncScopeID and writes it to \p SSID. |
| void CodeGenFunction::ProcessOrderScopeAMDGCN(Value *Order, Value *Scope, |
| llvm::AtomicOrdering &AO, |
| llvm::SyncScope::ID &SSID) { |
| int ord = cast<llvm::ConstantInt>(Order)->getZExtValue(); |
| |
| // Map C11/C++11 memory ordering to LLVM memory ordering |
| assert(llvm::isValidAtomicOrderingCABI(ord)); |
| switch (static_cast<llvm::AtomicOrderingCABI>(ord)) { |
| case llvm::AtomicOrderingCABI::acquire: |
| case llvm::AtomicOrderingCABI::consume: |
| AO = llvm::AtomicOrdering::Acquire; |
| break; |
| case llvm::AtomicOrderingCABI::release: |
| AO = llvm::AtomicOrdering::Release; |
| break; |
| case llvm::AtomicOrderingCABI::acq_rel: |
| AO = llvm::AtomicOrdering::AcquireRelease; |
| break; |
| case llvm::AtomicOrderingCABI::seq_cst: |
| AO = llvm::AtomicOrdering::SequentiallyConsistent; |
| break; |
| case llvm::AtomicOrderingCABI::relaxed: |
| AO = llvm::AtomicOrdering::Monotonic; |
| break; |
| } |
| |
| // Some of the atomic builtins take the scope as a string name. |
| StringRef scp; |
| if (llvm::getConstantStringInfo(Scope, scp)) { |
| SSID = getLLVMContext().getOrInsertSyncScopeID(scp); |
| return; |
| } |
| |
| // Older builtins had an enum argument for the memory scope. |
| int scope = cast<llvm::ConstantInt>(Scope)->getZExtValue(); |
| switch (scope) { |
| case 0: // __MEMORY_SCOPE_SYSTEM |
| SSID = llvm::SyncScope::System; |
| break; |
| case 1: // __MEMORY_SCOPE_DEVICE |
| SSID = getLLVMContext().getOrInsertSyncScopeID("agent"); |
| break; |
| case 2: // __MEMORY_SCOPE_WRKGRP |
| SSID = getLLVMContext().getOrInsertSyncScopeID("workgroup"); |
| break; |
| case 3: // __MEMORY_SCOPE_WVFRNT |
| SSID = getLLVMContext().getOrInsertSyncScopeID("wavefront"); |
| break; |
| case 4: // __MEMORY_SCOPE_SINGLE |
| SSID = llvm::SyncScope::SingleThread; |
| break; |
| default: |
| SSID = llvm::SyncScope::System; |
| break; |
| } |
| } |
| |
| llvm::Value *CodeGenFunction::EmitScalarOrConstFoldImmArg(unsigned ICEArguments, |
| unsigned Idx, |
| const CallExpr *E) { |
| llvm::Value *Arg = nullptr; |
| if ((ICEArguments & (1 << Idx)) == 0) { |
| Arg = EmitScalarExpr(E->getArg(Idx)); |
| } else { |
| // If this is required to be a constant, constant fold it so that we |
| // know that the generated intrinsic gets a ConstantInt. |
| std::optional<llvm::APSInt> Result = |
| E->getArg(Idx)->getIntegerConstantExpr(getContext()); |
| assert(Result && "Expected argument to be a constant"); |
| Arg = llvm::ConstantInt::get(getLLVMContext(), *Result); |
| } |
| return Arg; |
| } |
| |
| void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst, |
| const CallExpr *E) { |
| constexpr const char *Tag = "amdgpu-as"; |
| |
| LLVMContext &Ctx = Inst->getContext(); |
| SmallVector<MMRAMetadata::TagT, 3> MMRAs; |
| for (unsigned K = 2; K < E->getNumArgs(); ++K) { |
| llvm::Value *V = EmitScalarExpr(E->getArg(K)); |
| StringRef AS; |
| if (llvm::getConstantStringInfo(V, AS)) { |
| MMRAs.push_back({Tag, AS}); |
| // TODO: Delete the resulting unused constant? |
| continue; |
| } |
| CGM.Error(E->getExprLoc(), |
| "expected an address space name as a string literal"); |
| } |
| |
| llvm::sort(MMRAs); |
| MMRAs.erase(llvm::unique(MMRAs), MMRAs.end()); |
| Inst->setMetadata(LLVMContext::MD_mmra, MMRAMetadata::getMD(Ctx, MMRAs)); |
| } |
| |
| Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, |
| const CallExpr *E) { |
| llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent; |
| llvm::SyncScope::ID SSID; |
| switch (BuiltinID) { |
| case AMDGPU::BI__builtin_amdgcn_div_scale: |
| case AMDGPU::BI__builtin_amdgcn_div_scalef: { |
| // Translate from the intrinsics's struct return to the builtin's out |
| // argument. |
| |
| Address FlagOutPtr = EmitPointerWithAlignment(E->getArg(3)); |
| |
| llvm::Value *X = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Y = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *Z = EmitScalarExpr(E->getArg(2)); |
| |
| llvm::Function *Callee = CGM.getIntrinsic(Intrinsic::amdgcn_div_scale, |
| X->getType()); |
| |
| llvm::Value *Tmp = Builder.CreateCall(Callee, {X, Y, Z}); |
| |
| llvm::Value *Result = Builder.CreateExtractValue(Tmp, 0); |
| llvm::Value *Flag = Builder.CreateExtractValue(Tmp, 1); |
| |
| llvm::Type *RealFlagType = FlagOutPtr.getElementType(); |
| |
| llvm::Value *FlagExt = Builder.CreateZExt(Flag, RealFlagType); |
| Builder.CreateStore(FlagExt, FlagOutPtr); |
| return Result; |
| } |
| case AMDGPU::BI__builtin_amdgcn_div_fmas: |
| case AMDGPU::BI__builtin_amdgcn_div_fmasf: { |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); |
| llvm::Value *Src3 = EmitScalarExpr(E->getArg(3)); |
| |
| llvm::Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_div_fmas, |
| Src0->getType()); |
| llvm::Value *Src3ToBool = Builder.CreateIsNotNull(Src3); |
| return Builder.CreateCall(F, {Src0, Src1, Src2, Src3ToBool}); |
| } |
| |
| case AMDGPU::BI__builtin_amdgcn_ds_swizzle: |
| return emitBuiltinWithOneOverloadedType<2>(*this, E, |
| Intrinsic::amdgcn_ds_swizzle); |
| case AMDGPU::BI__builtin_amdgcn_mov_dpp8: |
| case AMDGPU::BI__builtin_amdgcn_mov_dpp: |
| case AMDGPU::BI__builtin_amdgcn_update_dpp: { |
| llvm::SmallVector<llvm::Value *, 6> Args; |
| // Find out if any arguments are required to be integer constant |
| // expressions. |
| unsigned ICEArguments = 0; |
| ASTContext::GetBuiltinTypeError Error; |
| getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments); |
| assert(Error == ASTContext::GE_None && "Should not codegen an error"); |
| llvm::Type *DataTy = ConvertType(E->getArg(0)->getType()); |
| unsigned Size = DataTy->getPrimitiveSizeInBits(); |
| llvm::Type *IntTy = |
| llvm::IntegerType::get(Builder.getContext(), std::max(Size, 32u)); |
| Function *F = |
| CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp8 |
| ? Intrinsic::amdgcn_mov_dpp8 |
| : Intrinsic::amdgcn_update_dpp, |
| IntTy); |
| assert(E->getNumArgs() == 5 || E->getNumArgs() == 6 || |
| E->getNumArgs() == 2); |
| bool InsertOld = BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp; |
| if (InsertOld) |
| Args.push_back(llvm::PoisonValue::get(IntTy)); |
| for (unsigned I = 0; I != E->getNumArgs(); ++I) { |
| llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E); |
| if (I < (BuiltinID == AMDGPU::BI__builtin_amdgcn_update_dpp ? 2u : 1u) && |
| Size < 32) { |
| if (!DataTy->isIntegerTy()) |
| V = Builder.CreateBitCast( |
| V, llvm::IntegerType::get(Builder.getContext(), Size)); |
| V = Builder.CreateZExtOrBitCast(V, IntTy); |
| } |
| llvm::Type *ExpTy = |
| F->getFunctionType()->getFunctionParamType(I + InsertOld); |
| Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy)); |
| } |
| Value *V = Builder.CreateCall(F, Args); |
| if (Size < 32 && !DataTy->isIntegerTy()) |
| V = Builder.CreateTrunc( |
| V, llvm::IntegerType::get(Builder.getContext(), Size)); |
| return Builder.CreateTruncOrBitCast(V, DataTy); |
| } |
| case AMDGPU::BI__builtin_amdgcn_permlane16: |
| case AMDGPU::BI__builtin_amdgcn_permlanex16: |
| return emitBuiltinWithOneOverloadedType<6>( |
| *this, E, |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16 |
| ? Intrinsic::amdgcn_permlane16 |
| : Intrinsic::amdgcn_permlanex16); |
| case AMDGPU::BI__builtin_amdgcn_permlane64: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_permlane64); |
| case AMDGPU::BI__builtin_amdgcn_readlane: |
| return emitBuiltinWithOneOverloadedType<2>(*this, E, |
| Intrinsic::amdgcn_readlane); |
| case AMDGPU::BI__builtin_amdgcn_readfirstlane: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_readfirstlane); |
| case AMDGPU::BI__builtin_amdgcn_div_fixup: |
| case AMDGPU::BI__builtin_amdgcn_div_fixupf: |
| case AMDGPU::BI__builtin_amdgcn_div_fixuph: |
| return emitBuiltinWithOneOverloadedType<3>(*this, E, |
| Intrinsic::amdgcn_div_fixup); |
| case AMDGPU::BI__builtin_amdgcn_trig_preop: |
| case AMDGPU::BI__builtin_amdgcn_trig_preopf: |
| return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop); |
| 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>(*this, E, Intrinsic::amdgcn_rcp); |
| 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>(*this, E, |
| Intrinsic::amdgcn_sqrt); |
| 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>(*this, E, Intrinsic::amdgcn_rsq); |
| case AMDGPU::BI__builtin_amdgcn_rsq_clamp: |
| case AMDGPU::BI__builtin_amdgcn_rsq_clampf: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_rsq_clamp); |
| case AMDGPU::BI__builtin_amdgcn_sinf: |
| case AMDGPU::BI__builtin_amdgcn_sinh: |
| case AMDGPU::BI__builtin_amdgcn_sin_bf16: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_sin); |
| case AMDGPU::BI__builtin_amdgcn_cosf: |
| case AMDGPU::BI__builtin_amdgcn_cosh: |
| case AMDGPU::BI__builtin_amdgcn_cos_bf16: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_cos); |
| case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: |
| return EmitAMDGPUDispatchPtr(*this, E); |
| case AMDGPU::BI__builtin_amdgcn_logf: |
| case AMDGPU::BI__builtin_amdgcn_log_bf16: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_log); |
| case AMDGPU::BI__builtin_amdgcn_exp2f: |
| case AMDGPU::BI__builtin_amdgcn_exp2_bf16: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_exp2); |
| case AMDGPU::BI__builtin_amdgcn_log_clampf: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_log_clamp); |
| case AMDGPU::BI__builtin_amdgcn_ldexp: |
| case AMDGPU::BI__builtin_amdgcn_ldexpf: { |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Function *F = |
| CGM.getIntrinsic(Intrinsic::ldexp, {Src0->getType(), Src1->getType()}); |
| return Builder.CreateCall(F, {Src0, Src1}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_ldexph: { |
| // The raw instruction has a different behavior for out of bounds exponent |
| // values (implicit truncation instead of saturate to short_min/short_max). |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Function *F = |
| CGM.getIntrinsic(Intrinsic::ldexp, {Src0->getType(), Int16Ty}); |
| return Builder.CreateCall(F, {Src0, Builder.CreateTrunc(Src1, Int16Ty)}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_frexp_mant: |
| case AMDGPU::BI__builtin_amdgcn_frexp_mantf: |
| case AMDGPU::BI__builtin_amdgcn_frexp_manth: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_frexp_mant); |
| case AMDGPU::BI__builtin_amdgcn_frexp_exp: |
| case AMDGPU::BI__builtin_amdgcn_frexp_expf: { |
| Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp, |
| { Builder.getInt32Ty(), Src0->getType() }); |
| return Builder.CreateCall(F, Src0); |
| } |
| case AMDGPU::BI__builtin_amdgcn_frexp_exph: { |
| Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp, |
| { Builder.getInt16Ty(), Src0->getType() }); |
| return Builder.CreateCall(F, Src0); |
| } |
| case AMDGPU::BI__builtin_amdgcn_fract: |
| case AMDGPU::BI__builtin_amdgcn_fractf: |
| case AMDGPU::BI__builtin_amdgcn_fracth: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_fract); |
| case AMDGPU::BI__builtin_amdgcn_lerp: |
| return emitBuiltinWithOneOverloadedType<3>(*this, E, |
| Intrinsic::amdgcn_lerp); |
| case AMDGPU::BI__builtin_amdgcn_ubfe: |
| return emitBuiltinWithOneOverloadedType<3>(*this, E, |
| Intrinsic::amdgcn_ubfe); |
| case AMDGPU::BI__builtin_amdgcn_sbfe: |
| return emitBuiltinWithOneOverloadedType<3>(*this, E, |
| Intrinsic::amdgcn_sbfe); |
| case AMDGPU::BI__builtin_amdgcn_ballot_w32: |
| case AMDGPU::BI__builtin_amdgcn_ballot_w64: { |
| llvm::Type *ResultType = ConvertType(E->getType()); |
| llvm::Value *Src = EmitScalarExpr(E->getArg(0)); |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, { ResultType }); |
| return Builder.CreateCall(F, { Src }); |
| } |
| case AMDGPU::BI__builtin_amdgcn_tanhf: |
| case AMDGPU::BI__builtin_amdgcn_tanhh: |
| case AMDGPU::BI__builtin_amdgcn_tanh_bf16: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::amdgcn_tanh); |
| case AMDGPU::BI__builtin_amdgcn_uicmp: |
| case AMDGPU::BI__builtin_amdgcn_uicmpl: |
| case AMDGPU::BI__builtin_amdgcn_sicmp: |
| case AMDGPU::BI__builtin_amdgcn_sicmpl: { |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); |
| |
| // FIXME-GFX10: How should 32 bit mask be handled? |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp, |
| { Builder.getInt64Ty(), Src0->getType() }); |
| return Builder.CreateCall(F, { Src0, Src1, Src2 }); |
| } |
| case AMDGPU::BI__builtin_amdgcn_fcmp: |
| case AMDGPU::BI__builtin_amdgcn_fcmpf: { |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); |
| |
| // FIXME-GFX10: How should 32 bit mask be handled? |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp, |
| { Builder.getInt64Ty(), Src0->getType() }); |
| return Builder.CreateCall(F, { Src0, Src1, Src2 }); |
| } |
| case AMDGPU::BI__builtin_amdgcn_class: |
| case AMDGPU::BI__builtin_amdgcn_classf: |
| case AMDGPU::BI__builtin_amdgcn_classh: |
| return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class); |
| case AMDGPU::BI__builtin_amdgcn_fmed3f: |
| case AMDGPU::BI__builtin_amdgcn_fmed3h: |
| return emitBuiltinWithOneOverloadedType<3>(*this, E, |
| Intrinsic::amdgcn_fmed3); |
| case AMDGPU::BI__builtin_amdgcn_ds_append: |
| case AMDGPU::BI__builtin_amdgcn_ds_consume: { |
| Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ? |
| Intrinsic::amdgcn_ds_append : Intrinsic::amdgcn_ds_consume; |
| Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() }); |
| return Builder.CreateCall(F, { Src0, Builder.getFalse() }); |
| } |
| 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_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_tr4_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: |
| 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: |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: |
| 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: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16: |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: { |
| Intrinsic::ID IID; |
| switch (BuiltinID) { |
| 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: |
| IID = Intrinsic::amdgcn_global_load_tr_b64; |
| break; |
| 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: |
| IID = Intrinsic::amdgcn_global_load_tr_b128; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32: |
| IID = Intrinsic::amdgcn_global_load_tr4_b64; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32: |
| IID = Intrinsic::amdgcn_global_load_tr6_b96; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32: |
| IID = Intrinsic::amdgcn_ds_load_tr4_b64; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32: |
| IID = Intrinsic::amdgcn_ds_load_tr6_b96; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32: |
| IID = Intrinsic::amdgcn_ds_load_tr8_b64; |
| break; |
| 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: |
| IID = Intrinsic::amdgcn_ds_load_tr16_b128; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32: |
| IID = Intrinsic::amdgcn_ds_read_tr4_b64; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32: |
| IID = Intrinsic::amdgcn_ds_read_tr8_b64; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32: |
| IID = Intrinsic::amdgcn_ds_read_tr6_b96; |
| break; |
| 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: |
| IID = Intrinsic::amdgcn_ds_read_tr16_b64; |
| break; |
| } |
| llvm::Type *LoadTy = ConvertType(E->getType()); |
| llvm::Value *Addr = EmitScalarExpr(E->getArg(0)); |
| llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy}); |
| return Builder.CreateCall(F, {Addr}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_load_to_lds: { |
| // Should this have asan instrumentation? |
| return emitBuiltinWithOneOverloadedType<5>(*this, E, |
| Intrinsic::amdgcn_load_to_lds); |
| } |
| case AMDGPU::BI__builtin_amdgcn_get_fpenv: { |
| Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv, |
| {llvm::Type::getInt64Ty(getLLVMContext())}); |
| return Builder.CreateCall(F); |
| } |
| case AMDGPU::BI__builtin_amdgcn_set_fpenv: { |
| Function *F = CGM.getIntrinsic(Intrinsic::set_fpenv, |
| {llvm::Type::getInt64Ty(getLLVMContext())}); |
| llvm::Value *Env = EmitScalarExpr(E->getArg(0)); |
| return Builder.CreateCall(F, {Env}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_read_exec: |
| return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false); |
| case AMDGPU::BI__builtin_amdgcn_read_exec_lo: |
| return EmitAMDGCNBallotForExec(*this, E, Int32Ty, Int32Ty, false); |
| case AMDGPU::BI__builtin_amdgcn_read_exec_hi: |
| return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, true); |
| 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: { |
| llvm::Value *NodePtr = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *RayExtent = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *RayOrigin = EmitScalarExpr(E->getArg(2)); |
| llvm::Value *RayDir = EmitScalarExpr(E->getArg(3)); |
| llvm::Value *RayInverseDir = EmitScalarExpr(E->getArg(4)); |
| llvm::Value *TextureDescr = EmitScalarExpr(E->getArg(5)); |
| |
| // The builtins take these arguments as vec4 where the last element is |
| // ignored. The intrinsic takes them as vec3. |
| RayOrigin = Builder.CreateShuffleVector(RayOrigin, RayOrigin, |
| {0, 1, 2}); |
| RayDir = |
| Builder.CreateShuffleVector(RayDir, RayDir, {0, 1, 2}); |
| RayInverseDir = Builder.CreateShuffleVector(RayInverseDir, RayInverseDir, |
| {0, 1, 2}); |
| |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_image_bvh_intersect_ray, |
| {NodePtr->getType(), RayDir->getType()}); |
| return Builder.CreateCall(F, {NodePtr, RayExtent, RayOrigin, RayDir, |
| RayInverseDir, TextureDescr}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray: |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: { |
| Intrinsic::ID IID; |
| switch (BuiltinID) { |
| case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray: |
| IID = Intrinsic::amdgcn_image_bvh8_intersect_ray; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: |
| IID = Intrinsic::amdgcn_image_bvh_dual_intersect_ray; |
| break; |
| } |
| llvm::Value *NodePtr = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *RayExtent = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *InstanceMask = EmitScalarExpr(E->getArg(2)); |
| llvm::Value *RayOrigin = EmitScalarExpr(E->getArg(3)); |
| llvm::Value *RayDir = EmitScalarExpr(E->getArg(4)); |
| llvm::Value *Offset = EmitScalarExpr(E->getArg(5)); |
| llvm::Value *TextureDescr = EmitScalarExpr(E->getArg(6)); |
| |
| Address RetRayOriginPtr = EmitPointerWithAlignment(E->getArg(7)); |
| Address RetRayDirPtr = EmitPointerWithAlignment(E->getArg(8)); |
| |
| llvm::Function *IntrinsicFunc = CGM.getIntrinsic(IID); |
| |
| llvm::CallInst *CI = Builder.CreateCall( |
| IntrinsicFunc, {NodePtr, RayExtent, InstanceMask, RayOrigin, RayDir, |
| Offset, TextureDescr}); |
| |
| llvm::Value *RetVData = Builder.CreateExtractValue(CI, 0); |
| llvm::Value *RetRayOrigin = Builder.CreateExtractValue(CI, 1); |
| llvm::Value *RetRayDir = Builder.CreateExtractValue(CI, 2); |
| |
| Builder.CreateStore(RetRayOrigin, RetRayOriginPtr); |
| Builder.CreateStore(RetRayDir, RetRayDirPtr); |
| |
| return RetVData; |
| } |
| |
| 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: { |
| Intrinsic::ID IID; |
| switch (BuiltinID) { |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn: |
| IID = Intrinsic::amdgcn_ds_bvh_stack_rtn; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn: |
| IID = Intrinsic::amdgcn_ds_bvh_stack_push4_pop1_rtn; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn: |
| IID = Intrinsic::amdgcn_ds_bvh_stack_push8_pop1_rtn; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: |
| IID = Intrinsic::amdgcn_ds_bvh_stack_push8_pop2_rtn; |
| break; |
| } |
| |
| SmallVector<Value *, 4> Args; |
| for (int i = 0, e = E->getNumArgs(); i != e; ++i) |
| Args.push_back(EmitScalarExpr(E->getArg(i))); |
| |
| Function *F = CGM.getIntrinsic(IID); |
| Value *Call = Builder.CreateCall(F, Args); |
| Value *Rtn = Builder.CreateExtractValue(Call, 0); |
| Value *A = Builder.CreateExtractValue(Call, 1); |
| llvm::Type *RetTy = ConvertType(E->getType()); |
| Value *I0 = Builder.CreateInsertElement(PoisonValue::get(RetTy), Rtn, |
| (uint64_t)0); |
| // ds_bvh_stack_push8_pop2_rtn returns {i64, i32} but the builtin returns |
| // <2 x i64>, zext the second value. |
| if (A->getType()->getPrimitiveSizeInBits() < |
| RetTy->getScalarType()->getPrimitiveSizeInBits()) |
| A = Builder.CreateZExt(A, RetTy->getScalarType()); |
| |
| return Builder.CreateInsertElement(I0, A, 1); |
| } |
| case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4: |
| case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: { |
| llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8); |
| Function *F = CGM.getIntrinsic( |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4 |
| ? Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4 |
| : Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4, |
| {VT, VT}); |
| |
| SmallVector<Value *, 9> Args; |
| for (unsigned I = 0, N = E->getNumArgs(); I != N; ++I) |
| Args.push_back(EmitScalarExpr(E->getArg(I))); |
| return Builder.CreateCall(F, Args); |
| } |
| 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: |
| 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: |
| // GFX1250 WMMA builtins |
| 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_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: { |
| |
| // These operations perform a matrix multiplication and accumulation of |
| // the form: |
| // D = A * B + C |
| // We need to specify one type for matrices AB and one for matrices CD. |
| // Sparse matrix operations can have different types for A and B as well as |
| // an additional type for sparsity index. |
| // Destination type should be put before types used for source operands. |
| SmallVector<unsigned, 2> ArgsForMatchingMatrixTypes; |
| // On GFX12, the intrinsics with 16-bit accumulator use a packed layout. |
| // There is no need for the variable opsel argument, so always set it to |
| // "false". |
| bool AppendFalseForOpselArg = false; |
| unsigned BuiltinWMMAOp; |
| // Need return type when D and C are of different types. |
| bool NeedReturnType = false; |
| |
| switch (BuiltinID) { |
| 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_f32_16x16x16_f16_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_f16; |
| break; |
| 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_bf16_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12: |
| AppendFalseForOpselArg = true; |
| [[fallthrough]]; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x16_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12: |
| AppendFalseForOpselArg = true; |
| [[fallthrough]]; |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32: |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x16_f16_tied; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32: |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16_tied; |
| break; |
| 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_i32_16x16x16_iu8_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x16_iu8; |
| break; |
| 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_iu4_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x16_iu4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12: |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12: |
| ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x32_iu4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x32_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16_16x16x32_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x32_iu8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x32_iu4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x64_iu4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32: |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_bf8; |
| break; |
| // GFX1250 WMMA builtins |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32: |
| ArgsForMatchingMatrixTypes = {5, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x4_f32; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16: |
| ArgsForMatchingMatrixTypes = {5, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x32_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16: |
| ArgsForMatchingMatrixTypes = {5, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x32_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16: |
| ArgsForMatchingMatrixTypes = {5, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x32_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16: |
| ArgsForMatchingMatrixTypes = {5, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16: |
| NeedReturnType = true; |
| ArgsForMatchingMatrixTypes = {1, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {3, 0}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8: |
| ArgsForMatchingMatrixTypes = {4, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x64_iu8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4: |
| ArgsForMatchingMatrixTypes = {5, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4: |
| ArgsForMatchingMatrixTypes = {3, 0, 1}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_32x16x128_f4; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x64_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x64_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x64_f16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16_16x16x64_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16f32_16x16x64_bf16; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_fp8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_fp8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_bf8_fp8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8: |
| ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_bf8_bf8; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: |
| ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; |
| BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8; |
| break; |
| } |
| |
| SmallVector<Value *, 6> Args; |
| for (int i = 0, e = E->getNumArgs(); i != e; ++i) |
| Args.push_back(EmitScalarExpr(E->getArg(i))); |
| if (AppendFalseForOpselArg) |
| Args.push_back(Builder.getFalse()); |
| |
| SmallVector<llvm::Type *, 6> ArgTypes; |
| if (NeedReturnType) |
| ArgTypes.push_back(ConvertType(E->getType())); |
| for (auto ArgIdx : ArgsForMatchingMatrixTypes) |
| ArgTypes.push_back(Args[ArgIdx]->getType()); |
| |
| Function *F = CGM.getIntrinsic(BuiltinWMMAOp, ArgTypes); |
| return Builder.CreateCall(F, Args); |
| } |
| // amdgcn workgroup size |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_x: |
| return EmitAMDGPUWorkGroupSize(*this, 0); |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_y: |
| return EmitAMDGPUWorkGroupSize(*this, 1); |
| case AMDGPU::BI__builtin_amdgcn_workgroup_size_z: |
| return EmitAMDGPUWorkGroupSize(*this, 2); |
| |
| // amdgcn grid size |
| case AMDGPU::BI__builtin_amdgcn_grid_size_x: |
| return EmitAMDGPUGridSize(*this, 0); |
| case AMDGPU::BI__builtin_amdgcn_grid_size_y: |
| return EmitAMDGPUGridSize(*this, 1); |
| case AMDGPU::BI__builtin_amdgcn_grid_size_z: |
| return EmitAMDGPUGridSize(*this, 2); |
| |
| // r600 intrinsics |
| case AMDGPU::BI__builtin_r600_recipsqrt_ieee: |
| case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: |
| return emitBuiltinWithOneOverloadedType<1>(*this, E, |
| Intrinsic::r600_recipsqrt_ieee); |
| case AMDGPU::BI__builtin_amdgcn_alignbit: { |
| llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); |
| llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); |
| Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType()); |
| return Builder.CreateCall(F, { Src0, Src1, Src2 }); |
| } |
| case AMDGPU::BI__builtin_amdgcn_fence: { |
| ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)), |
| EmitScalarExpr(E->getArg(1)), AO, SSID); |
| FenceInst *Fence = Builder.CreateFence(AO, SSID); |
| if (E->getNumArgs() > 2) |
| AddAMDGPUFenceAddressSpaceMMRA(Fence, E); |
| return Fence; |
| } |
| 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: { |
| llvm::AtomicRMWInst::BinOp BinOp; |
| switch (BuiltinID) { |
| case AMDGPU::BI__builtin_amdgcn_atomic_inc32: |
| case AMDGPU::BI__builtin_amdgcn_atomic_inc64: |
| BinOp = llvm::AtomicRMWInst::UIncWrap; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_atomic_dec32: |
| case AMDGPU::BI__builtin_amdgcn_atomic_dec64: |
| BinOp = llvm::AtomicRMWInst::UDecWrap; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_faddf: |
| 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_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: |
| BinOp = llvm::AtomicRMWInst::FAdd; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_ds_fminf: |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64: |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64: |
| BinOp = llvm::AtomicRMWInst::FMin; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64: |
| case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: |
| case AMDGPU::BI__builtin_amdgcn_ds_fmaxf: |
| BinOp = llvm::AtomicRMWInst::FMax; |
| break; |
| } |
| |
| Address Ptr = CheckAtomicAlignment(*this, E); |
| Value *Val = EmitScalarExpr(E->getArg(1)); |
| llvm::Type *OrigTy = Val->getType(); |
| QualType PtrTy = E->getArg(0)->IgnoreImpCasts()->getType(); |
| |
| bool Volatile; |
| |
| if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_faddf || |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fminf || |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fmaxf) { |
| // __builtin_amdgcn_ds_faddf/fminf/fmaxf has an explicit volatile argument |
| Volatile = |
| cast<ConstantInt>(EmitScalarExpr(E->getArg(4)))->getZExtValue(); |
| } else { |
| // Infer volatile from the passed type. |
| Volatile = |
| PtrTy->castAs<PointerType>()->getPointeeType().isVolatileQualified(); |
| } |
| |
| if (E->getNumArgs() >= 4) { |
| // Some of the builtins have explicit ordering and scope arguments. |
| ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(2)), |
| EmitScalarExpr(E->getArg(3)), AO, SSID); |
| } else { |
| // Most of the builtins do not have syncscope/order arguments. For DS |
| // atomics the scope doesn't really matter, as they implicitly operate at |
| // workgroup scope. |
| // |
| // The global/flat cases need to use agent scope to consistently produce |
| // the native instruction instead of a cmpxchg expansion. |
| SSID = getLLVMContext().getOrInsertSyncScopeID("agent"); |
| AO = AtomicOrdering::Monotonic; |
| |
| // The v2bf16 builtin uses i16 instead of a natural bfloat type. |
| if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16 || |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16 || |
| BuiltinID == AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16) { |
| llvm::Type *V2BF16Ty = FixedVectorType::get( |
| llvm::Type::getBFloatTy(Builder.getContext()), 2); |
| Val = Builder.CreateBitCast(Val, V2BF16Ty); |
| } |
| } |
| |
| llvm::AtomicRMWInst *RMW = |
| Builder.CreateAtomicRMW(BinOp, Ptr, Val, AO, SSID); |
| if (Volatile) |
| RMW->setVolatile(true); |
| |
| unsigned AddrSpace = Ptr.getType()->getAddressSpace(); |
| if (AddrSpace != llvm::AMDGPUAS::LOCAL_ADDRESS) { |
| // Most targets require "amdgpu.no.fine.grained.memory" to emit the native |
| // instruction for flat and global operations. |
| llvm::MDTuple *EmptyMD = MDNode::get(getLLVMContext(), {}); |
| RMW->setMetadata("amdgpu.no.fine.grained.memory", EmptyMD); |
| |
| // Most targets require "amdgpu.ignore.denormal.mode" to emit the native |
| // instruction, but this only matters for float fadd. |
| if (BinOp == llvm::AtomicRMWInst::FAdd && Val->getType()->isFloatTy()) |
| RMW->setMetadata("amdgpu.ignore.denormal.mode", EmptyMD); |
| } |
| |
| return Builder.CreateBitCast(RMW, OrigTy); |
| } |
| case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn: |
| case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: { |
| llvm::Value *Arg = EmitScalarExpr(E->getArg(0)); |
| llvm::Type *ResultType = ConvertType(E->getType()); |
| // s_sendmsg_rtn is mangled using return type only. |
| Function *F = |
| CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType}); |
| return Builder.CreateCall(F, {Arg}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_permlane16_swap: |
| case AMDGPU::BI__builtin_amdgcn_permlane32_swap: { |
| // Because builtin types are limited, and the intrinsic uses a struct/pair |
| // output, marshal the pair-of-i32 to <2 x i32>. |
| Value *VDstOld = EmitScalarExpr(E->getArg(0)); |
| Value *VSrcOld = EmitScalarExpr(E->getArg(1)); |
| Value *FI = EmitScalarExpr(E->getArg(2)); |
| Value *BoundCtrl = EmitScalarExpr(E->getArg(3)); |
| Function *F = |
| CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16_swap |
| ? Intrinsic::amdgcn_permlane16_swap |
| : Intrinsic::amdgcn_permlane32_swap); |
| llvm::CallInst *Call = |
| Builder.CreateCall(F, {VDstOld, VSrcOld, FI, BoundCtrl}); |
| |
| llvm::Value *Elt0 = Builder.CreateExtractValue(Call, 0); |
| llvm::Value *Elt1 = Builder.CreateExtractValue(Call, 1); |
| |
| llvm::Type *ResultType = ConvertType(E->getType()); |
| |
| llvm::Value *Insert0 = Builder.CreateInsertElement( |
| llvm::PoisonValue::get(ResultType), Elt0, UINT64_C(0)); |
| llvm::Value *AsVector = |
| Builder.CreateInsertElement(Insert0, Elt1, UINT64_C(1)); |
| return AsVector; |
| } |
| case AMDGPU::BI__builtin_amdgcn_bitop3_b32: |
| case AMDGPU::BI__builtin_amdgcn_bitop3_b16: |
| return emitBuiltinWithOneOverloadedType<4>(*this, E, |
| Intrinsic::amdgcn_bitop3); |
| case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: { |
| // TODO: LLVM has this overloaded to allow for fat pointers, but since |
| // those haven't been plumbed through to Clang yet, default to creating the |
| // resource type. |
| SmallVector<Value *, 4> Args; |
| for (unsigned I = 0; I < 4; ++I) |
| Args.push_back(EmitScalarExpr(E->getArg(I))); |
| llvm::PointerType *RetTy = llvm::PointerType::get( |
| Builder.getContext(), llvm::AMDGPUAS::BUFFER_RESOURCE); |
| Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_make_buffer_rsrc, |
| {RetTy, Args[0]->getType()}); |
| return Builder.CreateCall(F, Args); |
| } |
| 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: |
| return emitBuiltinWithOneOverloadedType<5>( |
| *this, E, Intrinsic::amdgcn_raw_ptr_buffer_store); |
| 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: { |
| llvm::Type *RetTy = nullptr; |
| switch (BuiltinID) { |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8: |
| RetTy = Int8Ty; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16: |
| RetTy = Int16Ty; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32: |
| RetTy = Int32Ty; |
| break; |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64: |
| RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/2); |
| break; |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96: |
| RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/3); |
| break; |
| case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: |
| RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/4); |
| break; |
| } |
| Function *F = |
| CGM.getIntrinsic(Intrinsic::amdgcn_raw_ptr_buffer_load, RetTy); |
| return Builder.CreateCall( |
| F, {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)), |
| EmitScalarExpr(E->getArg(2)), EmitScalarExpr(E->getArg(3))}); |
| } |
| case AMDGPU::BI__builtin_amdgcn_s_prefetch_data: |
| return emitBuiltinWithOneOverloadedType<2>( |
| *this, E, Intrinsic::amdgcn_s_prefetch_data); |
| case Builtin::BIlogbf: |
| case Builtin::BI__builtin_logbf: { |
| Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| Function *FrExpFunc = CGM.getIntrinsic( |
| Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()}); |
| CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0); |
| Value *Exp = Builder.CreateExtractValue(FrExp, 1); |
| Value *Add = Builder.CreateAdd( |
| Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true); |
| Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getFloatTy()); |
| Value *Fabs = |
| emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs); |
| Value *FCmpONE = Builder.CreateFCmpONE( |
| Fabs, ConstantFP::getInfinity(Builder.getFloatTy())); |
| Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs); |
| Value *FCmpOEQ = |
| Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getFloatTy())); |
| Value *Sel2 = Builder.CreateSelect( |
| FCmpOEQ, |
| ConstantFP::getInfinity(Builder.getFloatTy(), /*Negative=*/true), Sel1); |
| return Sel2; |
| } |
| case Builtin::BIlogb: |
| case Builtin::BI__builtin_logb: { |
| Value *Src0 = EmitScalarExpr(E->getArg(0)); |
| Function *FrExpFunc = CGM.getIntrinsic( |
| Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()}); |
| CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0); |
| Value *Exp = Builder.CreateExtractValue(FrExp, 1); |
| Value *Add = Builder.CreateAdd( |
| Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true); |
| Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getDoubleTy()); |
| Value *Fabs = |
| emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs); |
| Value *FCmpONE = Builder.CreateFCmpONE( |
| Fabs, ConstantFP::getInfinity(Builder.getDoubleTy())); |
| Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs); |
| Value *FCmpOEQ = |
| Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getDoubleTy())); |
| Value *Sel2 = Builder.CreateSelect( |
| FCmpOEQ, |
| ConstantFP::getInfinity(Builder.getDoubleTy(), /*Negative=*/true), |
| Sel1); |
| return Sel2; |
| } |
| case Builtin::BIscalbnf: |
| case Builtin::BI__builtin_scalbnf: |
| case Builtin::BIscalbn: |
| case Builtin::BI__builtin_scalbn: |
| return emitBinaryExpMaybeConstrainedFPBuiltin( |
| *this, E, Intrinsic::ldexp, Intrinsic::experimental_constrained_ldexp); |
| default: |
| return nullptr; |
| } |
| } |
