xref: /freebsd-src/contrib/llvm-project/llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp (revision 0eae32dcef82f6f06de6419a0d623d7def0cc8f6) 
1 //===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===//
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 /// \file
10 /// AMDGPU HSA Metadata Streamer.
11 ///
12 //
13 //===----------------------------------------------------------------------===//
14 
15 #include "AMDGPUHSAMetadataStreamer.h"
16 #include "AMDGPU.h"
17 #include "GCNSubtarget.h"
18 #include "MCTargetDesc/AMDGPUTargetStreamer.h"
19 #include "SIMachineFunctionInfo.h"
20 #include "SIProgramInfo.h"
21 #include "llvm/IR/Module.h"
22 using namespace llvm;
23 
24 static std::pair<Type *, Align> getArgumentTypeAlign(const Argument &Arg,
25                                                      const DataLayout &DL) {
26   Type *Ty = Arg.getType();
27   MaybeAlign ArgAlign;
28   if (Arg.hasByRefAttr()) {
29     Ty = Arg.getParamByRefType();
30     ArgAlign = Arg.getParamAlign();
31   }
32 
33   if (!ArgAlign)
34     ArgAlign = DL.getABITypeAlign(Ty);
35 
36   return std::make_pair(Ty, *ArgAlign);
37 }
38 
39 namespace llvm {
40 
41 static cl::opt<bool> DumpHSAMetadata(
42     "amdgpu-dump-hsa-metadata",
43     cl::desc("Dump AMDGPU HSA Metadata"));
44 static cl::opt<bool> VerifyHSAMetadata(
45     "amdgpu-verify-hsa-metadata",
46     cl::desc("Verify AMDGPU HSA Metadata"));
47 
48 namespace AMDGPU {
49 namespace HSAMD {
50 
51 //===----------------------------------------------------------------------===//
52 // HSAMetadataStreamerV2
53 //===----------------------------------------------------------------------===//
54 void MetadataStreamerV2::dump(StringRef HSAMetadataString) const {
55   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
56 }
57 
58 void MetadataStreamerV2::verify(StringRef HSAMetadataString) const {
59   errs() << "AMDGPU HSA Metadata Parser Test: ";
60 
61   HSAMD::Metadata FromHSAMetadataString;
62   if (fromString(HSAMetadataString, FromHSAMetadataString)) {
63     errs() << "FAIL\n";
64     return;
65   }
66 
67   std::string ToHSAMetadataString;
68   if (toString(FromHSAMetadataString, ToHSAMetadataString)) {
69     errs() << "FAIL\n";
70     return;
71   }
72 
73   errs() << (HSAMetadataString == ToHSAMetadataString ? "PASS" : "FAIL")
74          << '\n';
75   if (HSAMetadataString != ToHSAMetadataString) {
76     errs() << "Original input: " << HSAMetadataString << '\n'
77            << "Produced output: " << ToHSAMetadataString << '\n';
78   }
79 }
80 
81 AccessQualifier
82 MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
83   if (AccQual.empty())
84     return AccessQualifier::Unknown;
85 
86   return StringSwitch<AccessQualifier>(AccQual)
87              .Case("read_only",  AccessQualifier::ReadOnly)
88              .Case("write_only", AccessQualifier::WriteOnly)
89              .Case("read_write", AccessQualifier::ReadWrite)
90              .Default(AccessQualifier::Default);
91 }
92 
93 AddressSpaceQualifier
94 MetadataStreamerV2::getAddressSpaceQualifier(
95     unsigned AddressSpace) const {
96   switch (AddressSpace) {
97   case AMDGPUAS::PRIVATE_ADDRESS:
98     return AddressSpaceQualifier::Private;
99   case AMDGPUAS::GLOBAL_ADDRESS:
100     return AddressSpaceQualifier::Global;
101   case AMDGPUAS::CONSTANT_ADDRESS:
102     return AddressSpaceQualifier::Constant;
103   case AMDGPUAS::LOCAL_ADDRESS:
104     return AddressSpaceQualifier::Local;
105   case AMDGPUAS::FLAT_ADDRESS:
106     return AddressSpaceQualifier::Generic;
107   case AMDGPUAS::REGION_ADDRESS:
108     return AddressSpaceQualifier::Region;
109   default:
110     return AddressSpaceQualifier::Unknown;
111   }
112 }
113 
114 ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual,
115                                            StringRef BaseTypeName) const {
116   if (TypeQual.contains("pipe"))
117     return ValueKind::Pipe;
118 
119   return StringSwitch<ValueKind>(BaseTypeName)
120              .Case("image1d_t", ValueKind::Image)
121              .Case("image1d_array_t", ValueKind::Image)
122              .Case("image1d_buffer_t", ValueKind::Image)
123              .Case("image2d_t", ValueKind::Image)
124              .Case("image2d_array_t", ValueKind::Image)
125              .Case("image2d_array_depth_t", ValueKind::Image)
126              .Case("image2d_array_msaa_t", ValueKind::Image)
127              .Case("image2d_array_msaa_depth_t", ValueKind::Image)
128              .Case("image2d_depth_t", ValueKind::Image)
129              .Case("image2d_msaa_t", ValueKind::Image)
130              .Case("image2d_msaa_depth_t", ValueKind::Image)
131              .Case("image3d_t", ValueKind::Image)
132              .Case("sampler_t", ValueKind::Sampler)
133              .Case("queue_t", ValueKind::Queue)
134              .Default(isa<PointerType>(Ty) ?
135                           (Ty->getPointerAddressSpace() ==
136                            AMDGPUAS::LOCAL_ADDRESS ?
137                            ValueKind::DynamicSharedPointer :
138                            ValueKind::GlobalBuffer) :
139                       ValueKind::ByValue);
140 }
141 
142 std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const {
143   switch (Ty->getTypeID()) {
144   case Type::IntegerTyID: {
145     if (!Signed)
146       return (Twine('u') + getTypeName(Ty, true)).str();
147 
148     auto BitWidth = Ty->getIntegerBitWidth();
149     switch (BitWidth) {
150     case 8:
151       return "char";
152     case 16:
153       return "short";
154     case 32:
155       return "int";
156     case 64:
157       return "long";
158     default:
159       return (Twine('i') + Twine(BitWidth)).str();
160     }
161   }
162   case Type::HalfTyID:
163     return "half";
164   case Type::FloatTyID:
165     return "float";
166   case Type::DoubleTyID:
167     return "double";
168   case Type::FixedVectorTyID: {
169     auto VecTy = cast<FixedVectorType>(Ty);
170     auto ElTy = VecTy->getElementType();
171     auto NumElements = VecTy->getNumElements();
172     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
173   }
174   default:
175     return "unknown";
176   }
177 }
178 
179 std::vector<uint32_t>
180 MetadataStreamerV2::getWorkGroupDimensions(MDNode *Node) const {
181   std::vector<uint32_t> Dims;
182   if (Node->getNumOperands() != 3)
183     return Dims;
184 
185   for (auto &Op : Node->operands())
186     Dims.push_back(mdconst::extract<ConstantInt>(Op)->getZExtValue());
187   return Dims;
188 }
189 
190 Kernel::CodeProps::Metadata
191 MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF,
192                                     const SIProgramInfo &ProgramInfo) const {
193   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
194   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
195   HSAMD::Kernel::CodeProps::Metadata HSACodeProps;
196   const Function &F = MF.getFunction();
197 
198   assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
199          F.getCallingConv() == CallingConv::SPIR_KERNEL);
200 
201   Align MaxKernArgAlign;
202   HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F,
203                                                                MaxKernArgAlign);
204   HSACodeProps.mKernargSegmentAlign =
205     std::max(MaxKernArgAlign, Align(4)).value();
206 
207   HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
208   HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
209   HSACodeProps.mWavefrontSize = STM.getWavefrontSize();
210   HSACodeProps.mNumSGPRs = ProgramInfo.NumSGPR;
211   HSACodeProps.mNumVGPRs = ProgramInfo.NumVGPR;
212   HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize();
213   HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack;
214   HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled();
215   HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
216   HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
217 
218   return HSACodeProps;
219 }
220 
221 Kernel::DebugProps::Metadata
222 MetadataStreamerV2::getHSADebugProps(const MachineFunction &MF,
223                                      const SIProgramInfo &ProgramInfo) const {
224   return HSAMD::Kernel::DebugProps::Metadata();
225 }
226 
227 void MetadataStreamerV2::emitVersion() {
228   auto &Version = HSAMetadata.mVersion;
229 
230   Version.push_back(VersionMajorV2);
231   Version.push_back(VersionMinorV2);
232 }
233 
234 void MetadataStreamerV2::emitPrintf(const Module &Mod) {
235   auto &Printf = HSAMetadata.mPrintf;
236 
237   auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
238   if (!Node)
239     return;
240 
241   for (auto Op : Node->operands())
242     if (Op->getNumOperands())
243       Printf.push_back(
244           std::string(cast<MDString>(Op->getOperand(0))->getString()));
245 }
246 
247 void MetadataStreamerV2::emitKernelLanguage(const Function &Func) {
248   auto &Kernel = HSAMetadata.mKernels.back();
249 
250   // TODO: What about other languages?
251   auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
252   if (!Node || !Node->getNumOperands())
253     return;
254   auto Op0 = Node->getOperand(0);
255   if (Op0->getNumOperands() <= 1)
256     return;
257 
258   Kernel.mLanguage = "OpenCL C";
259   Kernel.mLanguageVersion.push_back(
260       mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
261   Kernel.mLanguageVersion.push_back(
262       mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
263 }
264 
265 void MetadataStreamerV2::emitKernelAttrs(const Function &Func) {
266   auto &Attrs = HSAMetadata.mKernels.back().mAttrs;
267 
268   if (auto Node = Func.getMetadata("reqd_work_group_size"))
269     Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
270   if (auto Node = Func.getMetadata("work_group_size_hint"))
271     Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
272   if (auto Node = Func.getMetadata("vec_type_hint")) {
273     Attrs.mVecTypeHint = getTypeName(
274         cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
275         mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
276   }
277   if (Func.hasFnAttribute("runtime-handle")) {
278     Attrs.mRuntimeHandle =
279         Func.getFnAttribute("runtime-handle").getValueAsString().str();
280   }
281 }
282 
283 void MetadataStreamerV2::emitKernelArgs(const Function &Func,
284                                         const GCNSubtarget &ST) {
285   for (auto &Arg : Func.args())
286     emitKernelArg(Arg);
287 
288   emitHiddenKernelArgs(Func, ST);
289 }
290 
291 void MetadataStreamerV2::emitKernelArg(const Argument &Arg) {
292   auto Func = Arg.getParent();
293   auto ArgNo = Arg.getArgNo();
294   const MDNode *Node;
295 
296   StringRef Name;
297   Node = Func->getMetadata("kernel_arg_name");
298   if (Node && ArgNo < Node->getNumOperands())
299     Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
300   else if (Arg.hasName())
301     Name = Arg.getName();
302 
303   StringRef TypeName;
304   Node = Func->getMetadata("kernel_arg_type");
305   if (Node && ArgNo < Node->getNumOperands())
306     TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
307 
308   StringRef BaseTypeName;
309   Node = Func->getMetadata("kernel_arg_base_type");
310   if (Node && ArgNo < Node->getNumOperands())
311     BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
312 
313   StringRef AccQual;
314   if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
315       Arg.hasNoAliasAttr()) {
316     AccQual = "read_only";
317   } else {
318     Node = Func->getMetadata("kernel_arg_access_qual");
319     if (Node && ArgNo < Node->getNumOperands())
320       AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
321   }
322 
323   StringRef TypeQual;
324   Node = Func->getMetadata("kernel_arg_type_qual");
325   if (Node && ArgNo < Node->getNumOperands())
326     TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
327 
328   const DataLayout &DL = Func->getParent()->getDataLayout();
329 
330   MaybeAlign PointeeAlign;
331   if (auto PtrTy = dyn_cast<PointerType>(Arg.getType())) {
332     if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
333       // FIXME: Should report this for all address spaces
334       PointeeAlign = DL.getValueOrABITypeAlignment(Arg.getParamAlign(),
335                                                    PtrTy->getElementType());
336     }
337   }
338 
339   Type *ArgTy;
340   Align ArgAlign;
341   std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
342 
343   emitKernelArg(DL, ArgTy, ArgAlign,
344                 getValueKind(ArgTy, TypeQual, BaseTypeName), PointeeAlign, Name,
345                 TypeName, BaseTypeName, AccQual, TypeQual);
346 }
347 
348 void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
349                                        Align Alignment, ValueKind ValueKind,
350                                        MaybeAlign PointeeAlign, StringRef Name,
351                                        StringRef TypeName,
352                                        StringRef BaseTypeName,
353                                        StringRef AccQual, StringRef TypeQual) {
354   HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
355   auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
356 
357   Arg.mName = std::string(Name);
358   Arg.mTypeName = std::string(TypeName);
359   Arg.mSize = DL.getTypeAllocSize(Ty);
360   Arg.mAlign = Alignment.value();
361   Arg.mValueKind = ValueKind;
362   Arg.mPointeeAlign = PointeeAlign ? PointeeAlign->value() : 0;
363 
364   if (auto PtrTy = dyn_cast<PointerType>(Ty))
365     Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
366 
367   Arg.mAccQual = getAccessQualifier(AccQual);
368 
369   // TODO: Emit Arg.mActualAccQual.
370 
371   SmallVector<StringRef, 1> SplitTypeQuals;
372   TypeQual.split(SplitTypeQuals, " ", -1, false);
373   for (StringRef Key : SplitTypeQuals) {
374     auto P = StringSwitch<bool*>(Key)
375                  .Case("const",    &Arg.mIsConst)
376                  .Case("restrict", &Arg.mIsRestrict)
377                  .Case("volatile", &Arg.mIsVolatile)
378                  .Case("pipe",     &Arg.mIsPipe)
379                  .Default(nullptr);
380     if (P)
381       *P = true;
382   }
383 }
384 
385 void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
386                                               const GCNSubtarget &ST) {
387   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
388   if (!HiddenArgNumBytes)
389     return;
390 
391   auto &DL = Func.getParent()->getDataLayout();
392   auto Int64Ty = Type::getInt64Ty(Func.getContext());
393 
394   if (HiddenArgNumBytes >= 8)
395     emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetX);
396   if (HiddenArgNumBytes >= 16)
397     emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetY);
398   if (HiddenArgNumBytes >= 24)
399     emitKernelArg(DL, Int64Ty, Align(8), ValueKind::HiddenGlobalOffsetZ);
400 
401   auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
402                                       AMDGPUAS::GLOBAL_ADDRESS);
403 
404   // Emit "printf buffer" argument if printf is used, otherwise emit dummy
405   // "none" argument.
406   if (HiddenArgNumBytes >= 32) {
407     if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
408       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer);
409     else if (Func.getParent()->getFunction("__ockl_hostcall_internal")) {
410       // The printf runtime binding pass should have ensured that hostcall and
411       // printf are not used in the same module.
412       assert(!Func.getParent()->getNamedMetadata("llvm.printf.fmts"));
413       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer);
414     } else
415       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
416   }
417 
418   // Emit "default queue" and "completion action" arguments if enqueue kernel is
419   // used, otherwise emit dummy "none" arguments.
420   if (HiddenArgNumBytes >= 48) {
421     if (Func.hasFnAttribute("calls-enqueue-kernel")) {
422       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenDefaultQueue);
423       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenCompletionAction);
424     } else {
425       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
426       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
427     }
428   }
429 
430   // Emit the pointer argument for multi-grid object.
431   if (HiddenArgNumBytes >= 56)
432     emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenMultiGridSyncArg);
433 }
434 
435 bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
436   return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
437 }
438 
439 void MetadataStreamerV2::begin(const Module &Mod,
440                                const IsaInfo::AMDGPUTargetID &TargetID) {
441   emitVersion();
442   emitPrintf(Mod);
443 }
444 
445 void MetadataStreamerV2::end() {
446   std::string HSAMetadataString;
447   if (toString(HSAMetadata, HSAMetadataString))
448     return;
449 
450   if (DumpHSAMetadata)
451     dump(HSAMetadataString);
452   if (VerifyHSAMetadata)
453     verify(HSAMetadataString);
454 }
455 
456 void MetadataStreamerV2::emitKernel(const MachineFunction &MF,
457                                     const SIProgramInfo &ProgramInfo) {
458   auto &Func = MF.getFunction();
459   if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
460     return;
461 
462   auto CodeProps = getHSACodeProps(MF, ProgramInfo);
463   auto DebugProps = getHSADebugProps(MF, ProgramInfo);
464 
465   HSAMetadata.mKernels.push_back(Kernel::Metadata());
466   auto &Kernel = HSAMetadata.mKernels.back();
467 
468   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
469   Kernel.mName = std::string(Func.getName());
470   Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str();
471   emitKernelLanguage(Func);
472   emitKernelAttrs(Func);
473   emitKernelArgs(Func, ST);
474   HSAMetadata.mKernels.back().mCodeProps = CodeProps;
475   HSAMetadata.mKernels.back().mDebugProps = DebugProps;
476 }
477 
478 //===----------------------------------------------------------------------===//
479 // HSAMetadataStreamerV3
480 //===----------------------------------------------------------------------===//
481 
482 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
483   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
484 }
485 
486 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
487   errs() << "AMDGPU HSA Metadata Parser Test: ";
488 
489   msgpack::Document FromHSAMetadataString;
490 
491   if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
492     errs() << "FAIL\n";
493     return;
494   }
495 
496   std::string ToHSAMetadataString;
497   raw_string_ostream StrOS(ToHSAMetadataString);
498   FromHSAMetadataString.toYAML(StrOS);
499 
500   errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
501   if (HSAMetadataString != ToHSAMetadataString) {
502     errs() << "Original input: " << HSAMetadataString << '\n'
503            << "Produced output: " << StrOS.str() << '\n';
504   }
505 }
506 
507 Optional<StringRef>
508 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
509   return StringSwitch<Optional<StringRef>>(AccQual)
510       .Case("read_only", StringRef("read_only"))
511       .Case("write_only", StringRef("write_only"))
512       .Case("read_write", StringRef("read_write"))
513       .Default(None);
514 }
515 
516 Optional<StringRef>
517 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
518   switch (AddressSpace) {
519   case AMDGPUAS::PRIVATE_ADDRESS:
520     return StringRef("private");
521   case AMDGPUAS::GLOBAL_ADDRESS:
522     return StringRef("global");
523   case AMDGPUAS::CONSTANT_ADDRESS:
524     return StringRef("constant");
525   case AMDGPUAS::LOCAL_ADDRESS:
526     return StringRef("local");
527   case AMDGPUAS::FLAT_ADDRESS:
528     return StringRef("generic");
529   case AMDGPUAS::REGION_ADDRESS:
530     return StringRef("region");
531   default:
532     return None;
533   }
534 }
535 
536 StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
537                                            StringRef BaseTypeName) const {
538   if (TypeQual.contains("pipe"))
539     return "pipe";
540 
541   return StringSwitch<StringRef>(BaseTypeName)
542       .Case("image1d_t", "image")
543       .Case("image1d_array_t", "image")
544       .Case("image1d_buffer_t", "image")
545       .Case("image2d_t", "image")
546       .Case("image2d_array_t", "image")
547       .Case("image2d_array_depth_t", "image")
548       .Case("image2d_array_msaa_t", "image")
549       .Case("image2d_array_msaa_depth_t", "image")
550       .Case("image2d_depth_t", "image")
551       .Case("image2d_msaa_t", "image")
552       .Case("image2d_msaa_depth_t", "image")
553       .Case("image3d_t", "image")
554       .Case("sampler_t", "sampler")
555       .Case("queue_t", "queue")
556       .Default(isa<PointerType>(Ty)
557                    ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
558                           ? "dynamic_shared_pointer"
559                           : "global_buffer")
560                    : "by_value");
561 }
562 
563 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
564   switch (Ty->getTypeID()) {
565   case Type::IntegerTyID: {
566     if (!Signed)
567       return (Twine('u') + getTypeName(Ty, true)).str();
568 
569     auto BitWidth = Ty->getIntegerBitWidth();
570     switch (BitWidth) {
571     case 8:
572       return "char";
573     case 16:
574       return "short";
575     case 32:
576       return "int";
577     case 64:
578       return "long";
579     default:
580       return (Twine('i') + Twine(BitWidth)).str();
581     }
582   }
583   case Type::HalfTyID:
584     return "half";
585   case Type::FloatTyID:
586     return "float";
587   case Type::DoubleTyID:
588     return "double";
589   case Type::FixedVectorTyID: {
590     auto VecTy = cast<FixedVectorType>(Ty);
591     auto ElTy = VecTy->getElementType();
592     auto NumElements = VecTy->getNumElements();
593     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
594   }
595   default:
596     return "unknown";
597   }
598 }
599 
600 msgpack::ArrayDocNode
601 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
602   auto Dims = HSAMetadataDoc->getArrayNode();
603   if (Node->getNumOperands() != 3)
604     return Dims;
605 
606   for (auto &Op : Node->operands())
607     Dims.push_back(Dims.getDocument()->getNode(
608         uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
609   return Dims;
610 }
611 
612 void MetadataStreamerV3::emitVersion() {
613   auto Version = HSAMetadataDoc->getArrayNode();
614   Version.push_back(Version.getDocument()->getNode(VersionMajorV3));
615   Version.push_back(Version.getDocument()->getNode(VersionMinorV3));
616   getRootMetadata("amdhsa.version") = Version;
617 }
618 
619 void MetadataStreamerV3::emitPrintf(const Module &Mod) {
620   auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
621   if (!Node)
622     return;
623 
624   auto Printf = HSAMetadataDoc->getArrayNode();
625   for (auto Op : Node->operands())
626     if (Op->getNumOperands())
627       Printf.push_back(Printf.getDocument()->getNode(
628           cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
629   getRootMetadata("amdhsa.printf") = Printf;
630 }
631 
632 void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
633                                             msgpack::MapDocNode Kern) {
634   // TODO: What about other languages?
635   auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
636   if (!Node || !Node->getNumOperands())
637     return;
638   auto Op0 = Node->getOperand(0);
639   if (Op0->getNumOperands() <= 1)
640     return;
641 
642   Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
643   auto LanguageVersion = Kern.getDocument()->getArrayNode();
644   LanguageVersion.push_back(Kern.getDocument()->getNode(
645       mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
646   LanguageVersion.push_back(Kern.getDocument()->getNode(
647       mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
648   Kern[".language_version"] = LanguageVersion;
649 }
650 
651 void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
652                                          msgpack::MapDocNode Kern) {
653 
654   if (auto Node = Func.getMetadata("reqd_work_group_size"))
655     Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
656   if (auto Node = Func.getMetadata("work_group_size_hint"))
657     Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
658   if (auto Node = Func.getMetadata("vec_type_hint")) {
659     Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
660         getTypeName(
661             cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
662             mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
663         /*Copy=*/true);
664   }
665   if (Func.hasFnAttribute("runtime-handle")) {
666     Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
667         Func.getFnAttribute("runtime-handle").getValueAsString().str(),
668         /*Copy=*/true);
669   }
670   if (Func.hasFnAttribute("device-init"))
671     Kern[".kind"] = Kern.getDocument()->getNode("init");
672   else if (Func.hasFnAttribute("device-fini"))
673     Kern[".kind"] = Kern.getDocument()->getNode("fini");
674 }
675 
676 void MetadataStreamerV3::emitKernelArgs(const Function &Func,
677                                         const GCNSubtarget &ST,
678                                         msgpack::MapDocNode Kern) {
679   unsigned Offset = 0;
680   auto Args = HSAMetadataDoc->getArrayNode();
681   for (auto &Arg : Func.args())
682     emitKernelArg(Arg, Offset, Args);
683 
684   emitHiddenKernelArgs(Func, ST, Offset, Args);
685 
686   Kern[".args"] = Args;
687 }
688 
689 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
690                                        msgpack::ArrayDocNode Args) {
691   auto Func = Arg.getParent();
692   auto ArgNo = Arg.getArgNo();
693   const MDNode *Node;
694 
695   StringRef Name;
696   Node = Func->getMetadata("kernel_arg_name");
697   if (Node && ArgNo < Node->getNumOperands())
698     Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
699   else if (Arg.hasName())
700     Name = Arg.getName();
701 
702   StringRef TypeName;
703   Node = Func->getMetadata("kernel_arg_type");
704   if (Node && ArgNo < Node->getNumOperands())
705     TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
706 
707   StringRef BaseTypeName;
708   Node = Func->getMetadata("kernel_arg_base_type");
709   if (Node && ArgNo < Node->getNumOperands())
710     BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
711 
712   StringRef AccQual;
713   if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
714       Arg.hasNoAliasAttr()) {
715     AccQual = "read_only";
716   } else {
717     Node = Func->getMetadata("kernel_arg_access_qual");
718     if (Node && ArgNo < Node->getNumOperands())
719       AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
720   }
721 
722   StringRef TypeQual;
723   Node = Func->getMetadata("kernel_arg_type_qual");
724   if (Node && ArgNo < Node->getNumOperands())
725     TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
726 
727   const DataLayout &DL = Func->getParent()->getDataLayout();
728 
729   MaybeAlign PointeeAlign;
730   Type *Ty = Arg.hasByRefAttr() ? Arg.getParamByRefType() : Arg.getType();
731 
732   // FIXME: Need to distinguish in memory alignment from pointer alignment.
733   if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
734     if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
735       PointeeAlign = DL.getValueOrABITypeAlignment(Arg.getParamAlign(),
736                                                    PtrTy->getElementType());
737     }
738   }
739 
740   // There's no distinction between byval aggregates and raw aggregates.
741   Type *ArgTy;
742   Align ArgAlign;
743   std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
744 
745   emitKernelArg(DL, ArgTy, ArgAlign,
746                 getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args,
747                 PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual);
748 }
749 
750 void MetadataStreamerV3::emitKernelArg(
751     const DataLayout &DL, Type *Ty, Align Alignment, StringRef ValueKind,
752     unsigned &Offset, msgpack::ArrayDocNode Args, MaybeAlign PointeeAlign,
753     StringRef Name, StringRef TypeName, StringRef BaseTypeName,
754     StringRef AccQual, StringRef TypeQual) {
755   auto Arg = Args.getDocument()->getMapNode();
756 
757   if (!Name.empty())
758     Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
759   if (!TypeName.empty())
760     Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
761   auto Size = DL.getTypeAllocSize(Ty);
762   Arg[".size"] = Arg.getDocument()->getNode(Size);
763   Offset = alignTo(Offset, Alignment);
764   Arg[".offset"] = Arg.getDocument()->getNode(Offset);
765   Offset += Size;
766   Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
767   if (PointeeAlign)
768     Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
769 
770   if (auto PtrTy = dyn_cast<PointerType>(Ty))
771     if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
772       Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
773 
774   if (auto AQ = getAccessQualifier(AccQual))
775     Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
776 
777   // TODO: Emit Arg[".actual_access"].
778 
779   SmallVector<StringRef, 1> SplitTypeQuals;
780   TypeQual.split(SplitTypeQuals, " ", -1, false);
781   for (StringRef Key : SplitTypeQuals) {
782     if (Key == "const")
783       Arg[".is_const"] = Arg.getDocument()->getNode(true);
784     else if (Key == "restrict")
785       Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
786     else if (Key == "volatile")
787       Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
788     else if (Key == "pipe")
789       Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
790   }
791 
792   Args.push_back(Arg);
793 }
794 
795 void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
796                                               const GCNSubtarget &ST,
797                                               unsigned &Offset,
798                                               msgpack::ArrayDocNode Args) {
799   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
800   if (!HiddenArgNumBytes)
801     return;
802 
803   const Module *M = Func.getParent();
804   auto &DL = M->getDataLayout();
805   auto Int64Ty = Type::getInt64Ty(Func.getContext());
806 
807   if (HiddenArgNumBytes >= 8)
808     emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_x", Offset,
809                   Args);
810   if (HiddenArgNumBytes >= 16)
811     emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_y", Offset,
812                   Args);
813   if (HiddenArgNumBytes >= 24)
814     emitKernelArg(DL, Int64Ty, Align(8), "hidden_global_offset_z", Offset,
815                   Args);
816 
817   auto Int8PtrTy =
818       Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
819 
820   // Emit "printf buffer" argument if printf is used, emit "hostcall buffer"
821   // if "hostcall" module flag is set, otherwise emit dummy "none" argument.
822   if (HiddenArgNumBytes >= 32) {
823     if (M->getNamedMetadata("llvm.printf.fmts"))
824       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
825                     Args);
826     else if (M->getModuleFlag("amdgpu_hostcall")) {
827       // The printf runtime binding pass should have ensured that hostcall and
828       // printf are not used in the same module.
829       assert(!M->getNamedMetadata("llvm.printf.fmts"));
830       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
831                     Args);
832     } else
833       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
834   }
835 
836   // Emit "default queue" and "completion action" arguments if enqueue kernel is
837   // used, otherwise emit dummy "none" arguments.
838   if (HiddenArgNumBytes >= 48) {
839     if (Func.hasFnAttribute("calls-enqueue-kernel")) {
840       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_default_queue", Offset,
841                     Args);
842       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_completion_action", Offset,
843                     Args);
844     } else {
845       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
846       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
847     }
848   }
849 
850   // Emit the pointer argument for multi-grid object.
851   if (HiddenArgNumBytes >= 56)
852     emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset,
853                   Args);
854 }
855 
856 msgpack::MapDocNode
857 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
858                                       const SIProgramInfo &ProgramInfo) const {
859   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
860   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
861   const Function &F = MF.getFunction();
862 
863   auto Kern = HSAMetadataDoc->getMapNode();
864 
865   Align MaxKernArgAlign;
866   Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
867       STM.getKernArgSegmentSize(F, MaxKernArgAlign));
868   Kern[".group_segment_fixed_size"] =
869       Kern.getDocument()->getNode(ProgramInfo.LDSSize);
870   Kern[".private_segment_fixed_size"] =
871       Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
872 
873   // FIXME: The metadata treats the minimum as 16?
874   Kern[".kernarg_segment_align"] =
875       Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
876   Kern[".wavefront_size"] =
877       Kern.getDocument()->getNode(STM.getWavefrontSize());
878   Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
879   Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
880   Kern[".max_flat_workgroup_size"] =
881       Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
882   Kern[".sgpr_spill_count"] =
883       Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
884   Kern[".vgpr_spill_count"] =
885       Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
886 
887   return Kern;
888 }
889 
890 bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
891   return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
892 }
893 
894 void MetadataStreamerV3::begin(const Module &Mod,
895                                const IsaInfo::AMDGPUTargetID &TargetID) {
896   emitVersion();
897   emitPrintf(Mod);
898   getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
899 }
900 
901 void MetadataStreamerV3::end() {
902   std::string HSAMetadataString;
903   raw_string_ostream StrOS(HSAMetadataString);
904   HSAMetadataDoc->toYAML(StrOS);
905 
906   if (DumpHSAMetadata)
907     dump(StrOS.str());
908   if (VerifyHSAMetadata)
909     verify(StrOS.str());
910 }
911 
912 void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
913                                     const SIProgramInfo &ProgramInfo) {
914   auto &Func = MF.getFunction();
915   auto Kern = getHSAKernelProps(MF, ProgramInfo);
916   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
917 
918   assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
919          Func.getCallingConv() == CallingConv::SPIR_KERNEL);
920 
921   auto Kernels =
922       getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
923 
924   {
925     Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
926     Kern[".symbol"] = Kern.getDocument()->getNode(
927         (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
928     emitKernelLanguage(Func, Kern);
929     emitKernelAttrs(Func, Kern);
930     emitKernelArgs(Func, ST, Kern);
931   }
932 
933   Kernels.push_back(Kern);
934 }
935 
936 //===----------------------------------------------------------------------===//
937 // HSAMetadataStreamerV4
938 //===----------------------------------------------------------------------===//
939 
940 void MetadataStreamerV4::emitVersion() {
941   auto Version = HSAMetadataDoc->getArrayNode();
942   Version.push_back(Version.getDocument()->getNode(VersionMajorV4));
943   Version.push_back(Version.getDocument()->getNode(VersionMinorV4));
944   getRootMetadata("amdhsa.version") = Version;
945 }
946 
947 void MetadataStreamerV4::emitTargetID(const IsaInfo::AMDGPUTargetID &TargetID) {
948   getRootMetadata("amdhsa.target") =
949       HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true);
950 }
951 
952 void MetadataStreamerV4::begin(const Module &Mod,
953                                const IsaInfo::AMDGPUTargetID &TargetID) {
954   emitVersion();
955   emitTargetID(TargetID);
956   emitPrintf(Mod);
957   getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
958 }
959 
960 } // end namespace HSAMD
961 } // end namespace AMDGPU
962 } // end namespace llvm
963