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