1//===- AMDGPU.cpp ---------------------------------------------------------===//
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#include "ABIInfoImpl.h"
10#include "TargetInfo.h"
11#include "clang/AST/DeclCXX.h"
12#include "llvm/ADT/StringExtras.h"
13#include "llvm/IR/LLVMContext.h"
14#include "llvm/IR/MemoryModelRelaxationAnnotations.h"
15#include "llvm/Support/AMDGPUAddrSpace.h"
16
17using namespace clang;
18using namespace clang::CodeGen;
19
20//===----------------------------------------------------------------------===//
21// AMDGPU ABI Implementation
22//===----------------------------------------------------------------------===//
23
24namespace {
25
26class AMDGPUABIInfo final : public DefaultABIInfo {
27private:
28 static const unsigned MaxNumRegsForArgsRet = 16;
29
30 uint64_t numRegsForType(QualType Ty) const;
31
32 bool isHomogeneousAggregateBaseType(QualType Ty) const override;
33 bool isHomogeneousAggregateSmallEnough(const Type *Base,
34 uint64_t Members) const override;
35
36 // Coerce HIP scalar pointer arguments from generic pointers to global ones.
37 llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS,
38 unsigned ToAS) const {
39 // Single value types.
40 auto *PtrTy = llvm::dyn_cast<llvm::PointerType>(Val: Ty);
41 if (PtrTy && PtrTy->getAddressSpace() == FromAS)
42 return llvm::PointerType::get(C&: Ty->getContext(), AddressSpace: ToAS);
43 return Ty;
44 }
45
46public:
47 explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) :
48 DefaultABIInfo(CGT) {}
49
50 ABIArgInfo classifyReturnType(QualType RetTy) const;
51 ABIArgInfo classifyKernelArgumentType(QualType Ty) const;
52 ABIArgInfo classifyArgumentType(QualType Ty, bool Variadic,
53 unsigned &NumRegsLeft) const;
54
55 void computeInfo(CGFunctionInfo &FI) const override;
56 RValue EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, QualType Ty,
57 AggValueSlot Slot) const override;
58
59 llvm::FixedVectorType *
60 getOptimalVectorMemoryType(llvm::FixedVectorType *T,
61 const LangOptions &Opt) const override {
62 // We have legal instructions for 96-bit so 3x32 can be supported.
63 // FIXME: This check should be a subtarget feature as technically SI doesn't
64 // support it.
65 if (T->getNumElements() == 3 && getDataLayout().getTypeSizeInBits(Ty: T) == 96)
66 return T;
67 return DefaultABIInfo::getOptimalVectorMemoryType(T, Opt);
68 }
69};
70
71bool AMDGPUABIInfo::isHomogeneousAggregateBaseType(QualType Ty) const {
72 return true;
73}
74
75bool AMDGPUABIInfo::isHomogeneousAggregateSmallEnough(
76 const Type *Base, uint64_t Members) const {
77 uint32_t NumRegs = (getContext().getTypeSize(T: Base) + 31) / 32;
78
79 // Homogeneous Aggregates may occupy at most 16 registers.
80 return Members * NumRegs <= MaxNumRegsForArgsRet;
81}
82
83/// Estimate number of registers the type will use when passed in registers.
84uint64_t AMDGPUABIInfo::numRegsForType(QualType Ty) const {
85 uint64_t NumRegs = 0;
86
87 if (const VectorType *VT = Ty->getAs<VectorType>()) {
88 // Compute from the number of elements. The reported size is based on the
89 // in-memory size, which includes the padding 4th element for 3-vectors.
90 QualType EltTy = VT->getElementType();
91 uint64_t EltSize = getContext().getTypeSize(T: EltTy);
92
93 // 16-bit element vectors should be passed as packed.
94 if (EltSize == 16)
95 return (VT->getNumElements() + 1) / 2;
96
97 uint64_t EltNumRegs = (EltSize + 31) / 32;
98 return EltNumRegs * VT->getNumElements();
99 }
100
101 if (const auto *RD = Ty->getAsRecordDecl()) {
102 assert(!RD->hasFlexibleArrayMember());
103
104 for (const FieldDecl *Field : RD->fields()) {
105 QualType FieldTy = Field->getType();
106 NumRegs += numRegsForType(Ty: FieldTy);
107 }
108
109 return NumRegs;
110 }
111
112 return (getContext().getTypeSize(T: Ty) + 31) / 32;
113}
114
115void AMDGPUABIInfo::computeInfo(CGFunctionInfo &FI) const {
116 llvm::CallingConv::ID CC = FI.getCallingConvention();
117
118 if (!getCXXABI().classifyReturnType(FI))
119 FI.getReturnInfo() = classifyReturnType(RetTy: FI.getReturnType());
120
121 unsigned ArgumentIndex = 0;
122 const unsigned numFixedArguments = FI.getNumRequiredArgs();
123
124 unsigned NumRegsLeft = MaxNumRegsForArgsRet;
125 for (auto &Arg : FI.arguments()) {
126 if (CC == llvm::CallingConv::AMDGPU_KERNEL) {
127 Arg.info = classifyKernelArgumentType(Ty: Arg.type);
128 } else {
129 bool FixedArgument = ArgumentIndex++ < numFixedArguments;
130 Arg.info = classifyArgumentType(Ty: Arg.type, Variadic: !FixedArgument, NumRegsLeft);
131 }
132 }
133}
134
135RValue AMDGPUABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr,
136 QualType Ty, AggValueSlot Slot) const {
137 const bool IsIndirect = false;
138 const bool AllowHigherAlign = false;
139 return emitVoidPtrVAArg(CGF, VAListAddr, ValueTy: Ty, IsIndirect,
140 ValueInfo: getContext().getTypeInfoInChars(T: Ty),
141 SlotSizeAndAlign: CharUnits::fromQuantity(Quantity: 4), AllowHigherAlign, Slot);
142}
143
144ABIArgInfo AMDGPUABIInfo::classifyReturnType(QualType RetTy) const {
145 if (isAggregateTypeForABI(T: RetTy)) {
146 // Records with non-trivial destructors/copy-constructors should not be
147 // returned by value.
148 if (!getRecordArgABI(T: RetTy, CXXABI&: getCXXABI())) {
149 // Ignore empty structs/unions.
150 if (isEmptyRecord(Context&: getContext(), T: RetTy, AllowArrays: true))
151 return ABIArgInfo::getIgnore();
152
153 // Lower single-element structs to just return a regular value.
154 if (const Type *SeltTy = isSingleElementStruct(T: RetTy, Context&: getContext()))
155 return ABIArgInfo::getDirect(T: CGT.ConvertType(T: QualType(SeltTy, 0)));
156
157 if (const auto *RD = RetTy->getAsRecordDecl();
158 RD && RD->hasFlexibleArrayMember())
159 return DefaultABIInfo::classifyReturnType(RetTy);
160
161 // Pack aggregates <= 4 bytes into single VGPR or pair.
162 uint64_t Size = getContext().getTypeSize(T: RetTy);
163 if (Size <= 16)
164 return ABIArgInfo::getDirect(T: llvm::Type::getInt16Ty(C&: getVMContext()));
165
166 if (Size <= 32)
167 return ABIArgInfo::getDirect(T: llvm::Type::getInt32Ty(C&: getVMContext()));
168
169 if (Size <= 64) {
170 llvm::Type *I32Ty = llvm::Type::getInt32Ty(C&: getVMContext());
171 return ABIArgInfo::getDirect(T: llvm::ArrayType::get(ElementType: I32Ty, NumElements: 2));
172 }
173
174 if (numRegsForType(Ty: RetTy) <= MaxNumRegsForArgsRet)
175 return ABIArgInfo::getDirect();
176 }
177 }
178
179 // Otherwise just do the default thing.
180 return DefaultABIInfo::classifyReturnType(RetTy);
181}
182
183/// For kernels all parameters are really passed in a special buffer. It doesn't
184/// make sense to pass anything byval, so everything must be direct.
185ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const {
186 Ty = useFirstFieldIfTransparentUnion(Ty);
187
188 // TODO: Can we omit empty structs?
189
190 if (const Type *SeltTy = isSingleElementStruct(T: Ty, Context&: getContext()))
191 Ty = QualType(SeltTy, 0);
192
193 llvm::Type *OrigLTy = CGT.ConvertType(T: Ty);
194 llvm::Type *LTy = OrigLTy;
195 if (getContext().getLangOpts().HIP) {
196 LTy = coerceKernelArgumentType(
197 Ty: OrigLTy, /*FromAS=*/getContext().getTargetAddressSpace(AS: LangAS::Default),
198 /*ToAS=*/getContext().getTargetAddressSpace(AS: LangAS::cuda_device));
199 }
200
201 // FIXME: This doesn't apply the optimization of coercing pointers in structs
202 // to global address space when using byref. This would require implementing a
203 // new kind of coercion of the in-memory type when for indirect arguments.
204 if (LTy == OrigLTy && isAggregateTypeForABI(T: Ty)) {
205 return ABIArgInfo::getIndirectAliased(
206 Alignment: getContext().getTypeAlignInChars(T: Ty),
207 AddrSpace: getContext().getTargetAddressSpace(AS: LangAS::opencl_constant),
208 Realign: false /*Realign*/, Padding: nullptr /*Padding*/);
209 }
210
211 // If we set CanBeFlattened to true, CodeGen will expand the struct to its
212 // individual elements, which confuses the Clover OpenCL backend; therefore we
213 // have to set it to false here. Other args of getDirect() are just defaults.
214 return ABIArgInfo::getDirect(T: LTy, Offset: 0, Padding: nullptr, CanBeFlattened: false);
215}
216
217ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty, bool Variadic,
218 unsigned &NumRegsLeft) const {
219 assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow");
220
221 Ty = useFirstFieldIfTransparentUnion(Ty);
222
223 if (Variadic) {
224 return ABIArgInfo::getDirect(/*T=*/nullptr,
225 /*Offset=*/0,
226 /*Padding=*/nullptr,
227 /*CanBeFlattened=*/false,
228 /*Align=*/0);
229 }
230
231 if (isAggregateTypeForABI(T: Ty)) {
232 // Records with non-trivial destructors/copy-constructors should not be
233 // passed by value.
234 if (auto RAA = getRecordArgABI(T: Ty, CXXABI&: getCXXABI()))
235 return getNaturalAlignIndirect(Ty, AddrSpace: getDataLayout().getAllocaAddrSpace(),
236 ByVal: RAA == CGCXXABI::RAA_DirectInMemory);
237
238 // Ignore empty structs/unions.
239 if (isEmptyRecord(Context&: getContext(), T: Ty, AllowArrays: true))
240 return ABIArgInfo::getIgnore();
241
242 // Lower single-element structs to just pass a regular value. TODO: We
243 // could do reasonable-size multiple-element structs too, using getExpand(),
244 // though watch out for things like bitfields.
245 if (const Type *SeltTy = isSingleElementStruct(T: Ty, Context&: getContext()))
246 return ABIArgInfo::getDirect(T: CGT.ConvertType(T: QualType(SeltTy, 0)));
247
248 if (const auto *RD = Ty->getAsRecordDecl();
249 RD && RD->hasFlexibleArrayMember())
250 return DefaultABIInfo::classifyArgumentType(RetTy: Ty);
251
252 // Pack aggregates <= 8 bytes into single VGPR or pair.
253 uint64_t Size = getContext().getTypeSize(T: Ty);
254 if (Size <= 64) {
255 unsigned NumRegs = (Size + 31) / 32;
256 NumRegsLeft -= std::min(a: NumRegsLeft, b: NumRegs);
257
258 if (Size <= 16)
259 return ABIArgInfo::getDirect(T: llvm::Type::getInt16Ty(C&: getVMContext()));
260
261 if (Size <= 32)
262 return ABIArgInfo::getDirect(T: llvm::Type::getInt32Ty(C&: getVMContext()));
263
264 // XXX: Should this be i64 instead, and should the limit increase?
265 llvm::Type *I32Ty = llvm::Type::getInt32Ty(C&: getVMContext());
266 return ABIArgInfo::getDirect(T: llvm::ArrayType::get(ElementType: I32Ty, NumElements: 2));
267 }
268
269 if (NumRegsLeft > 0) {
270 uint64_t NumRegs = numRegsForType(Ty);
271 if (NumRegsLeft >= NumRegs) {
272 NumRegsLeft -= NumRegs;
273 return ABIArgInfo::getDirect();
274 }
275 }
276
277 // Use pass-by-reference in stead of pass-by-value for struct arguments in
278 // function ABI.
279 return ABIArgInfo::getIndirectAliased(
280 Alignment: getContext().getTypeAlignInChars(T: Ty),
281 AddrSpace: getContext().getTargetAddressSpace(AS: LangAS::opencl_private));
282 }
283
284 // Otherwise just do the default thing.
285 ABIArgInfo ArgInfo = DefaultABIInfo::classifyArgumentType(RetTy: Ty);
286 if (!ArgInfo.isIndirect()) {
287 uint64_t NumRegs = numRegsForType(Ty);
288 NumRegsLeft -= std::min(a: NumRegs, b: uint64_t{NumRegsLeft});
289 }
290
291 return ArgInfo;
292}
293
294class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
295public:
296 AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT)
297 : TargetCodeGenInfo(std::make_unique<AMDGPUABIInfo>(args&: CGT)) {}
298
299 bool supportsLibCall() const override { return false; }
300 void setFunctionDeclAttributes(const FunctionDecl *FD, llvm::Function *F,
301 CodeGenModule &CGM) const;
302
303 void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
304 CodeGen::CodeGenModule &M) const override;
305 unsigned getDeviceKernelCallingConv() const override;
306
307 llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
308 llvm::PointerType *T, QualType QT) const override;
309
310 LangAS getSRetAddrSpace(const CXXRecordDecl *RD) const override;
311
312 LangAS getGlobalVarAddressSpace(CodeGenModule &CGM,
313 const VarDecl *D) const override;
314 StringRef getLLVMSyncScopeStr(const LangOptions &LangOpts, SyncScope Scope,
315 llvm::AtomicOrdering Ordering) const override;
316 void setTargetAtomicMetadata(CodeGenFunction &CGF,
317 llvm::Instruction &AtomicInst,
318 const AtomicExpr *Expr = nullptr) const override;
319 llvm::Value *createEnqueuedBlockKernel(CodeGenFunction &CGF,
320 llvm::Function *BlockInvokeFunc,
321 llvm::Type *BlockTy) const override;
322 bool shouldEmitStaticExternCAliases() const override;
323 bool shouldEmitDWARFBitFieldSeparators() const override;
324 void setCUDAKernelCallingConvention(const FunctionType *&FT) const override;
325};
326}
327
328static bool requiresAMDGPUProtectedVisibility(const Decl *D,
329 llvm::GlobalValue *GV) {
330 if (GV->getVisibility() != llvm::GlobalValue::HiddenVisibility)
331 return false;
332
333 return !D->hasAttr<OMPDeclareTargetDeclAttr>() &&
334 (D->hasAttr<DeviceKernelAttr>() ||
335 (isa<FunctionDecl>(Val: D) && D->hasAttr<CUDAGlobalAttr>()) ||
336 (isa<VarDecl>(Val: D) &&
337 (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
338 cast<VarDecl>(Val: D)->getType()->isCUDADeviceBuiltinSurfaceType() ||
339 cast<VarDecl>(Val: D)->getType()->isCUDADeviceBuiltinTextureType())));
340}
341
342void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes(
343 const FunctionDecl *FD, llvm::Function *F, CodeGenModule &M) const {
344 const auto *ReqdWGS =
345 M.getLangOpts().OpenCL ? FD->getAttr<ReqdWorkGroupSizeAttr>() : nullptr;
346 const bool IsOpenCLKernel =
347 M.getLangOpts().OpenCL && FD->hasAttr<DeviceKernelAttr>();
348 const bool IsHIPKernel = M.getLangOpts().HIP && FD->hasAttr<CUDAGlobalAttr>();
349
350 const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
351
352 // __launch_bounds__ only takes effect on kernels and is silently ignored on
353 // other functions The arguments are honored only if the equivalent native
354 // amdgpu_flat_work_group_size / amdgpu_waves_per_eu attribute was not also
355 // used out; those take precedence.
356 const auto *LaunchBounds =
357 IsHIPKernel ? FD->getAttr<CUDALaunchBoundsAttr>() : nullptr;
358 unsigned LBMaxThreads = 0;
359 unsigned LBMinWaves = 0;
360 if (LaunchBounds) {
361 LBMaxThreads = LaunchBounds->getMaxThreads()
362 ->EvaluateKnownConstInt(Ctx: M.getContext())
363 .getExtValue();
364 if (const Expr *MinBlocks = LaunchBounds->getMinBlocks()) {
365 LBMinWaves =
366 MinBlocks->EvaluateKnownConstInt(Ctx: M.getContext()).getExtValue();
367 }
368 }
369
370 if (ReqdWGS || FlatWGS) {
371 M.handleAMDGPUFlatWorkGroupSizeAttr(F, A: FlatWGS, ReqdWGS);
372 } else if (LBMaxThreads > 0) {
373 F->addFnAttr(Kind: "amdgpu-flat-work-group-size",
374 Val: "1," + llvm::utostr(X: LBMaxThreads));
375 } else if (IsOpenCLKernel || IsHIPKernel) {
376 // By default, restrict the maximum size to a value specified by
377 // --gpu-max-threads-per-block=n or its default value for HIP.
378 const unsigned OpenCLDefaultMaxWorkGroupSize = 256;
379 const unsigned DefaultMaxWorkGroupSize =
380 IsOpenCLKernel ? OpenCLDefaultMaxWorkGroupSize
381 : M.getLangOpts().GPUMaxThreadsPerBlock;
382 std::string AttrVal =
383 std::string("1,") + llvm::utostr(X: DefaultMaxWorkGroupSize);
384 F->addFnAttr(Kind: "amdgpu-flat-work-group-size", Val: AttrVal);
385 }
386
387 if (const auto *Attr = FD->getAttr<AMDGPUWavesPerEUAttr>()) {
388 M.handleAMDGPUWavesPerEUAttr(F, A: Attr);
389 } else if (LBMinWaves > 0) {
390 // HIP reinterprets the second argument as the minimum waves per EU.
391 //
392 // TODO: The third argument (maxclusterrank) could be used if the AMDGPU
393 // "clusters" feature is supported for the current subtarget.
394 F->addFnAttr(Kind: "amdgpu-waves-per-eu", Val: llvm::utostr(X: LBMinWaves));
395 }
396
397 if (const auto *Attr = FD->getAttr<AMDGPUNumSGPRAttr>()) {
398 unsigned NumSGPR = Attr->getNumSGPR();
399
400 if (NumSGPR != 0)
401 F->addFnAttr(Kind: "amdgpu-num-sgpr", Val: llvm::utostr(X: NumSGPR));
402 }
403
404 if (const auto *Attr = FD->getAttr<AMDGPUNumVGPRAttr>()) {
405 uint32_t NumVGPR = Attr->getNumVGPR();
406
407 if (NumVGPR != 0)
408 F->addFnAttr(Kind: "amdgpu-num-vgpr", Val: llvm::utostr(X: NumVGPR));
409 }
410
411 if (const auto *Attr = FD->getAttr<AMDGPUMaxNumWorkGroupsAttr>()) {
412 uint32_t X = Attr->getMaxNumWorkGroupsX()
413 ->EvaluateKnownConstInt(Ctx: M.getContext())
414 .getExtValue();
415 // Y and Z dimensions default to 1 if not specified
416 uint32_t Y = Attr->getMaxNumWorkGroupsY()
417 ? Attr->getMaxNumWorkGroupsY()
418 ->EvaluateKnownConstInt(Ctx: M.getContext())
419 .getExtValue()
420 : 1;
421 uint32_t Z = Attr->getMaxNumWorkGroupsZ()
422 ? Attr->getMaxNumWorkGroupsZ()
423 ->EvaluateKnownConstInt(Ctx: M.getContext())
424 .getExtValue()
425 : 1;
426
427 llvm::SmallString<32> AttrVal;
428 llvm::raw_svector_ostream OS(AttrVal);
429 OS << X << ',' << Y << ',' << Z;
430
431 F->addFnAttr(Kind: "amdgpu-max-num-workgroups", Val: AttrVal.str());
432 }
433
434 if (auto *Attr = FD->getAttr<CUDAClusterDimsAttr>()) {
435 auto GetExprVal = [&](const auto &E) {
436 return E ? E->EvaluateKnownConstInt(M.getContext()).getExtValue() : 1;
437 };
438 unsigned X = GetExprVal(Attr->getX());
439 unsigned Y = GetExprVal(Attr->getY());
440 unsigned Z = GetExprVal(Attr->getZ());
441 llvm::SmallString<32> AttrVal;
442 llvm::raw_svector_ostream OS(AttrVal);
443 OS << X << ',' << Y << ',' << Z;
444 F->addFnAttr(Kind: "amdgpu-cluster-dims", Val: AttrVal.str());
445 }
446
447 // OpenCL doesn't support cluster feature.
448 const TargetInfo &TTI = M.getContext().getTargetInfo();
449 if ((IsOpenCLKernel &&
450 TTI.hasFeatureEnabled(Features: TTI.getTargetOpts().FeatureMap, Name: "clusters")) ||
451 FD->hasAttr<CUDANoClusterAttr>())
452 F->addFnAttr(Kind: "amdgpu-cluster-dims", Val: "0,0,0");
453}
454
455void AMDGPUTargetCodeGenInfo::setTargetAttributes(
456 const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
457 if (requiresAMDGPUProtectedVisibility(D, GV)) {
458 GV->setVisibility(llvm::GlobalValue::ProtectedVisibility);
459 GV->setDSOLocal(true);
460 }
461
462 if (GV->isDeclaration())
463 return;
464
465 llvm::Function *F = dyn_cast<llvm::Function>(Val: GV);
466 if (!F)
467 return;
468
469 const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(Val: D);
470 if (FD)
471 setFunctionDeclAttributes(FD, F, M);
472 if (!getABIInfo().getCodeGenOpts().EmitIEEENaNCompliantInsts)
473 F->addFnAttr(Kind: "amdgpu-ieee", Val: "false");
474 if (getABIInfo().getCodeGenOpts().AMDGPUExpandWaitcntProfiling)
475 F->addFnAttr(Kind: "amdgpu-expand-waitcnt-profiling");
476}
477
478unsigned AMDGPUTargetCodeGenInfo::getDeviceKernelCallingConv() const {
479 return llvm::CallingConv::AMDGPU_KERNEL;
480}
481
482// Currently LLVM assumes null pointers always have value 0,
483// which results in incorrectly transformed IR. Therefore, instead of
484// emitting null pointers in private and local address spaces, a null
485// pointer in generic address space is emitted which is casted to a
486// pointer in local or private address space.
487llvm::Constant *AMDGPUTargetCodeGenInfo::getNullPointer(
488 const CodeGen::CodeGenModule &CGM, llvm::PointerType *PT,
489 QualType QT) const {
490 if (CGM.getContext().getTargetNullPointerValue(QT) == 0)
491 return llvm::ConstantPointerNull::get(T: PT);
492
493 auto &Ctx = CGM.getContext();
494 auto NPT = llvm::PointerType::get(
495 C&: PT->getContext(), AddressSpace: Ctx.getTargetAddressSpace(AS: LangAS::opencl_generic));
496 return llvm::ConstantExpr::getAddrSpaceCast(
497 C: llvm::ConstantPointerNull::get(T: NPT), Ty: PT);
498}
499
500LangAS
501AMDGPUTargetCodeGenInfo::getSRetAddrSpace(const CXXRecordDecl *RD) const {
502 // Types with no viable copy/move must be constructed in-place , use the
503 // default AS so the sret pointer matches the "this" convention.
504 if (RD && !RD->canPassInRegisters())
505 return LangAS::Default;
506 return getLangASFromTargetAS(
507 TargetAS: getABIInfo().getDataLayout().getAllocaAddrSpace());
508}
509
510LangAS
511AMDGPUTargetCodeGenInfo::getGlobalVarAddressSpace(CodeGenModule &CGM,
512 const VarDecl *D) const {
513 assert(!CGM.getLangOpts().OpenCL &&
514 !(CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice) &&
515 "Address space agnostic languages only");
516 LangAS DefaultGlobalAS = getLangASFromTargetAS(
517 TargetAS: CGM.getContext().getTargetAddressSpace(AS: LangAS::opencl_global));
518 if (!D)
519 return DefaultGlobalAS;
520
521 LangAS AddrSpace = D->getType().getAddressSpace();
522 if (AddrSpace != LangAS::Default)
523 return AddrSpace;
524
525 // Only promote to address space 4 if VarDecl has constant initialization.
526 if (D->getType().isConstantStorage(Ctx: CGM.getContext(), ExcludeCtor: false, ExcludeDtor: false) &&
527 D->hasConstantInitialization()) {
528 if (auto ConstAS = CGM.getTarget().getConstantAddressSpace())
529 return *ConstAS;
530 }
531 return DefaultGlobalAS;
532}
533
534StringRef AMDGPUTargetCodeGenInfo::getLLVMSyncScopeStr(
535 const LangOptions &LangOpts, SyncScope Scope,
536 llvm::AtomicOrdering Ordering) const {
537
538 // OpenCL assumes by default that atomic scopes are per-address space for
539 // non-sequentially consistent operations.
540 bool IsOneAs = (Scope >= SyncScope::OpenCLWorkGroup &&
541 Scope <= SyncScope::OpenCLSubGroup &&
542 Ordering != llvm::AtomicOrdering::SequentiallyConsistent);
543
544 llvm::AtomicScope AS = getAtomicScope(S: Scope);
545 assert((AS != llvm::AtomicScope::Cluster || !IsOneAs) &&
546 "OpenCL does not have cluster scope");
547 return *llvm::getAtomicScopeIRString(T: getABIInfo().getTarget().getTriple(), S: AS,
548 IsSingleAddressSpace: IsOneAs);
549}
550
551void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata(
552 CodeGenFunction &CGF, llvm::Instruction &AtomicInst,
553 const AtomicExpr *AE) const {
554 auto *RMW = dyn_cast<llvm::AtomicRMWInst>(Val: &AtomicInst);
555 auto *CmpX = dyn_cast<llvm::AtomicCmpXchgInst>(Val: &AtomicInst);
556
557 // OpenCL and old style HIP atomics consider atomics targeting thread private
558 // memory to be undefined.
559 //
560 // TODO: This is probably undefined for atomic load/store, but there's not
561 // much direct codegen benefit to knowing this.
562 if (((RMW && RMW->getPointerAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS) ||
563 (CmpX &&
564 CmpX->getPointerAddressSpace() == llvm::AMDGPUAS::FLAT_ADDRESS)) &&
565 AE && AE->threadPrivateMemoryAtomicsAreUndefined()) {
566 llvm::MDBuilder MDHelper(CGF.getLLVMContext());
567 llvm::MDNode *ASRange = MDHelper.createRange(
568 Lo: llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS),
569 Hi: llvm::APInt(32, llvm::AMDGPUAS::PRIVATE_ADDRESS + 1));
570 AtomicInst.setMetadata(KindID: llvm::LLVMContext::MD_noalias_addrspace, Node: ASRange);
571 }
572
573 CGF.AddAMDGPUAvailableVisibleMMRA(Inst: &AtomicInst);
574
575 if (!RMW)
576 return;
577
578 AtomicOptions AO = CGF.CGM.getAtomicOpts();
579 llvm::MDNode *Empty = llvm::MDNode::get(Context&: CGF.getLLVMContext(), MDs: {});
580 if (!AO.getOption(Kind: clang::AtomicOptionKind::FineGrainedMemory))
581 RMW->setMetadata(Kind: "amdgpu.no.fine.grained.memory", Node: Empty);
582 if (!AO.getOption(Kind: clang::AtomicOptionKind::RemoteMemory))
583 RMW->setMetadata(Kind: "amdgpu.no.remote.memory", Node: Empty);
584 if (AO.getOption(Kind: clang::AtomicOptionKind::IgnoreDenormalMode) &&
585 RMW->getOperation() == llvm::AtomicRMWInst::FAdd &&
586 RMW->getType()->isFloatTy())
587 RMW->setMetadata(KindID: llvm::LLVMContext::MD_atomic_ignore_denormal_mode, Node: Empty);
588}
589
590bool AMDGPUTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
591 return false;
592}
593
594bool AMDGPUTargetCodeGenInfo::shouldEmitDWARFBitFieldSeparators() const {
595 return true;
596}
597
598void AMDGPUTargetCodeGenInfo::setCUDAKernelCallingConvention(
599 const FunctionType *&FT) const {
600 FT = getABIInfo().getContext().adjustFunctionType(
601 Fn: FT, EInfo: FT->getExtInfo().withCallingConv(cc: CC_DeviceKernel));
602}
603
604/// Return IR struct type for rtinfo struct in rocm-device-libs used for device
605/// enqueue.
606///
607/// ptr addrspace(1) kernel_object, i32 private_segment_size,
608/// i32 group_segment_size
609
610static llvm::StructType *
611getAMDGPURuntimeHandleType(llvm::LLVMContext &C,
612 llvm::Type *KernelDescriptorPtrTy) {
613 llvm::Type *Int32 = llvm::Type::getInt32Ty(C);
614 return llvm::StructType::create(Context&: C, Elements: {KernelDescriptorPtrTy, Int32, Int32},
615 Name: "block.runtime.handle.t");
616}
617
618/// Create an OpenCL kernel for an enqueued block.
619///
620/// The type of the first argument (the block literal) is the struct type
621/// of the block literal instead of a pointer type. The first argument
622/// (block literal) is passed directly by value to the kernel. The kernel
623/// allocates the same type of struct on stack and stores the block literal
624/// to it and passes its pointer to the block invoke function. The kernel
625/// has "enqueued-block" function attribute and kernel argument metadata.
626llvm::Value *AMDGPUTargetCodeGenInfo::createEnqueuedBlockKernel(
627 CodeGenFunction &CGF, llvm::Function *Invoke, llvm::Type *BlockTy) const {
628 auto &Builder = CGF.Builder;
629 auto &C = CGF.getLLVMContext();
630
631 auto *InvokeFT = Invoke->getFunctionType();
632 llvm::SmallVector<llvm::Type *, 2> ArgTys;
633 llvm::SmallVector<llvm::Metadata *, 8> AddressQuals;
634 llvm::SmallVector<llvm::Metadata *, 8> AccessQuals;
635 llvm::SmallVector<llvm::Metadata *, 8> ArgTypeNames;
636 llvm::SmallVector<llvm::Metadata *, 8> ArgBaseTypeNames;
637 llvm::SmallVector<llvm::Metadata *, 8> ArgTypeQuals;
638 llvm::SmallVector<llvm::Metadata *, 8> ArgNames;
639
640 ArgTys.push_back(Elt: BlockTy);
641 ArgTypeNames.push_back(Elt: llvm::MDString::get(Context&: C, Str: "__block_literal"));
642 AddressQuals.push_back(Elt: llvm::ConstantAsMetadata::get(C: Builder.getInt32(C: 0)));
643 ArgBaseTypeNames.push_back(Elt: llvm::MDString::get(Context&: C, Str: "__block_literal"));
644 ArgTypeQuals.push_back(Elt: llvm::MDString::get(Context&: C, Str: ""));
645 AccessQuals.push_back(Elt: llvm::MDString::get(Context&: C, Str: "none"));
646 ArgNames.push_back(Elt: llvm::MDString::get(Context&: C, Str: "block_literal"));
647 for (unsigned I = 1, E = InvokeFT->getNumParams(); I < E; ++I) {
648 ArgTys.push_back(Elt: InvokeFT->getParamType(i: I));
649 ArgTypeNames.push_back(Elt: llvm::MDString::get(Context&: C, Str: "void*"));
650 AddressQuals.push_back(Elt: llvm::ConstantAsMetadata::get(C: Builder.getInt32(C: 3)));
651 AccessQuals.push_back(Elt: llvm::MDString::get(Context&: C, Str: "none"));
652 ArgBaseTypeNames.push_back(Elt: llvm::MDString::get(Context&: C, Str: "void*"));
653 ArgTypeQuals.push_back(Elt: llvm::MDString::get(Context&: C, Str: ""));
654 ArgNames.push_back(
655 Elt: llvm::MDString::get(Context&: C, Str: (Twine("local_arg") + Twine(I)).str()));
656 }
657
658 llvm::Module &Mod = CGF.CGM.getModule();
659 const llvm::DataLayout &DL = Mod.getDataLayout();
660
661 llvm::Twine Name = Invoke->getName() + "_kernel";
662 auto *FT = llvm::FunctionType::get(Result: llvm::Type::getVoidTy(C), Params: ArgTys, isVarArg: false);
663
664 // The kernel itself can be internal, the runtime does not directly access the
665 // kernel address (only the kernel descriptor).
666 auto *F = llvm::Function::Create(Ty: FT, Linkage: llvm::GlobalValue::InternalLinkage, N: Name,
667 M: &Mod);
668 F->setCallingConv(getDeviceKernelCallingConv());
669
670 llvm::AttrBuilder KernelAttrs(C);
671 // FIXME: The invoke isn't applying the right attributes either
672 // FIXME: This is missing setTargetAttributes
673 CGF.CGM.addDefaultFunctionDefinitionAttributes(attrs&: KernelAttrs);
674 F->addFnAttrs(Attrs: KernelAttrs);
675
676 auto IP = CGF.Builder.saveIP();
677 auto *BB = llvm::BasicBlock::Create(Context&: C, Name: "entry", Parent: F);
678 Builder.SetInsertPoint(BB);
679 const auto BlockAlign = DL.getPrefTypeAlign(Ty: BlockTy);
680 auto *BlockPtr = Builder.CreateAlloca(Ty: BlockTy, ArraySize: nullptr);
681 BlockPtr->setAlignment(BlockAlign);
682 Builder.CreateAlignedStore(Val: F->arg_begin(), Ptr: BlockPtr, Align: BlockAlign);
683 auto *Cast = Builder.CreatePointerCast(V: BlockPtr, DestTy: InvokeFT->getParamType(i: 0));
684 llvm::SmallVector<llvm::Value *, 2> Args;
685 Args.push_back(Elt: Cast);
686 for (llvm::Argument &A : llvm::drop_begin(RangeOrContainer: F->args()))
687 Args.push_back(Elt: &A);
688 llvm::CallInst *call = Builder.CreateCall(Callee: Invoke, Args);
689 call->setCallingConv(Invoke->getCallingConv());
690 Builder.CreateRetVoid();
691 Builder.restoreIP(IP);
692
693 F->setMetadata(Kind: "kernel_arg_addr_space", Node: llvm::MDNode::get(Context&: C, MDs: AddressQuals));
694 F->setMetadata(Kind: "kernel_arg_access_qual", Node: llvm::MDNode::get(Context&: C, MDs: AccessQuals));
695 F->setMetadata(Kind: "kernel_arg_type", Node: llvm::MDNode::get(Context&: C, MDs: ArgTypeNames));
696 F->setMetadata(Kind: "kernel_arg_base_type",
697 Node: llvm::MDNode::get(Context&: C, MDs: ArgBaseTypeNames));
698 F->setMetadata(Kind: "kernel_arg_type_qual", Node: llvm::MDNode::get(Context&: C, MDs: ArgTypeQuals));
699 if (CGF.CGM.getCodeGenOpts().EmitOpenCLArgMetadata)
700 F->setMetadata(Kind: "kernel_arg_name", Node: llvm::MDNode::get(Context&: C, MDs: ArgNames));
701
702 llvm::StructType *HandleTy = getAMDGPURuntimeHandleType(
703 C, KernelDescriptorPtrTy: llvm::PointerType::get(C, AddressSpace: DL.getDefaultGlobalsAddressSpace()));
704 llvm::Constant *RuntimeHandleInitializer =
705 llvm::ConstantAggregateZero::get(Ty: HandleTy);
706
707 llvm::Twine RuntimeHandleName = F->getName() + ".runtime.handle";
708
709 // The runtime needs access to the runtime handle as an external symbol. The
710 // runtime handle will need to be made external later, in
711 // AMDGPUExportOpenCLEnqueuedBlocks. The kernel itself has a hidden reference
712 // inside the runtime handle, and is not directly referenced.
713
714 // TODO: We would initialize the first field by declaring F->getName() + ".kd"
715 // to reference the kernel descriptor. The runtime wouldn't need to bother
716 // setting it. We would need to have a final symbol name though.
717 // TODO: Can we directly use an external symbol with getGlobalIdentifier?
718 auto *RuntimeHandle = new llvm::GlobalVariable(
719 Mod, HandleTy,
720 /*isConstant=*/true, llvm::GlobalValue::InternalLinkage,
721 /*Initializer=*/RuntimeHandleInitializer, RuntimeHandleName,
722 /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
723 DL.getDefaultGlobalsAddressSpace(),
724 /*isExternallyInitialized=*/true);
725
726 llvm::MDNode *HandleAsMD =
727 llvm::MDNode::get(Context&: C, MDs: llvm::ValueAsMetadata::get(V: RuntimeHandle));
728 F->setMetadata(KindID: llvm::LLVMContext::MD_associated, Node: HandleAsMD);
729
730 RuntimeHandle->setSection(".amdgpu.kernel.runtime.handle");
731
732 CGF.CGM.addUsedGlobal(GV: F);
733 CGF.CGM.addUsedGlobal(GV: RuntimeHandle);
734 return RuntimeHandle;
735}
736
737void CodeGenModule::handleAMDGPUFlatWorkGroupSizeAttr(
738 llvm::Function *F, const AMDGPUFlatWorkGroupSizeAttr *FlatWGS,
739 const ReqdWorkGroupSizeAttr *ReqdWGS, int32_t *MinThreadsVal,
740 int32_t *MaxThreadsVal) {
741 unsigned Min = 0;
742 unsigned Max = 0;
743 auto Eval = [&](Expr *E) {
744 return E->EvaluateKnownConstInt(Ctx: getContext()).getExtValue();
745 };
746 if (ReqdWGS) {
747 Min = Max = Eval(ReqdWGS->getXDim()) * Eval(ReqdWGS->getYDim()) *
748 Eval(ReqdWGS->getZDim());
749 } else if (FlatWGS) {
750 Min = Eval(FlatWGS->getMin());
751 Max = Eval(FlatWGS->getMax());
752 }
753
754 if (Min != 0 || ReqdWGS) {
755 assert(Min <= Max && "Min must be less than or equal Max");
756
757 if (MinThreadsVal)
758 *MinThreadsVal = Min;
759 if (MaxThreadsVal)
760 *MaxThreadsVal = Max;
761 std::string AttrVal = llvm::utostr(X: Min) + "," + llvm::utostr(X: Max);
762 if (F)
763 F->addFnAttr(Kind: "amdgpu-flat-work-group-size", Val: AttrVal);
764 } else
765 assert(Max == 0 && "Max must be zero");
766}
767
768void CodeGenModule::handleAMDGPUWavesPerEUAttr(
769 llvm::Function *F, const AMDGPUWavesPerEUAttr *Attr) {
770 unsigned Min =
771 Attr->getMin()->EvaluateKnownConstInt(Ctx: getContext()).getExtValue();
772 unsigned Max =
773 Attr->getMax()
774 ? Attr->getMax()->EvaluateKnownConstInt(Ctx: getContext()).getExtValue()
775 : 0;
776
777 if (Min != 0) {
778 assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max");
779
780 std::string AttrVal = llvm::utostr(X: Min);
781 if (Max != 0)
782 AttrVal = AttrVal + "," + llvm::utostr(X: Max);
783 F->addFnAttr(Kind: "amdgpu-waves-per-eu", Val: AttrVal);
784 } else
785 assert(Max == 0 && "Max must be zero");
786}
787
788std::unique_ptr<TargetCodeGenInfo>
789CodeGen::createAMDGPUTargetCodeGenInfo(CodeGenModule &CGM) {
790 return std::make_unique<AMDGPUTargetCodeGenInfo>(args&: CGM.getTypes());
791}
792