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