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