2022 lines · cpp
1//===------- AMDCPU.cpp - Emit LLVM Code for builtins ---------------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// This contains code to emit Builtin calls as LLVM code.10//11//===----------------------------------------------------------------------===//12 13#include "CGBuiltin.h"14#include "CodeGenFunction.h"15#include "clang/Basic/DiagnosticFrontend.h"16#include "clang/Basic/SyncScope.h"17#include "clang/Basic/TargetBuiltins.h"18#include "llvm/Analysis/ValueTracking.h"19#include "llvm/CodeGen/MachineFunction.h"20#include "llvm/IR/IntrinsicsAMDGPU.h"21#include "llvm/IR/IntrinsicsR600.h"22#include "llvm/IR/MemoryModelRelaxationAnnotations.h"23#include "llvm/Support/AMDGPUAddrSpace.h"24 25using namespace clang;26using namespace CodeGen;27using namespace llvm;28 29namespace {30 31// Has second type mangled argument.32static Value *33emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E,34 Intrinsic::ID IntrinsicID,35 Intrinsic::ID ConstrainedIntrinsicID) {36 llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));37 llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));38 39 CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, E);40 if (CGF.Builder.getIsFPConstrained()) {41 Function *F = CGF.CGM.getIntrinsic(ConstrainedIntrinsicID,42 {Src0->getType(), Src1->getType()});43 return CGF.Builder.CreateConstrainedFPCall(F, {Src0, Src1});44 }45 46 Function *F =47 CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()});48 return CGF.Builder.CreateCall(F, {Src0, Src1});49}50 51// If \p E is not null pointer, insert address space cast to match return52// type of \p E if necessary.53Value *EmitAMDGPUDispatchPtr(CodeGenFunction &CGF,54 const CallExpr *E = nullptr) {55 auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_dispatch_ptr);56 auto *Call = CGF.Builder.CreateCall(F);57 Call->addRetAttr(58 Attribute::getWithDereferenceableBytes(Call->getContext(), 64));59 Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(4)));60 if (!E)61 return Call;62 QualType BuiltinRetType = E->getType();63 auto *RetTy = cast<llvm::PointerType>(CGF.ConvertType(BuiltinRetType));64 if (RetTy == Call->getType())65 return Call;66 return CGF.Builder.CreateAddrSpaceCast(Call, RetTy);67}68 69Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {70 auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_implicitarg_ptr);71 auto *Call = CGF.Builder.CreateCall(F);72 Call->addRetAttr(73 Attribute::getWithDereferenceableBytes(Call->getContext(), 256));74 Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(8)));75 return Call;76}77 78// \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.79/// Emit code based on Code Object ABI version.80/// COV_4 : Emit code to use dispatch ptr81/// COV_5+ : Emit code to use implicitarg ptr82/// COV_NONE : Emit code to load a global variable "__oclc_ABI_version"83/// and use its value for COV_4 or COV_5+ approach. It is used for84/// compiling device libraries in an ABI-agnostic way.85Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {86 llvm::LoadInst *LD;87 88 auto Cov = CGF.getTarget().getTargetOpts().CodeObjectVersion;89 90 if (Cov == CodeObjectVersionKind::COV_None) {91 StringRef Name = "__oclc_ABI_version";92 auto *ABIVersionC = CGF.CGM.getModule().getNamedGlobal(Name);93 if (!ABIVersionC)94 ABIVersionC = new llvm::GlobalVariable(95 CGF.CGM.getModule(), CGF.Int32Ty, false,96 llvm::GlobalValue::ExternalLinkage, nullptr, Name, nullptr,97 llvm::GlobalVariable::NotThreadLocal,98 CGF.CGM.getContext().getTargetAddressSpace(LangAS::opencl_constant));99 100 // This load will be eliminated by the IPSCCP because it is constant101 // weak_odr without externally_initialized. Either changing it to weak or102 // adding externally_initialized will keep the load.103 Value *ABIVersion = CGF.Builder.CreateAlignedLoad(CGF.Int32Ty, ABIVersionC,104 CGF.CGM.getIntAlign());105 106 Value *IsCOV5 = CGF.Builder.CreateICmpSGE(107 ABIVersion,108 llvm::ConstantInt::get(CGF.Int32Ty, CodeObjectVersionKind::COV_5));109 110 // Indexing the implicit kernarg segment.111 Value *ImplicitGEP = CGF.Builder.CreateConstGEP1_32(112 CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);113 114 // Indexing the HSA kernel_dispatch_packet struct.115 Value *DispatchGEP = CGF.Builder.CreateConstGEP1_32(116 CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);117 118 auto Result = CGF.Builder.CreateSelect(IsCOV5, ImplicitGEP, DispatchGEP);119 LD = CGF.Builder.CreateLoad(120 Address(Result, CGF.Int16Ty, CharUnits::fromQuantity(2)));121 } else {122 Value *GEP = nullptr;123 if (Cov >= CodeObjectVersionKind::COV_5) {124 // Indexing the implicit kernarg segment.125 GEP = CGF.Builder.CreateConstGEP1_32(126 CGF.Int8Ty, EmitAMDGPUImplicitArgPtr(CGF), 12 + Index * 2);127 } else {128 // Indexing the HSA kernel_dispatch_packet struct.129 GEP = CGF.Builder.CreateConstGEP1_32(130 CGF.Int8Ty, EmitAMDGPUDispatchPtr(CGF), 4 + Index * 2);131 }132 LD = CGF.Builder.CreateLoad(133 Address(GEP, CGF.Int16Ty, CharUnits::fromQuantity(2)));134 }135 136 llvm::MDBuilder MDHelper(CGF.getLLVMContext());137 llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),138 APInt(16, CGF.getTarget().getMaxOpenCLWorkGroupSize() + 1));139 LD->setMetadata(llvm::LLVMContext::MD_range, RNode);140 LD->setMetadata(llvm::LLVMContext::MD_noundef,141 llvm::MDNode::get(CGF.getLLVMContext(), {}));142 LD->setMetadata(llvm::LLVMContext::MD_invariant_load,143 llvm::MDNode::get(CGF.getLLVMContext(), {}));144 return LD;145}146 147// \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.148Value *EmitAMDGPUGridSize(CodeGenFunction &CGF, unsigned Index) {149 const unsigned XOffset = 12;150 auto *DP = EmitAMDGPUDispatchPtr(CGF);151 // Indexing the HSA kernel_dispatch_packet struct.152 auto *Offset = llvm::ConstantInt::get(CGF.Int32Ty, XOffset + Index * 4);153 auto *GEP = CGF.Builder.CreateGEP(CGF.Int8Ty, DP, Offset);154 auto *LD = CGF.Builder.CreateLoad(155 Address(GEP, CGF.Int32Ty, CharUnits::fromQuantity(4)));156 157 llvm::MDBuilder MDB(CGF.getLLVMContext());158 159 // Known non-zero.160 LD->setMetadata(llvm::LLVMContext::MD_range,161 MDB.createRange(APInt(32, 1), APInt::getZero(32)));162 LD->setMetadata(llvm::LLVMContext::MD_invariant_load,163 llvm::MDNode::get(CGF.getLLVMContext(), {}));164 return LD;165}166} // namespace167 168// Generates the IR for __builtin_read_exec_*.169// Lowers the builtin to amdgcn_ballot intrinsic.170static Value *EmitAMDGCNBallotForExec(CodeGenFunction &CGF, const CallExpr *E,171 llvm::Type *RegisterType,172 llvm::Type *ValueType, bool isExecHi) {173 CodeGen::CGBuilderTy &Builder = CGF.Builder;174 CodeGen::CodeGenModule &CGM = CGF.CGM;175 176 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {RegisterType});177 llvm::Value *Call = Builder.CreateCall(F, {Builder.getInt1(true)});178 179 if (isExecHi) {180 Value *Rt2 = Builder.CreateLShr(Call, 32);181 Rt2 = Builder.CreateTrunc(Rt2, CGF.Int32Ty);182 return Rt2;183 }184 185 return Call;186}187 188static llvm::Value *loadTextureDescPtorAsVec8I32(CodeGenFunction &CGF,189 llvm::Value *RsrcPtr) {190 auto &B = CGF.Builder;191 auto *VecTy = llvm::FixedVectorType::get(B.getInt32Ty(), 8);192 193 if (RsrcPtr->getType() == VecTy)194 return RsrcPtr;195 196 if (RsrcPtr->getType()->isIntegerTy(32)) {197 llvm::PointerType *VecPtrTy =198 llvm::PointerType::get(CGF.getLLVMContext(), 8);199 llvm::Value *Ptr = B.CreateIntToPtr(RsrcPtr, VecPtrTy, "tex.rsrc.from.int");200 return B.CreateAlignedLoad(VecTy, Ptr, llvm::Align(32), "tex.rsrc.val");201 }202 203 if (RsrcPtr->getType()->isPointerTy()) {204 auto *VecPtrTy = llvm::PointerType::get(205 CGF.getLLVMContext(), RsrcPtr->getType()->getPointerAddressSpace());206 llvm::Value *Typed = B.CreateBitCast(RsrcPtr, VecPtrTy, "tex.rsrc.typed");207 return B.CreateAlignedLoad(VecTy, Typed, llvm::Align(32), "tex.rsrc.val");208 }209 210 const auto &DL = CGF.CGM.getDataLayout();211 if (DL.getTypeSizeInBits(RsrcPtr->getType()) == 256)212 return B.CreateBitCast(RsrcPtr, VecTy, "tex.rsrc.val");213 214 llvm::report_fatal_error("Unexpected texture resource argument form");215}216 217llvm::CallInst *218emitAMDGCNImageOverloadedReturnType(clang::CodeGen::CodeGenFunction &CGF,219 const clang::CallExpr *E,220 unsigned IntrinsicID, bool IsImageStore) {221 auto findTextureDescIndex = [&CGF](const CallExpr *E) -> unsigned {222 QualType TexQT = CGF.getContext().AMDGPUTextureTy;223 for (unsigned I = 0, N = E->getNumArgs(); I < N; ++I) {224 QualType ArgTy = E->getArg(I)->getType();225 if (ArgTy == TexQT) {226 return I;227 }228 229 if (ArgTy.getCanonicalType() == TexQT.getCanonicalType()) {230 return I;231 }232 }233 234 return ~0U;235 };236 237 clang::SmallVector<llvm::Value *, 10> Args;238 unsigned RsrcIndex = findTextureDescIndex(E);239 240 if (RsrcIndex == ~0U) {241 llvm::report_fatal_error("Invalid argument count for image builtin");242 }243 244 for (unsigned I = 0; I < E->getNumArgs(); ++I) {245 llvm::Value *V = CGF.EmitScalarExpr(E->getArg(I));246 if (I == RsrcIndex)247 V = loadTextureDescPtorAsVec8I32(CGF, V);248 Args.push_back(V);249 }250 251 llvm::Type *RetTy = IsImageStore ? CGF.VoidTy : CGF.ConvertType(E->getType());252 llvm::CallInst *Call = CGF.Builder.CreateIntrinsic(RetTy, IntrinsicID, Args);253 return Call;254}255 256// Emit an intrinsic that has 1 float or double operand, and 1 integer.257static Value *emitFPIntBuiltin(CodeGenFunction &CGF,258 const CallExpr *E,259 unsigned IntrinsicID) {260 llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));261 llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));262 263 Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());264 return CGF.Builder.CreateCall(F, {Src0, Src1});265}266 267static inline StringRef mapScopeToSPIRV(StringRef AMDGCNScope) {268 if (AMDGCNScope == "agent")269 return "device";270 if (AMDGCNScope == "wavefront")271 return "subgroup";272 return AMDGCNScope;273}274 275// For processing memory ordering and memory scope arguments of various276// amdgcn builtins.277// \p Order takes a C++11 compatible memory-ordering specifier and converts278// it into LLVM's memory ordering specifier using atomic C ABI, and writes279// to \p AO. \p Scope takes a const char * and converts it into AMDGCN280// specific SyncScopeID and writes it to \p SSID.281void CodeGenFunction::ProcessOrderScopeAMDGCN(Value *Order, Value *Scope,282 llvm::AtomicOrdering &AO,283 llvm::SyncScope::ID &SSID) {284 int ord = cast<llvm::ConstantInt>(Order)->getZExtValue();285 286 // Map C11/C++11 memory ordering to LLVM memory ordering287 assert(llvm::isValidAtomicOrderingCABI(ord));288 switch (static_cast<llvm::AtomicOrderingCABI>(ord)) {289 case llvm::AtomicOrderingCABI::acquire:290 case llvm::AtomicOrderingCABI::consume:291 AO = llvm::AtomicOrdering::Acquire;292 break;293 case llvm::AtomicOrderingCABI::release:294 AO = llvm::AtomicOrdering::Release;295 break;296 case llvm::AtomicOrderingCABI::acq_rel:297 AO = llvm::AtomicOrdering::AcquireRelease;298 break;299 case llvm::AtomicOrderingCABI::seq_cst:300 AO = llvm::AtomicOrdering::SequentiallyConsistent;301 break;302 case llvm::AtomicOrderingCABI::relaxed:303 AO = llvm::AtomicOrdering::Monotonic;304 break;305 }306 307 // Some of the atomic builtins take the scope as a string name.308 StringRef scp;309 if (llvm::getConstantStringInfo(Scope, scp)) {310 if (getTarget().getTriple().isSPIRV())311 scp = mapScopeToSPIRV(scp);312 SSID = getLLVMContext().getOrInsertSyncScopeID(scp);313 return;314 }315 316 // Older builtins had an enum argument for the memory scope.317 const char *SSN = nullptr;318 int scope = cast<llvm::ConstantInt>(Scope)->getZExtValue();319 switch (scope) {320 case AtomicScopeGenericModel::System: // __MEMORY_SCOPE_SYSTEM321 SSID = llvm::SyncScope::System;322 break;323 case AtomicScopeGenericModel::Device: // __MEMORY_SCOPE_DEVICE324 SSN = getTarget().getTriple().isSPIRV() ? "device" : "agent";325 break;326 case AtomicScopeGenericModel::Workgroup: // __MEMORY_SCOPE_WRKGRP327 SSN = "workgroup";328 break;329 case AtomicScopeGenericModel::Cluster: // __MEMORY_SCOPE_CLUSTR330 SSN = getTarget().getTriple().isSPIRV() ? "workgroup" : "cluster";331 break;332 case AtomicScopeGenericModel::Wavefront: // __MEMORY_SCOPE_WVFRNT333 SSN = getTarget().getTriple().isSPIRV() ? "subgroup" : "wavefront";334 break;335 case AtomicScopeGenericModel::Single: // __MEMORY_SCOPE_SINGLE336 SSID = llvm::SyncScope::SingleThread;337 break;338 default:339 SSID = llvm::SyncScope::System;340 break;341 }342 if (SSN)343 SSID = getLLVMContext().getOrInsertSyncScopeID(SSN);344}345 346llvm::Value *CodeGenFunction::EmitScalarOrConstFoldImmArg(unsigned ICEArguments,347 unsigned Idx,348 const CallExpr *E) {349 llvm::Value *Arg = nullptr;350 if ((ICEArguments & (1 << Idx)) == 0) {351 Arg = EmitScalarExpr(E->getArg(Idx));352 } else {353 // If this is required to be a constant, constant fold it so that we354 // know that the generated intrinsic gets a ConstantInt.355 std::optional<llvm::APSInt> Result =356 E->getArg(Idx)->getIntegerConstantExpr(getContext());357 assert(Result && "Expected argument to be a constant");358 Arg = llvm::ConstantInt::get(getLLVMContext(), *Result);359 }360 return Arg;361}362 363void CodeGenFunction::AddAMDGPUFenceAddressSpaceMMRA(llvm::Instruction *Inst,364 const CallExpr *E) {365 constexpr const char *Tag = "amdgpu-synchronize-as";366 367 LLVMContext &Ctx = Inst->getContext();368 SmallVector<MMRAMetadata::TagT, 3> MMRAs;369 for (unsigned K = 2; K < E->getNumArgs(); ++K) {370 llvm::Value *V = EmitScalarExpr(E->getArg(K));371 StringRef AS;372 if (llvm::getConstantStringInfo(V, AS)) {373 MMRAs.push_back({Tag, AS});374 // TODO: Delete the resulting unused constant?375 continue;376 }377 CGM.Error(E->getExprLoc(),378 "expected an address space name as a string literal");379 }380 381 llvm::sort(MMRAs);382 MMRAs.erase(llvm::unique(MMRAs), MMRAs.end());383 Inst->setMetadata(LLVMContext::MD_mmra, MMRAMetadata::getMD(Ctx, MMRAs));384}385 386static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID) {387 switch (BuiltinID) {388 default:389 llvm_unreachable("Unknown BuiltinID for wave reduction");390 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:391 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:392 return Intrinsic::amdgcn_wave_reduce_add;393 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32:394 return Intrinsic::amdgcn_wave_reduce_fadd;395 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:396 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:397 return Intrinsic::amdgcn_wave_reduce_sub;398 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f32:399 return Intrinsic::amdgcn_wave_reduce_fsub;400 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:401 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:402 return Intrinsic::amdgcn_wave_reduce_min;403 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f32:404 return Intrinsic::amdgcn_wave_reduce_fmin;405 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:406 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:407 return Intrinsic::amdgcn_wave_reduce_umin;408 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:409 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:410 return Intrinsic::amdgcn_wave_reduce_max;411 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f32:412 return Intrinsic::amdgcn_wave_reduce_fmax;413 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:414 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:415 return Intrinsic::amdgcn_wave_reduce_umax;416 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:417 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:418 return Intrinsic::amdgcn_wave_reduce_and;419 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:420 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:421 return Intrinsic::amdgcn_wave_reduce_or;422 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:423 case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64:424 return Intrinsic::amdgcn_wave_reduce_xor;425 }426}427 428Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,429 const CallExpr *E) {430 llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;431 llvm::SyncScope::ID SSID;432 switch (BuiltinID) {433 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32:434 case AMDGPU::BI__builtin_amdgcn_wave_reduce_fadd_f32:435 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32:436 case AMDGPU::BI__builtin_amdgcn_wave_reduce_fsub_f32:437 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32:438 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32:439 case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmin_f32:440 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32:441 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32:442 case AMDGPU::BI__builtin_amdgcn_wave_reduce_fmax_f32:443 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b32:444 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b32:445 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b32:446 case AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64:447 case AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64:448 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64:449 case AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64:450 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64:451 case AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64:452 case AMDGPU::BI__builtin_amdgcn_wave_reduce_and_b64:453 case AMDGPU::BI__builtin_amdgcn_wave_reduce_or_b64:454 case AMDGPU::BI__builtin_amdgcn_wave_reduce_xor_b64: {455 Intrinsic::ID IID = getIntrinsicIDforWaveReduction(BuiltinID);456 llvm::Value *Value = EmitScalarExpr(E->getArg(0));457 llvm::Value *Strategy = EmitScalarExpr(E->getArg(1));458 llvm::Function *F = CGM.getIntrinsic(IID, {Value->getType()});459 return Builder.CreateCall(F, {Value, Strategy});460 }461 case AMDGPU::BI__builtin_amdgcn_div_scale:462 case AMDGPU::BI__builtin_amdgcn_div_scalef: {463 // Translate from the intrinsics's struct return to the builtin's out464 // argument.465 466 Address FlagOutPtr = EmitPointerWithAlignment(E->getArg(3));467 468 llvm::Value *X = EmitScalarExpr(E->getArg(0));469 llvm::Value *Y = EmitScalarExpr(E->getArg(1));470 llvm::Value *Z = EmitScalarExpr(E->getArg(2));471 472 llvm::Function *Callee = CGM.getIntrinsic(Intrinsic::amdgcn_div_scale,473 X->getType());474 475 llvm::Value *Tmp = Builder.CreateCall(Callee, {X, Y, Z});476 477 llvm::Value *Result = Builder.CreateExtractValue(Tmp, 0);478 llvm::Value *Flag = Builder.CreateExtractValue(Tmp, 1);479 480 llvm::Type *RealFlagType = FlagOutPtr.getElementType();481 482 llvm::Value *FlagExt = Builder.CreateZExt(Flag, RealFlagType);483 Builder.CreateStore(FlagExt, FlagOutPtr);484 return Result;485 }486 case AMDGPU::BI__builtin_amdgcn_div_fmas:487 case AMDGPU::BI__builtin_amdgcn_div_fmasf: {488 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));489 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));490 llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));491 llvm::Value *Src3 = EmitScalarExpr(E->getArg(3));492 493 llvm::Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_div_fmas,494 Src0->getType());495 llvm::Value *Src3ToBool = Builder.CreateIsNotNull(Src3);496 return Builder.CreateCall(F, {Src0, Src1, Src2, Src3ToBool});497 }498 499 case AMDGPU::BI__builtin_amdgcn_ds_swizzle:500 return emitBuiltinWithOneOverloadedType<2>(*this, E,501 Intrinsic::amdgcn_ds_swizzle);502 case AMDGPU::BI__builtin_amdgcn_mov_dpp8:503 case AMDGPU::BI__builtin_amdgcn_mov_dpp:504 case AMDGPU::BI__builtin_amdgcn_update_dpp: {505 llvm::SmallVector<llvm::Value *, 6> Args;506 // Find out if any arguments are required to be integer constant507 // expressions.508 unsigned ICEArguments = 0;509 ASTContext::GetBuiltinTypeError Error;510 getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);511 assert(Error == ASTContext::GE_None && "Should not codegen an error");512 llvm::Type *DataTy = ConvertType(E->getArg(0)->getType());513 unsigned Size = DataTy->getPrimitiveSizeInBits();514 llvm::Type *IntTy =515 llvm::IntegerType::get(Builder.getContext(), std::max(Size, 32u));516 Function *F =517 CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp8518 ? Intrinsic::amdgcn_mov_dpp8519 : Intrinsic::amdgcn_update_dpp,520 IntTy);521 assert(E->getNumArgs() == 5 || E->getNumArgs() == 6 ||522 E->getNumArgs() == 2);523 bool InsertOld = BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp;524 if (InsertOld)525 Args.push_back(llvm::PoisonValue::get(IntTy));526 for (unsigned I = 0; I != E->getNumArgs(); ++I) {527 llvm::Value *V = EmitScalarOrConstFoldImmArg(ICEArguments, I, E);528 if (I < (BuiltinID == AMDGPU::BI__builtin_amdgcn_update_dpp ? 2u : 1u) &&529 Size < 32) {530 if (!DataTy->isIntegerTy())531 V = Builder.CreateBitCast(532 V, llvm::IntegerType::get(Builder.getContext(), Size));533 V = Builder.CreateZExtOrBitCast(V, IntTy);534 }535 llvm::Type *ExpTy =536 F->getFunctionType()->getFunctionParamType(I + InsertOld);537 Args.push_back(Builder.CreateTruncOrBitCast(V, ExpTy));538 }539 Value *V = Builder.CreateCall(F, Args);540 if (Size < 32 && !DataTy->isIntegerTy())541 V = Builder.CreateTrunc(542 V, llvm::IntegerType::get(Builder.getContext(), Size));543 return Builder.CreateTruncOrBitCast(V, DataTy);544 }545 case AMDGPU::BI__builtin_amdgcn_permlane16:546 case AMDGPU::BI__builtin_amdgcn_permlanex16:547 return emitBuiltinWithOneOverloadedType<6>(548 *this, E,549 BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16550 ? Intrinsic::amdgcn_permlane16551 : Intrinsic::amdgcn_permlanex16);552 case AMDGPU::BI__builtin_amdgcn_permlane64:553 return emitBuiltinWithOneOverloadedType<1>(*this, E,554 Intrinsic::amdgcn_permlane64);555 case AMDGPU::BI__builtin_amdgcn_readlane:556 return emitBuiltinWithOneOverloadedType<2>(*this, E,557 Intrinsic::amdgcn_readlane);558 case AMDGPU::BI__builtin_amdgcn_readfirstlane:559 return emitBuiltinWithOneOverloadedType<1>(*this, E,560 Intrinsic::amdgcn_readfirstlane);561 case AMDGPU::BI__builtin_amdgcn_div_fixup:562 case AMDGPU::BI__builtin_amdgcn_div_fixupf:563 case AMDGPU::BI__builtin_amdgcn_div_fixuph:564 return emitBuiltinWithOneOverloadedType<3>(*this, E,565 Intrinsic::amdgcn_div_fixup);566 case AMDGPU::BI__builtin_amdgcn_trig_preop:567 case AMDGPU::BI__builtin_amdgcn_trig_preopf:568 return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop);569 case AMDGPU::BI__builtin_amdgcn_rcp:570 case AMDGPU::BI__builtin_amdgcn_rcpf:571 case AMDGPU::BI__builtin_amdgcn_rcph:572 case AMDGPU::BI__builtin_amdgcn_rcp_bf16:573 return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_rcp);574 case AMDGPU::BI__builtin_amdgcn_sqrt:575 case AMDGPU::BI__builtin_amdgcn_sqrtf:576 case AMDGPU::BI__builtin_amdgcn_sqrth:577 case AMDGPU::BI__builtin_amdgcn_sqrt_bf16:578 return emitBuiltinWithOneOverloadedType<1>(*this, E,579 Intrinsic::amdgcn_sqrt);580 case AMDGPU::BI__builtin_amdgcn_rsq:581 case AMDGPU::BI__builtin_amdgcn_rsqf:582 case AMDGPU::BI__builtin_amdgcn_rsqh:583 case AMDGPU::BI__builtin_amdgcn_rsq_bf16:584 return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_rsq);585 case AMDGPU::BI__builtin_amdgcn_rsq_clamp:586 case AMDGPU::BI__builtin_amdgcn_rsq_clampf:587 return emitBuiltinWithOneOverloadedType<1>(*this, E,588 Intrinsic::amdgcn_rsq_clamp);589 case AMDGPU::BI__builtin_amdgcn_sinf:590 case AMDGPU::BI__builtin_amdgcn_sinh:591 case AMDGPU::BI__builtin_amdgcn_sin_bf16:592 return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_sin);593 case AMDGPU::BI__builtin_amdgcn_cosf:594 case AMDGPU::BI__builtin_amdgcn_cosh:595 case AMDGPU::BI__builtin_amdgcn_cos_bf16:596 return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_cos);597 case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:598 return EmitAMDGPUDispatchPtr(*this, E);599 case AMDGPU::BI__builtin_amdgcn_logf:600 case AMDGPU::BI__builtin_amdgcn_log_bf16:601 return emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::amdgcn_log);602 case AMDGPU::BI__builtin_amdgcn_exp2f:603 case AMDGPU::BI__builtin_amdgcn_exp2_bf16:604 return emitBuiltinWithOneOverloadedType<1>(*this, E,605 Intrinsic::amdgcn_exp2);606 case AMDGPU::BI__builtin_amdgcn_log_clampf:607 return emitBuiltinWithOneOverloadedType<1>(*this, E,608 Intrinsic::amdgcn_log_clamp);609 case AMDGPU::BI__builtin_amdgcn_ldexp:610 case AMDGPU::BI__builtin_amdgcn_ldexpf: {611 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));612 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));613 llvm::Function *F =614 CGM.getIntrinsic(Intrinsic::ldexp, {Src0->getType(), Src1->getType()});615 return Builder.CreateCall(F, {Src0, Src1});616 }617 case AMDGPU::BI__builtin_amdgcn_ldexph: {618 // The raw instruction has a different behavior for out of bounds exponent619 // values (implicit truncation instead of saturate to short_min/short_max).620 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));621 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));622 llvm::Function *F =623 CGM.getIntrinsic(Intrinsic::ldexp, {Src0->getType(), Int16Ty});624 return Builder.CreateCall(F, {Src0, Builder.CreateTrunc(Src1, Int16Ty)});625 }626 case AMDGPU::BI__builtin_amdgcn_frexp_mant:627 case AMDGPU::BI__builtin_amdgcn_frexp_mantf:628 case AMDGPU::BI__builtin_amdgcn_frexp_manth:629 return emitBuiltinWithOneOverloadedType<1>(*this, E,630 Intrinsic::amdgcn_frexp_mant);631 case AMDGPU::BI__builtin_amdgcn_frexp_exp:632 case AMDGPU::BI__builtin_amdgcn_frexp_expf: {633 Value *Src0 = EmitScalarExpr(E->getArg(0));634 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,635 { Builder.getInt32Ty(), Src0->getType() });636 return Builder.CreateCall(F, Src0);637 }638 case AMDGPU::BI__builtin_amdgcn_frexp_exph: {639 Value *Src0 = EmitScalarExpr(E->getArg(0));640 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp,641 { Builder.getInt16Ty(), Src0->getType() });642 return Builder.CreateCall(F, Src0);643 }644 case AMDGPU::BI__builtin_amdgcn_fract:645 case AMDGPU::BI__builtin_amdgcn_fractf:646 case AMDGPU::BI__builtin_amdgcn_fracth:647 return emitBuiltinWithOneOverloadedType<1>(*this, E,648 Intrinsic::amdgcn_fract);649 case AMDGPU::BI__builtin_amdgcn_lerp:650 return emitBuiltinWithOneOverloadedType<3>(*this, E,651 Intrinsic::amdgcn_lerp);652 case AMDGPU::BI__builtin_amdgcn_ubfe:653 return emitBuiltinWithOneOverloadedType<3>(*this, E,654 Intrinsic::amdgcn_ubfe);655 case AMDGPU::BI__builtin_amdgcn_sbfe:656 return emitBuiltinWithOneOverloadedType<3>(*this, E,657 Intrinsic::amdgcn_sbfe);658 case AMDGPU::BI__builtin_amdgcn_ballot_w32:659 case AMDGPU::BI__builtin_amdgcn_ballot_w64: {660 llvm::Type *ResultType = ConvertType(E->getType());661 llvm::Value *Src = EmitScalarExpr(E->getArg(0));662 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_ballot, {ResultType});663 return Builder.CreateCall(F, {Src});664 }665 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:666 case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {667 llvm::Value *Src = EmitScalarExpr(E->getArg(0));668 Function *F =669 CGM.getIntrinsic(Intrinsic::amdgcn_inverse_ballot, {Src->getType()});670 return Builder.CreateCall(F, {Src});671 }672 case AMDGPU::BI__builtin_amdgcn_tanhf:673 case AMDGPU::BI__builtin_amdgcn_tanhh:674 case AMDGPU::BI__builtin_amdgcn_tanh_bf16:675 return emitBuiltinWithOneOverloadedType<1>(*this, E,676 Intrinsic::amdgcn_tanh);677 case AMDGPU::BI__builtin_amdgcn_uicmp:678 case AMDGPU::BI__builtin_amdgcn_uicmpl:679 case AMDGPU::BI__builtin_amdgcn_sicmp:680 case AMDGPU::BI__builtin_amdgcn_sicmpl: {681 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));682 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));683 llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));684 685 // FIXME-GFX10: How should 32 bit mask be handled?686 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp,687 { Builder.getInt64Ty(), Src0->getType() });688 return Builder.CreateCall(F, { Src0, Src1, Src2 });689 }690 case AMDGPU::BI__builtin_amdgcn_fcmp:691 case AMDGPU::BI__builtin_amdgcn_fcmpf: {692 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));693 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));694 llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));695 696 // FIXME-GFX10: How should 32 bit mask be handled?697 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp,698 { Builder.getInt64Ty(), Src0->getType() });699 return Builder.CreateCall(F, { Src0, Src1, Src2 });700 }701 case AMDGPU::BI__builtin_amdgcn_class:702 case AMDGPU::BI__builtin_amdgcn_classf:703 case AMDGPU::BI__builtin_amdgcn_classh:704 return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);705 case AMDGPU::BI__builtin_amdgcn_fmed3f:706 case AMDGPU::BI__builtin_amdgcn_fmed3h:707 return emitBuiltinWithOneOverloadedType<3>(*this, E,708 Intrinsic::amdgcn_fmed3);709 case AMDGPU::BI__builtin_amdgcn_ds_append:710 case AMDGPU::BI__builtin_amdgcn_ds_consume: {711 Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?712 Intrinsic::amdgcn_ds_append : Intrinsic::amdgcn_ds_consume;713 Value *Src0 = EmitScalarExpr(E->getArg(0));714 Function *F = CGM.getIntrinsic(Intrin, { Src0->getType() });715 return Builder.CreateCall(F, { Src0, Builder.getFalse() });716 }717 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:718 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:719 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:720 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:721 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:722 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:723 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:724 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:725 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:726 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:727 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:728 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:729 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:730 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:731 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:732 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:733 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:734 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:735 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:736 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:737 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:738 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:739 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:740 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:741 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:742 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {743 Intrinsic::ID IID;744 switch (BuiltinID) {745 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:746 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:747 case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:748 IID = Intrinsic::amdgcn_global_load_tr_b64;749 break;750 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:751 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:752 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:753 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:754 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:755 case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:756 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:757 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:758 case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:759 IID = Intrinsic::amdgcn_global_load_tr_b128;760 break;761 case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:762 IID = Intrinsic::amdgcn_global_load_tr4_b64;763 break;764 case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:765 IID = Intrinsic::amdgcn_global_load_tr6_b96;766 break;767 case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:768 IID = Intrinsic::amdgcn_ds_load_tr4_b64;769 break;770 case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:771 IID = Intrinsic::amdgcn_ds_load_tr6_b96;772 break;773 case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:774 IID = Intrinsic::amdgcn_ds_load_tr8_b64;775 break;776 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:777 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:778 case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:779 IID = Intrinsic::amdgcn_ds_load_tr16_b128;780 break;781 case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:782 IID = Intrinsic::amdgcn_ds_read_tr4_b64;783 break;784 case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:785 IID = Intrinsic::amdgcn_ds_read_tr8_b64;786 break;787 case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:788 IID = Intrinsic::amdgcn_ds_read_tr6_b96;789 break;790 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16:791 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:792 case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:793 IID = Intrinsic::amdgcn_ds_read_tr16_b64;794 break;795 }796 llvm::Type *LoadTy = ConvertType(E->getType());797 llvm::Value *Addr = EmitScalarExpr(E->getArg(0));798 llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy});799 return Builder.CreateCall(F, {Addr});800 }801 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:802 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:803 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:804 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:805 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:806 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128: {807 808 Intrinsic::ID IID;809 switch (BuiltinID) {810 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b32:811 IID = Intrinsic::amdgcn_global_load_monitor_b32;812 break;813 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:814 IID = Intrinsic::amdgcn_global_load_monitor_b64;815 break;816 case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:817 IID = Intrinsic::amdgcn_global_load_monitor_b128;818 break;819 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b32:820 IID = Intrinsic::amdgcn_flat_load_monitor_b32;821 break;822 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b64:823 IID = Intrinsic::amdgcn_flat_load_monitor_b64;824 break;825 case AMDGPU::BI__builtin_amdgcn_flat_load_monitor_b128:826 IID = Intrinsic::amdgcn_flat_load_monitor_b128;827 break;828 }829 830 llvm::Type *LoadTy = ConvertType(E->getType());831 llvm::Value *Addr = EmitScalarExpr(E->getArg(0));832 llvm::Value *Val = EmitScalarExpr(E->getArg(1));833 llvm::Function *F = CGM.getIntrinsic(IID, {LoadTy});834 return Builder.CreateCall(F, {Addr, Val});835 }836 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:837 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:838 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128: {839 Intrinsic::ID IID;840 switch (BuiltinID) {841 case AMDGPU::BI__builtin_amdgcn_cluster_load_b32:842 IID = Intrinsic::amdgcn_cluster_load_b32;843 break;844 case AMDGPU::BI__builtin_amdgcn_cluster_load_b64:845 IID = Intrinsic::amdgcn_cluster_load_b64;846 break;847 case AMDGPU::BI__builtin_amdgcn_cluster_load_b128:848 IID = Intrinsic::amdgcn_cluster_load_b128;849 break;850 }851 SmallVector<Value *, 3> Args;852 for (int i = 0, e = E->getNumArgs(); i != e; ++i)853 Args.push_back(EmitScalarExpr(E->getArg(i)));854 llvm::Function *F = CGM.getIntrinsic(IID, {ConvertType(E->getType())});855 return Builder.CreateCall(F, {Args});856 }857 case AMDGPU::BI__builtin_amdgcn_load_to_lds: {858 // Should this have asan instrumentation?859 return emitBuiltinWithOneOverloadedType<5>(*this, E,860 Intrinsic::amdgcn_load_to_lds);861 }862 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:863 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:864 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:865 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:866 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:867 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B: {868 Intrinsic::ID IID;869 switch (BuiltinID) {870 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_32x4B:871 IID = Intrinsic::amdgcn_cooperative_atomic_load_32x4B;872 break;873 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_32x4B:874 IID = Intrinsic::amdgcn_cooperative_atomic_store_32x4B;875 break;876 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_16x8B:877 IID = Intrinsic::amdgcn_cooperative_atomic_load_16x8B;878 break;879 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:880 IID = Intrinsic::amdgcn_cooperative_atomic_store_16x8B;881 break;882 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_load_8x16B:883 IID = Intrinsic::amdgcn_cooperative_atomic_load_8x16B;884 break;885 case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:886 IID = Intrinsic::amdgcn_cooperative_atomic_store_8x16B;887 break;888 }889 890 LLVMContext &Ctx = CGM.getLLVMContext();891 SmallVector<Value *, 5> Args;892 // last argument is a MD string893 const unsigned ScopeArg = E->getNumArgs() - 1;894 for (unsigned i = 0; i != ScopeArg; ++i)895 Args.push_back(EmitScalarExpr(E->getArg(i)));896 StringRef Arg = cast<StringLiteral>(E->getArg(ScopeArg)->IgnoreParenCasts())897 ->getString();898 llvm::MDNode *MD = llvm::MDNode::get(Ctx, {llvm::MDString::get(Ctx, Arg)});899 Args.push_back(llvm::MetadataAsValue::get(Ctx, MD));900 // Intrinsic is typed based on the pointer AS. Pointer is always the first901 // argument.902 llvm::Function *F = CGM.getIntrinsic(IID, {Args[0]->getType()});903 return Builder.CreateCall(F, {Args});904 }905 case AMDGPU::BI__builtin_amdgcn_get_fpenv: {906 Function *F = CGM.getIntrinsic(Intrinsic::get_fpenv,907 {llvm::Type::getInt64Ty(getLLVMContext())});908 return Builder.CreateCall(F);909 }910 case AMDGPU::BI__builtin_amdgcn_set_fpenv: {911 Function *F = CGM.getIntrinsic(Intrinsic::set_fpenv,912 {llvm::Type::getInt64Ty(getLLVMContext())});913 llvm::Value *Env = EmitScalarExpr(E->getArg(0));914 return Builder.CreateCall(F, {Env});915 }916 case AMDGPU::BI__builtin_amdgcn_read_exec:917 return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, false);918 case AMDGPU::BI__builtin_amdgcn_read_exec_lo:919 return EmitAMDGCNBallotForExec(*this, E, Int32Ty, Int32Ty, false);920 case AMDGPU::BI__builtin_amdgcn_read_exec_hi:921 return EmitAMDGCNBallotForExec(*this, E, Int64Ty, Int64Ty, true);922 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:923 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:924 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:925 case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {926 llvm::Value *NodePtr = EmitScalarExpr(E->getArg(0));927 llvm::Value *RayExtent = EmitScalarExpr(E->getArg(1));928 llvm::Value *RayOrigin = EmitScalarExpr(E->getArg(2));929 llvm::Value *RayDir = EmitScalarExpr(E->getArg(3));930 llvm::Value *RayInverseDir = EmitScalarExpr(E->getArg(4));931 llvm::Value *TextureDescr = EmitScalarExpr(E->getArg(5));932 933 // The builtins take these arguments as vec4 where the last element is934 // ignored. The intrinsic takes them as vec3.935 RayOrigin = Builder.CreateShuffleVector(RayOrigin, RayOrigin,936 {0, 1, 2});937 RayDir =938 Builder.CreateShuffleVector(RayDir, RayDir, {0, 1, 2});939 RayInverseDir = Builder.CreateShuffleVector(RayInverseDir, RayInverseDir,940 {0, 1, 2});941 942 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_image_bvh_intersect_ray,943 {NodePtr->getType(), RayDir->getType()});944 return Builder.CreateCall(F, {NodePtr, RayExtent, RayOrigin, RayDir,945 RayInverseDir, TextureDescr});946 }947 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:948 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {949 Intrinsic::ID IID;950 switch (BuiltinID) {951 case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:952 IID = Intrinsic::amdgcn_image_bvh8_intersect_ray;953 break;954 case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray:955 IID = Intrinsic::amdgcn_image_bvh_dual_intersect_ray;956 break;957 }958 llvm::Value *NodePtr = EmitScalarExpr(E->getArg(0));959 llvm::Value *RayExtent = EmitScalarExpr(E->getArg(1));960 llvm::Value *InstanceMask = EmitScalarExpr(E->getArg(2));961 llvm::Value *RayOrigin = EmitScalarExpr(E->getArg(3));962 llvm::Value *RayDir = EmitScalarExpr(E->getArg(4));963 llvm::Value *Offset = EmitScalarExpr(E->getArg(5));964 llvm::Value *TextureDescr = EmitScalarExpr(E->getArg(6));965 966 Address RetRayOriginPtr = EmitPointerWithAlignment(E->getArg(7));967 Address RetRayDirPtr = EmitPointerWithAlignment(E->getArg(8));968 969 llvm::Function *IntrinsicFunc = CGM.getIntrinsic(IID);970 971 llvm::CallInst *CI = Builder.CreateCall(972 IntrinsicFunc, {NodePtr, RayExtent, InstanceMask, RayOrigin, RayDir,973 Offset, TextureDescr});974 975 llvm::Value *RetVData = Builder.CreateExtractValue(CI, 0);976 llvm::Value *RetRayOrigin = Builder.CreateExtractValue(CI, 1);977 llvm::Value *RetRayDir = Builder.CreateExtractValue(CI, 2);978 979 Builder.CreateStore(RetRayOrigin, RetRayOriginPtr);980 Builder.CreateStore(RetRayDir, RetRayDirPtr);981 982 return RetVData;983 }984 985 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:986 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:987 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:988 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {989 Intrinsic::ID IID;990 switch (BuiltinID) {991 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:992 IID = Intrinsic::amdgcn_ds_bvh_stack_rtn;993 break;994 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:995 IID = Intrinsic::amdgcn_ds_bvh_stack_push4_pop1_rtn;996 break;997 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:998 IID = Intrinsic::amdgcn_ds_bvh_stack_push8_pop1_rtn;999 break;1000 case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn:1001 IID = Intrinsic::amdgcn_ds_bvh_stack_push8_pop2_rtn;1002 break;1003 }1004 1005 SmallVector<Value *, 4> Args;1006 for (int i = 0, e = E->getNumArgs(); i != e; ++i)1007 Args.push_back(EmitScalarExpr(E->getArg(i)));1008 1009 Function *F = CGM.getIntrinsic(IID);1010 Value *Call = Builder.CreateCall(F, Args);1011 Value *Rtn = Builder.CreateExtractValue(Call, 0);1012 Value *A = Builder.CreateExtractValue(Call, 1);1013 llvm::Type *RetTy = ConvertType(E->getType());1014 Value *I0 = Builder.CreateInsertElement(PoisonValue::get(RetTy), Rtn,1015 (uint64_t)0);1016 // ds_bvh_stack_push8_pop2_rtn returns {i64, i32} but the builtin returns1017 // <2 x i64>, zext the second value.1018 if (A->getType()->getPrimitiveSizeInBits() <1019 RetTy->getScalarType()->getPrimitiveSizeInBits())1020 A = Builder.CreateZExt(A, RetTy->getScalarType());1021 1022 return Builder.CreateInsertElement(I0, A, 1);1023 }1024 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:1025 case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:1026 return emitAMDGCNImageOverloadedReturnType(1027 *this, E, Intrinsic::amdgcn_image_load_1d, false);1028 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:1029 case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:1030 return emitAMDGCNImageOverloadedReturnType(1031 *this, E, Intrinsic::amdgcn_image_load_1darray, false);1032 case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:1033 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:1034 case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:1035 return emitAMDGCNImageOverloadedReturnType(1036 *this, E, Intrinsic::amdgcn_image_load_2d, false);1037 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:1038 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:1039 case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:1040 return emitAMDGCNImageOverloadedReturnType(1041 *this, E, Intrinsic::amdgcn_image_load_2darray, false);1042 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:1043 case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:1044 return emitAMDGCNImageOverloadedReturnType(1045 *this, E, Intrinsic::amdgcn_image_load_3d, false);1046 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:1047 case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:1048 return emitAMDGCNImageOverloadedReturnType(1049 *this, E, Intrinsic::amdgcn_image_load_cube, false);1050 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:1051 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:1052 return emitAMDGCNImageOverloadedReturnType(1053 *this, E, Intrinsic::amdgcn_image_load_mip_1d, false);1054 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:1055 case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:1056 return emitAMDGCNImageOverloadedReturnType(1057 *this, E, Intrinsic::amdgcn_image_load_mip_1darray, false);1058 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:1059 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:1060 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:1061 return emitAMDGCNImageOverloadedReturnType(1062 *this, E, Intrinsic::amdgcn_image_load_mip_2d, false);1063 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:1064 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:1065 case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:1066 return emitAMDGCNImageOverloadedReturnType(1067 *this, E, Intrinsic::amdgcn_image_load_mip_2darray, false);1068 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:1069 case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:1070 return emitAMDGCNImageOverloadedReturnType(1071 *this, E, Intrinsic::amdgcn_image_load_mip_3d, false);1072 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:1073 case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:1074 return emitAMDGCNImageOverloadedReturnType(1075 *this, E, Intrinsic::amdgcn_image_load_mip_cube, false);1076 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:1077 case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:1078 return emitAMDGCNImageOverloadedReturnType(1079 *this, E, Intrinsic::amdgcn_image_store_1d, true);1080 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:1081 case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:1082 return emitAMDGCNImageOverloadedReturnType(1083 *this, E, Intrinsic::amdgcn_image_store_1darray, true);1084 case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:1085 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:1086 case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:1087 return emitAMDGCNImageOverloadedReturnType(1088 *this, E, Intrinsic::amdgcn_image_store_2d, true);1089 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:1090 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:1091 case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:1092 return emitAMDGCNImageOverloadedReturnType(1093 *this, E, Intrinsic::amdgcn_image_store_2darray, true);1094 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:1095 case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:1096 return emitAMDGCNImageOverloadedReturnType(1097 *this, E, Intrinsic::amdgcn_image_store_3d, true);1098 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:1099 case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:1100 return emitAMDGCNImageOverloadedReturnType(1101 *this, E, Intrinsic::amdgcn_image_store_cube, true);1102 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:1103 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:1104 return emitAMDGCNImageOverloadedReturnType(1105 *this, E, Intrinsic::amdgcn_image_store_mip_1d, true);1106 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:1107 case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:1108 return emitAMDGCNImageOverloadedReturnType(1109 *this, E, Intrinsic::amdgcn_image_store_mip_1darray, true);1110 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:1111 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:1112 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:1113 return emitAMDGCNImageOverloadedReturnType(1114 *this, E, Intrinsic::amdgcn_image_store_mip_2d, true);1115 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:1116 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:1117 case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:1118 return emitAMDGCNImageOverloadedReturnType(1119 *this, E, Intrinsic::amdgcn_image_store_mip_2darray, true);1120 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:1121 case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:1122 return emitAMDGCNImageOverloadedReturnType(1123 *this, E, Intrinsic::amdgcn_image_store_mip_3d, true);1124 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:1125 case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32:1126 return emitAMDGCNImageOverloadedReturnType(1127 *this, E, Intrinsic::amdgcn_image_store_mip_cube, true);1128 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:1129 case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:1130 return emitAMDGCNImageOverloadedReturnType(1131 *this, E, Intrinsic::amdgcn_image_sample_1d, false);1132 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:1133 case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:1134 return emitAMDGCNImageOverloadedReturnType(1135 *this, E, Intrinsic::amdgcn_image_sample_1darray, false);1136 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:1137 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:1138 case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:1139 return emitAMDGCNImageOverloadedReturnType(1140 *this, E, Intrinsic::amdgcn_image_sample_2d, false);1141 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:1142 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:1143 case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:1144 return emitAMDGCNImageOverloadedReturnType(1145 *this, E, Intrinsic::amdgcn_image_sample_2darray, false);1146 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:1147 case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:1148 return emitAMDGCNImageOverloadedReturnType(1149 *this, E, Intrinsic::amdgcn_image_sample_3d, false);1150 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:1151 case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32:1152 return emitAMDGCNImageOverloadedReturnType(1153 *this, E, Intrinsic::amdgcn_image_sample_cube, false);1154 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f32_f32:1155 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_1d_v4f16_f32:1156 return emitAMDGCNImageOverloadedReturnType(1157 *this, E, Intrinsic::amdgcn_image_sample_lz_1d, false);1158 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f32_f32:1159 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_1d_v4f16_f32:1160 return emitAMDGCNImageOverloadedReturnType(1161 *this, E, Intrinsic::amdgcn_image_sample_l_1d, false);1162 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f32_f32:1163 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_1d_v4f16_f32:1164 return emitAMDGCNImageOverloadedReturnType(1165 *this, E, Intrinsic::amdgcn_image_sample_d_1d, false);1166 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f32_f32:1167 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_v4f16_f32:1168 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2d_f32_f32:1169 return emitAMDGCNImageOverloadedReturnType(1170 *this, E, Intrinsic::amdgcn_image_sample_lz_2d, false);1171 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f32_f32:1172 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_v4f16_f32:1173 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2d_f32_f32:1174 return emitAMDGCNImageOverloadedReturnType(1175 *this, E, Intrinsic::amdgcn_image_sample_l_2d, false);1176 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f32_f32:1177 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_v4f16_f32:1178 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2d_f32_f32:1179 return emitAMDGCNImageOverloadedReturnType(1180 *this, E, Intrinsic::amdgcn_image_sample_d_2d, false);1181 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f32_f32:1182 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_3d_v4f16_f32:1183 return emitAMDGCNImageOverloadedReturnType(1184 *this, E, Intrinsic::amdgcn_image_sample_lz_3d, false);1185 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f32_f32:1186 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_3d_v4f16_f32:1187 return emitAMDGCNImageOverloadedReturnType(1188 *this, E, Intrinsic::amdgcn_image_sample_l_3d, false);1189 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:1190 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:1191 return emitAMDGCNImageOverloadedReturnType(1192 *this, E, Intrinsic::amdgcn_image_sample_d_3d, false);1193 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f32_f32:1194 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_cube_v4f16_f32:1195 return emitAMDGCNImageOverloadedReturnType(1196 *this, E, Intrinsic::amdgcn_image_sample_lz_cube, false);1197 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f32_f32:1198 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_cube_v4f16_f32:1199 return emitAMDGCNImageOverloadedReturnType(1200 *this, E, Intrinsic::amdgcn_image_sample_l_cube, false);1201 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f32_f32:1202 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_1darray_v4f16_f32:1203 return emitAMDGCNImageOverloadedReturnType(1204 *this, E, Intrinsic::amdgcn_image_sample_lz_1darray, false);1205 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f32_f32:1206 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_1darray_v4f16_f32:1207 return emitAMDGCNImageOverloadedReturnType(1208 *this, E, Intrinsic::amdgcn_image_sample_l_1darray, false);1209 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f32_f32:1210 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_1darray_v4f16_f32:1211 return emitAMDGCNImageOverloadedReturnType(1212 *this, E, Intrinsic::amdgcn_image_sample_d_1darray, false);1213 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f32_f32:1214 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_v4f16_f32:1215 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_lz_2darray_f32_f32:1216 return emitAMDGCNImageOverloadedReturnType(1217 *this, E, Intrinsic::amdgcn_image_sample_lz_2darray, false);1218 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f32_f32:1219 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_v4f16_f32:1220 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_l_2darray_f32_f32:1221 return emitAMDGCNImageOverloadedReturnType(1222 *this, E, Intrinsic::amdgcn_image_sample_l_2darray, false);1223 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f32_f32:1224 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:1225 case clang::AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_f32_f32:1226 return emitAMDGCNImageOverloadedReturnType(1227 *this, E, Intrinsic::amdgcn_image_sample_d_2darray, false);1228 case clang::AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:1229 return emitAMDGCNImageOverloadedReturnType(1230 *this, E, Intrinsic::amdgcn_image_gather4_lz_2d, false);1231 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:1232 case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {1233 llvm::FixedVectorType *VT = FixedVectorType::get(Builder.getInt32Ty(), 8);1234 Function *F = CGM.getIntrinsic(1235 BuiltinID == AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f41236 ? Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f41237 : Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4,1238 {VT, VT});1239 1240 SmallVector<Value *, 9> Args;1241 for (unsigned I = 0, N = E->getNumArgs(); I != N; ++I)1242 Args.push_back(EmitScalarExpr(E->getArg(I)));1243 return Builder.CreateCall(F, Args);1244 }1245 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:1246 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:1247 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:1248 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:1249 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:1250 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:1251 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:1252 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:1253 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:1254 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:1255 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:1256 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:1257 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:1258 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:1259 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:1260 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:1261 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:1262 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:1263 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:1264 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:1265 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:1266 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:1267 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:1268 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:1269 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:1270 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:1271 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:1272 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:1273 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:1274 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:1275 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:1276 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:1277 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:1278 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:1279 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:1280 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:1281 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:1282 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12:1283 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:1284 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:1285 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:1286 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:1287 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:1288 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:1289 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:1290 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:1291 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:1292 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:1293 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:1294 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:1295 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:1296 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:1297 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:1298 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:1299 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:1300 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:1301 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:1302 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:1303 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:1304 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64:1305 // GFX1250 WMMA builtins1306 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:1307 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:1308 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:1309 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:1310 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:1311 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:1312 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:1313 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:1314 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:1315 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:1316 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:1317 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:1318 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:1319 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:1320 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:1321 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:1322 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:1323 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:1324 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:1325 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:1326 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:1327 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:1328 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:1329 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:1330 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:1331 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:1332 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:1333 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:1334 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4:1335 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:1336 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:1337 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:1338 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:1339 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:1340 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:1341 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:1342 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:1343 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:1344 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:1345 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:1346 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:1347 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:1348 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8: {1349 1350 // These operations perform a matrix multiplication and accumulation of1351 // the form:1352 // D = A * B + C1353 // We need to specify one type for matrices AB and one for matrices CD.1354 // Sparse matrix operations can have different types for A and B as well as1355 // an additional type for sparsity index.1356 // Destination type should be put before types used for source operands.1357 SmallVector<unsigned, 2> ArgsForMatchingMatrixTypes;1358 // On GFX12, the intrinsics with 16-bit accumulator use a packed layout.1359 // There is no need for the variable opsel argument, so always set it to1360 // "false".1361 bool AppendFalseForOpselArg = false;1362 unsigned BuiltinWMMAOp;1363 // Need return type when D and C are of different types.1364 bool NeedReturnType = false;1365 1366 switch (BuiltinID) {1367 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:1368 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:1369 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:1370 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:1371 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1372 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_f16;1373 break;1374 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:1375 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:1376 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:1377 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:1378 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1379 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf16;1380 break;1381 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:1382 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:1383 AppendFalseForOpselArg = true;1384 [[fallthrough]];1385 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:1386 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:1387 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1388 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x16_f16;1389 break;1390 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:1391 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:1392 AppendFalseForOpselArg = true;1393 [[fallthrough]];1394 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:1395 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:1396 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1397 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16;1398 break;1399 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:1400 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:1401 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1402 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x16_f16_tied;1403 break;1404 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:1405 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:1406 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1407 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x16_bf16_tied;1408 break;1409 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:1410 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:1411 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:1412 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:1413 ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB1414 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x16_iu8;1415 break;1416 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:1417 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:1418 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:1419 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:1420 ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB1421 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x16_iu4;1422 break;1423 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:1424 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:1425 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1426 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_fp8_fp8;1427 break;1428 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:1429 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:1430 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1431 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_fp8_bf8;1432 break;1433 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:1434 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:1435 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1436 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf8_fp8;1437 break;1438 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:1439 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:1440 ArgsForMatchingMatrixTypes = {2, 0}; // CD, AB1441 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x16_bf8_bf8;1442 break;1443 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:1444 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12:1445 ArgsForMatchingMatrixTypes = {4, 1}; // CD, AB1446 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x32_iu4;1447 break;1448 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:1449 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:1450 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1451 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_f16;1452 break;1453 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:1454 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:1455 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1456 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf16;1457 break;1458 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:1459 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:1460 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1461 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x32_f16;1462 break;1463 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:1464 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:1465 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1466 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16_16x16x32_bf16;1467 break;1468 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:1469 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:1470 ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index1471 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x32_iu8;1472 break;1473 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:1474 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:1475 ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index1476 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x32_iu4;1477 break;1478 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:1479 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:1480 ArgsForMatchingMatrixTypes = {4, 1, 3, 5}; // CD, A, B, Index1481 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x64_iu4;1482 break;1483 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:1484 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:1485 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1486 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_fp8_fp8;1487 break;1488 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:1489 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:1490 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1491 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_fp8_bf8;1492 break;1493 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:1494 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:1495 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1496 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_fp8;1497 break;1498 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:1499 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64:1500 ArgsForMatchingMatrixTypes = {2, 0, 1, 3}; // CD, A, B, Index1501 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x32_bf8_bf8;1502 break;1503 // GFX1250 WMMA builtins1504 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x4_f32:1505 ArgsForMatchingMatrixTypes = {5, 1};1506 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x4_f32;1507 break;1508 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_bf16:1509 ArgsForMatchingMatrixTypes = {5, 1};1510 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x32_bf16;1511 break;1512 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x32_f16:1513 ArgsForMatchingMatrixTypes = {5, 1};1514 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x32_f16;1515 break;1516 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x32_f16:1517 ArgsForMatchingMatrixTypes = {5, 1};1518 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x32_f16;1519 break;1520 case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x32_bf16:1521 ArgsForMatchingMatrixTypes = {5, 1};1522 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16_16x16x32_bf16;1523 break;1524 case AMDGPU::BI__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16:1525 NeedReturnType = true;1526 ArgsForMatchingMatrixTypes = {1, 5};1527 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_bf16f32_16x16x32_bf16;1528 break;1529 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_fp8:1530 ArgsForMatchingMatrixTypes = {3, 0};1531 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_fp8_fp8;1532 break;1533 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_fp8_bf8:1534 ArgsForMatchingMatrixTypes = {3, 0};1535 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_fp8_bf8;1536 break;1537 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_fp8:1538 ArgsForMatchingMatrixTypes = {3, 0};1539 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_bf8_fp8;1540 break;1541 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x64_bf8_bf8:1542 ArgsForMatchingMatrixTypes = {3, 0};1543 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x64_bf8_bf8;1544 break;1545 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_fp8:1546 ArgsForMatchingMatrixTypes = {3, 0};1547 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_fp8_fp8;1548 break;1549 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_fp8_bf8:1550 ArgsForMatchingMatrixTypes = {3, 0};1551 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_fp8_bf8;1552 break;1553 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_fp8:1554 ArgsForMatchingMatrixTypes = {3, 0};1555 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_bf8_fp8;1556 break;1557 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x64_bf8_bf8:1558 ArgsForMatchingMatrixTypes = {3, 0};1559 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x64_bf8_bf8;1560 break;1561 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_fp8:1562 ArgsForMatchingMatrixTypes = {3, 0};1563 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_fp8_fp8;1564 break;1565 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_fp8_bf8:1566 ArgsForMatchingMatrixTypes = {3, 0};1567 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_fp8_bf8;1568 break;1569 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_fp8:1570 ArgsForMatchingMatrixTypes = {3, 0};1571 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_bf8_fp8;1572 break;1573 case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x128_bf8_bf8:1574 ArgsForMatchingMatrixTypes = {3, 0};1575 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f16_16x16x128_bf8_bf8;1576 break;1577 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8:1578 ArgsForMatchingMatrixTypes = {3, 0};1579 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_fp8_fp8;1580 break;1581 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8:1582 ArgsForMatchingMatrixTypes = {3, 0};1583 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_fp8_bf8;1584 break;1585 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8:1586 ArgsForMatchingMatrixTypes = {3, 0};1587 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_bf8_fp8;1588 break;1589 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8:1590 ArgsForMatchingMatrixTypes = {3, 0};1591 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_bf8_bf8;1592 break;1593 case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x64_iu8:1594 ArgsForMatchingMatrixTypes = {4, 1};1595 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_i32_16x16x64_iu8;1596 break;1597 case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x128_f8f6f4:1598 ArgsForMatchingMatrixTypes = {5, 1, 3};1599 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_16x16x128_f8f6f4;1600 break;1601 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4:1602 ArgsForMatchingMatrixTypes = {5, 1, 3};1603 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale_f32_16x16x128_f8f6f4;1604 break;1605 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4:1606 ArgsForMatchingMatrixTypes = {5, 1, 3};1607 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale16_f32_16x16x128_f8f6f4;1608 break;1609 case AMDGPU::BI__builtin_amdgcn_wmma_f32_32x16x128_f4:1610 ArgsForMatchingMatrixTypes = {3, 0, 1};1611 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_f32_32x16x128_f4;1612 break;1613 case AMDGPU::BI__builtin_amdgcn_wmma_scale_f32_32x16x128_f4:1614 ArgsForMatchingMatrixTypes = {3, 0, 1};1615 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale_f32_32x16x128_f4;1616 break;1617 case AMDGPU::BI__builtin_amdgcn_wmma_scale16_f32_32x16x128_f4:1618 ArgsForMatchingMatrixTypes = {3, 0, 1};1619 BuiltinWMMAOp = Intrinsic::amdgcn_wmma_scale16_f32_32x16x128_f4;1620 break;1621 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_f16:1622 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1623 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x64_f16;1624 break;1625 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x64_bf16:1626 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1627 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x64_bf16;1628 break;1629 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x64_f16:1630 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1631 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x64_f16;1632 break;1633 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x64_bf16:1634 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1635 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16_16x16x64_bf16;1636 break;1637 case AMDGPU::BI__builtin_amdgcn_swmmac_bf16f32_16x16x64_bf16:1638 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1639 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_bf16f32_16x16x64_bf16;1640 break;1641 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_fp8:1642 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1643 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_fp8_fp8;1644 break;1645 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_fp8_bf8:1646 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1647 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_fp8_bf8;1648 break;1649 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_fp8:1650 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1651 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_bf8_fp8;1652 break;1653 case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x128_bf8_bf8:1654 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1655 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f32_16x16x128_bf8_bf8;1656 break;1657 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_fp8:1658 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1659 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_fp8_fp8;1660 break;1661 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_fp8_bf8:1662 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1663 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_fp8_bf8;1664 break;1665 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_fp8:1666 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1667 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_bf8_fp8;1668 break;1669 case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x128_bf8_bf8:1670 ArgsForMatchingMatrixTypes = {2, 0, 1, 3};1671 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_f16_16x16x128_bf8_bf8;1672 break;1673 case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x128_iu8:1674 ArgsForMatchingMatrixTypes = {4, 1, 3, 5};1675 BuiltinWMMAOp = Intrinsic::amdgcn_swmmac_i32_16x16x128_iu8;1676 break;1677 }1678 1679 SmallVector<Value *, 6> Args;1680 for (int i = 0, e = E->getNumArgs(); i != e; ++i)1681 Args.push_back(EmitScalarExpr(E->getArg(i)));1682 if (AppendFalseForOpselArg)1683 Args.push_back(Builder.getFalse());1684 1685 SmallVector<llvm::Type *, 6> ArgTypes;1686 if (NeedReturnType)1687 ArgTypes.push_back(ConvertType(E->getType()));1688 for (auto ArgIdx : ArgsForMatchingMatrixTypes)1689 ArgTypes.push_back(Args[ArgIdx]->getType());1690 1691 Function *F = CGM.getIntrinsic(BuiltinWMMAOp, ArgTypes);1692 return Builder.CreateCall(F, Args);1693 }1694 // amdgcn workgroup size1695 case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:1696 return EmitAMDGPUWorkGroupSize(*this, 0);1697 case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:1698 return EmitAMDGPUWorkGroupSize(*this, 1);1699 case AMDGPU::BI__builtin_amdgcn_workgroup_size_z:1700 return EmitAMDGPUWorkGroupSize(*this, 2);1701 1702 // amdgcn grid size1703 case AMDGPU::BI__builtin_amdgcn_grid_size_x:1704 return EmitAMDGPUGridSize(*this, 0);1705 case AMDGPU::BI__builtin_amdgcn_grid_size_y:1706 return EmitAMDGPUGridSize(*this, 1);1707 case AMDGPU::BI__builtin_amdgcn_grid_size_z:1708 return EmitAMDGPUGridSize(*this, 2);1709 1710 // r600 intrinsics1711 case AMDGPU::BI__builtin_r600_recipsqrt_ieee:1712 case AMDGPU::BI__builtin_r600_recipsqrt_ieeef:1713 return emitBuiltinWithOneOverloadedType<1>(*this, E,1714 Intrinsic::r600_recipsqrt_ieee);1715 case AMDGPU::BI__builtin_amdgcn_alignbit: {1716 llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));1717 llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));1718 llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));1719 Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());1720 return Builder.CreateCall(F, { Src0, Src1, Src2 });1721 }1722 case AMDGPU::BI__builtin_amdgcn_fence: {1723 ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),1724 EmitScalarExpr(E->getArg(1)), AO, SSID);1725 FenceInst *Fence = Builder.CreateFence(AO, SSID);1726 if (E->getNumArgs() > 2)1727 AddAMDGPUFenceAddressSpaceMMRA(Fence, E);1728 return Fence;1729 }1730 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:1731 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:1732 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:1733 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:1734 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:1735 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:1736 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:1737 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:1738 case AMDGPU::BI__builtin_amdgcn_ds_faddf:1739 case AMDGPU::BI__builtin_amdgcn_ds_fminf:1740 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:1741 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:1742 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:1743 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:1744 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:1745 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:1746 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:1747 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:1748 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:1749 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:1750 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:1751 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:1752 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {1753 llvm::AtomicRMWInst::BinOp BinOp;1754 switch (BuiltinID) {1755 case AMDGPU::BI__builtin_amdgcn_atomic_inc32:1756 case AMDGPU::BI__builtin_amdgcn_atomic_inc64:1757 BinOp = llvm::AtomicRMWInst::UIncWrap;1758 break;1759 case AMDGPU::BI__builtin_amdgcn_atomic_dec32:1760 case AMDGPU::BI__builtin_amdgcn_atomic_dec64:1761 BinOp = llvm::AtomicRMWInst::UDecWrap;1762 break;1763 case AMDGPU::BI__builtin_amdgcn_ds_faddf:1764 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:1765 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:1766 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:1767 case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:1768 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:1769 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:1770 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:1771 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:1772 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:1773 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:1774 case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:1775 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:1776 BinOp = llvm::AtomicRMWInst::FAdd;1777 break;1778 case AMDGPU::BI__builtin_amdgcn_ds_fminf:1779 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:1780 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:1781 BinOp = llvm::AtomicRMWInst::FMin;1782 break;1783 case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:1784 case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:1785 case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:1786 BinOp = llvm::AtomicRMWInst::FMax;1787 break;1788 }1789 1790 Address Ptr = CheckAtomicAlignment(*this, E);1791 Value *Val = EmitScalarExpr(E->getArg(1));1792 llvm::Type *OrigTy = Val->getType();1793 QualType PtrTy = E->getArg(0)->IgnoreImpCasts()->getType();1794 1795 bool Volatile;1796 1797 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_faddf ||1798 BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fminf ||1799 BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fmaxf) {1800 // __builtin_amdgcn_ds_faddf/fminf/fmaxf has an explicit volatile argument1801 Volatile =1802 cast<ConstantInt>(EmitScalarExpr(E->getArg(4)))->getZExtValue();1803 } else {1804 // Infer volatile from the passed type.1805 Volatile =1806 PtrTy->castAs<PointerType>()->getPointeeType().isVolatileQualified();1807 }1808 1809 if (E->getNumArgs() >= 4) {1810 // Some of the builtins have explicit ordering and scope arguments.1811 ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(2)),1812 EmitScalarExpr(E->getArg(3)), AO, SSID);1813 } else {1814 // Most of the builtins do not have syncscope/order arguments. For DS1815 // atomics the scope doesn't really matter, as they implicitly operate at1816 // workgroup scope.1817 //1818 // The global/flat cases need to use agent scope to consistently produce1819 // the native instruction instead of a cmpxchg expansion.1820 if (getTarget().getTriple().isSPIRV())1821 SSID = getLLVMContext().getOrInsertSyncScopeID("device");1822 else1823 SSID = getLLVMContext().getOrInsertSyncScopeID("agent");1824 AO = AtomicOrdering::Monotonic;1825 1826 // The v2bf16 builtin uses i16 instead of a natural bfloat type.1827 if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16 ||1828 BuiltinID == AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16 ||1829 BuiltinID == AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16) {1830 llvm::Type *V2BF16Ty = FixedVectorType::get(1831 llvm::Type::getBFloatTy(Builder.getContext()), 2);1832 Val = Builder.CreateBitCast(Val, V2BF16Ty);1833 }1834 }1835 1836 llvm::AtomicRMWInst *RMW =1837 Builder.CreateAtomicRMW(BinOp, Ptr, Val, AO, SSID);1838 if (Volatile)1839 RMW->setVolatile(true);1840 1841 unsigned AddrSpace = Ptr.getType()->getAddressSpace();1842 if (AddrSpace != llvm::AMDGPUAS::LOCAL_ADDRESS) {1843 // Most targets require "amdgpu.no.fine.grained.memory" to emit the native1844 // instruction for flat and global operations.1845 llvm::MDTuple *EmptyMD = MDNode::get(getLLVMContext(), {});1846 RMW->setMetadata("amdgpu.no.fine.grained.memory", EmptyMD);1847 1848 // Most targets require "amdgpu.ignore.denormal.mode" to emit the native1849 // instruction, but this only matters for float fadd.1850 if (BinOp == llvm::AtomicRMWInst::FAdd && Val->getType()->isFloatTy())1851 RMW->setMetadata("amdgpu.ignore.denormal.mode", EmptyMD);1852 }1853 1854 return Builder.CreateBitCast(RMW, OrigTy);1855 }1856 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:1857 case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {1858 llvm::Value *Arg = EmitScalarExpr(E->getArg(0));1859 llvm::Type *ResultType = ConvertType(E->getType());1860 // s_sendmsg_rtn is mangled using return type only.1861 Function *F =1862 CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType});1863 return Builder.CreateCall(F, {Arg});1864 }1865 case AMDGPU::BI__builtin_amdgcn_permlane16_swap:1866 case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {1867 // Because builtin types are limited, and the intrinsic uses a struct/pair1868 // output, marshal the pair-of-i32 to <2 x i32>.1869 Value *VDstOld = EmitScalarExpr(E->getArg(0));1870 Value *VSrcOld = EmitScalarExpr(E->getArg(1));1871 Value *FI = EmitScalarExpr(E->getArg(2));1872 Value *BoundCtrl = EmitScalarExpr(E->getArg(3));1873 Function *F =1874 CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16_swap1875 ? Intrinsic::amdgcn_permlane16_swap1876 : Intrinsic::amdgcn_permlane32_swap);1877 llvm::CallInst *Call =1878 Builder.CreateCall(F, {VDstOld, VSrcOld, FI, BoundCtrl});1879 1880 llvm::Value *Elt0 = Builder.CreateExtractValue(Call, 0);1881 llvm::Value *Elt1 = Builder.CreateExtractValue(Call, 1);1882 1883 llvm::Type *ResultType = ConvertType(E->getType());1884 1885 llvm::Value *Insert0 = Builder.CreateInsertElement(1886 llvm::PoisonValue::get(ResultType), Elt0, UINT64_C(0));1887 llvm::Value *AsVector =1888 Builder.CreateInsertElement(Insert0, Elt1, UINT64_C(1));1889 return AsVector;1890 }1891 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:1892 case AMDGPU::BI__builtin_amdgcn_bitop3_b16:1893 return emitBuiltinWithOneOverloadedType<4>(*this, E,1894 Intrinsic::amdgcn_bitop3);1895 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {1896 // TODO: LLVM has this overloaded to allow for fat pointers, but since1897 // those haven't been plumbed through to Clang yet, default to creating the1898 // resource type.1899 SmallVector<Value *, 4> Args;1900 for (unsigned I = 0; I < 4; ++I)1901 Args.push_back(EmitScalarExpr(E->getArg(I)));1902 llvm::PointerType *RetTy = llvm::PointerType::get(1903 Builder.getContext(), llvm::AMDGPUAS::BUFFER_RESOURCE);1904 Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_make_buffer_rsrc,1905 {RetTy, Args[0]->getType()});1906 return Builder.CreateCall(F, Args);1907 }1908 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:1909 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:1910 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:1911 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:1912 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:1913 case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128:1914 return emitBuiltinWithOneOverloadedType<5>(1915 *this, E, Intrinsic::amdgcn_raw_ptr_buffer_store);1916 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:1917 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:1918 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:1919 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:1920 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:1921 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {1922 llvm::Type *RetTy = nullptr;1923 switch (BuiltinID) {1924 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:1925 RetTy = Int8Ty;1926 break;1927 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:1928 RetTy = Int16Ty;1929 break;1930 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:1931 RetTy = Int32Ty;1932 break;1933 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:1934 RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/2);1935 break;1936 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:1937 RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/3);1938 break;1939 case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128:1940 RetTy = llvm::FixedVectorType::get(Int32Ty, /*NumElements=*/4);1941 break;1942 }1943 Function *F =1944 CGM.getIntrinsic(Intrinsic::amdgcn_raw_ptr_buffer_load, RetTy);1945 return Builder.CreateCall(1946 F, {EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)),1947 EmitScalarExpr(E->getArg(2)), EmitScalarExpr(E->getArg(3))});1948 }1949 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32:1950 return emitBuiltinWithOneOverloadedType<5>(1951 *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_add);1952 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32:1953 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16:1954 return emitBuiltinWithOneOverloadedType<5>(1955 *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fadd);1956 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32:1957 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64:1958 return emitBuiltinWithOneOverloadedType<5>(1959 *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmin);1960 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32:1961 case AMDGPU::BI__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64:1962 return emitBuiltinWithOneOverloadedType<5>(1963 *this, E, Intrinsic::amdgcn_raw_ptr_buffer_atomic_fmax);1964 case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:1965 return emitBuiltinWithOneOverloadedType<2>(1966 *this, E, Intrinsic::amdgcn_s_prefetch_data);1967 case Builtin::BIlogbf:1968 case Builtin::BI__builtin_logbf: {1969 Value *Src0 = EmitScalarExpr(E->getArg(0));1970 Function *FrExpFunc = CGM.getIntrinsic(1971 Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()});1972 CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0);1973 Value *Exp = Builder.CreateExtractValue(FrExp, 1);1974 Value *Add = Builder.CreateAdd(1975 Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true);1976 Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getFloatTy());1977 Value *Fabs =1978 emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs);1979 Value *FCmpONE = Builder.CreateFCmpONE(1980 Fabs, ConstantFP::getInfinity(Builder.getFloatTy()));1981 Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs);1982 Value *FCmpOEQ =1983 Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getFloatTy()));1984 Value *Sel2 = Builder.CreateSelect(1985 FCmpOEQ,1986 ConstantFP::getInfinity(Builder.getFloatTy(), /*Negative=*/true), Sel1);1987 return Sel2;1988 }1989 case Builtin::BIlogb:1990 case Builtin::BI__builtin_logb: {1991 Value *Src0 = EmitScalarExpr(E->getArg(0));1992 Function *FrExpFunc = CGM.getIntrinsic(1993 Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()});1994 CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0);1995 Value *Exp = Builder.CreateExtractValue(FrExp, 1);1996 Value *Add = Builder.CreateAdd(1997 Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true);1998 Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getDoubleTy());1999 Value *Fabs =2000 emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs);2001 Value *FCmpONE = Builder.CreateFCmpONE(2002 Fabs, ConstantFP::getInfinity(Builder.getDoubleTy()));2003 Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs);2004 Value *FCmpOEQ =2005 Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getDoubleTy()));2006 Value *Sel2 = Builder.CreateSelect(2007 FCmpOEQ,2008 ConstantFP::getInfinity(Builder.getDoubleTy(), /*Negative=*/true),2009 Sel1);2010 return Sel2;2011 }2012 case Builtin::BIscalbnf:2013 case Builtin::BI__builtin_scalbnf:2014 case Builtin::BIscalbn:2015 case Builtin::BI__builtin_scalbn:2016 return emitBinaryExpMaybeConstrainedFPBuiltin(2017 *this, E, Intrinsic::ldexp, Intrinsic::experimental_constrained_ldexp);2018 default:2019 return nullptr;2020 }2021}2022