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