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