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