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