LLVM  10.0.0svn
AMDGPUHSAMetadataStreamer.cpp
Go to the documentation of this file.
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 
16 #include "AMDGPU.h"
17 #include "AMDGPUSubtarget.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"
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(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 
70 MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
71  if (AccQual.empty())
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 
82 MetadataStreamerV2::getAddressSpaceQualifier(
83  unsigned AddressSpace) const {
84  switch (AddressSpace) {
97  default:
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() ==
128 }
129 
130 ValueType MetadataStreamerV2::getValueType(Type *Ty, StringRef TypeName) const {
131  switch (Ty->getTypeID()) {
132  case Type::IntegerTyID: {
133  auto Signed = !TypeName.startswith("u");
134  switch (Ty->getIntegerBitWidth()) {
135  case 8:
137  case 16:
139  case 32:
141  case 64:
143  default:
144  return ValueType::Struct;
145  }
146  }
147  case Type::HalfTyID:
148  return ValueType::F16;
149  case Type::FloatTyID:
150  return ValueType::F32;
151  case Type::DoubleTyID:
152  return ValueType::F64;
153  case Type::PointerTyID:
154  return getValueType(Ty->getPointerElementType(), TypeName);
155  case Type::VectorTyID:
156  return getValueType(Ty->getVectorElementType(), TypeName);
157  default:
158  return ValueType::Struct;
159  }
160 }
161 
162 std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const {
163  switch (Ty->getTypeID()) {
164  case Type::IntegerTyID: {
165  if (!Signed)
166  return (Twine('u') + getTypeName(Ty, true)).str();
167 
168  auto BitWidth = Ty->getIntegerBitWidth();
169  switch (BitWidth) {
170  case 8:
171  return "char";
172  case 16:
173  return "short";
174  case 32:
175  return "int";
176  case 64:
177  return "long";
178  default:
179  return (Twine('i') + Twine(BitWidth)).str();
180  }
181  }
182  case Type::HalfTyID:
183  return "half";
184  case Type::FloatTyID:
185  return "float";
186  case Type::DoubleTyID:
187  return "double";
188  case Type::VectorTyID: {
189  auto VecTy = cast<VectorType>(Ty);
190  auto ElTy = VecTy->getElementType();
191  auto NumElements = VecTy->getVectorNumElements();
192  return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
193  }
194  default:
195  return "unknown";
196  }
197 }
198 
199 std::vector<uint32_t>
200 MetadataStreamerV2::getWorkGroupDimensions(MDNode *Node) const {
201  std::vector<uint32_t> Dims;
202  if (Node->getNumOperands() != 3)
203  return Dims;
204 
205  for (auto &Op : Node->operands())
206  Dims.push_back(mdconst::extract<ConstantInt>(Op)->getZExtValue());
207  return Dims;
208 }
209 
210 Kernel::CodeProps::Metadata
211 MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF,
212  const SIProgramInfo &ProgramInfo) const {
213  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
214  const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
215  HSAMD::Kernel::CodeProps::Metadata HSACodeProps;
216  const Function &F = MF.getFunction();
217 
218  assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
219  F.getCallingConv() == CallingConv::SPIR_KERNEL);
220 
221  Align MaxKernArgAlign;
222  HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F,
223  MaxKernArgAlign);
224  HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
225  HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
226  HSACodeProps.mKernargSegmentAlign =
227  std::max(MaxKernArgAlign, Align(4)).value();
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  return HSAMD::Kernel::DebugProps::Metadata();
244 }
245 
246 void MetadataStreamerV2::emitVersion() {
247  auto &Version = HSAMetadata.mVersion;
248 
249  Version.push_back(VersionMajor);
250  Version.push_back(VersionMinor);
251 }
252 
253 void MetadataStreamerV2::emitPrintf(const Module &Mod) {
254  auto &Printf = HSAMetadata.mPrintf;
255 
256  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
257  if (!Node)
258  return;
259 
260  for (auto Op : Node->operands())
261  if (Op->getNumOperands())
262  Printf.push_back(cast<MDString>(Op->getOperand(0))->getString());
263 }
264 
265 void MetadataStreamerV2::emitKernelLanguage(const Function &Func) {
266  auto &Kernel = HSAMetadata.mKernels.back();
267 
268  // TODO: What about other languages?
269  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
270  if (!Node || !Node->getNumOperands())
271  return;
272  auto Op0 = Node->getOperand(0);
273  if (Op0->getNumOperands() <= 1)
274  return;
275 
276  Kernel.mLanguage = "OpenCL C";
277  Kernel.mLanguageVersion.push_back(
278  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
279  Kernel.mLanguageVersion.push_back(
280  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
281 }
282 
283 void MetadataStreamerV2::emitKernelAttrs(const Function &Func) {
284  auto &Attrs = HSAMetadata.mKernels.back().mAttrs;
285 
286  if (auto Node = Func.getMetadata("reqd_work_group_size"))
287  Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
288  if (auto Node = Func.getMetadata("work_group_size_hint"))
289  Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
290  if (auto Node = Func.getMetadata("vec_type_hint")) {
291  Attrs.mVecTypeHint = getTypeName(
292  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
293  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
294  }
295  if (Func.hasFnAttribute("runtime-handle")) {
296  Attrs.mRuntimeHandle =
297  Func.getFnAttribute("runtime-handle").getValueAsString().str();
298  }
299 }
300 
301 void MetadataStreamerV2::emitKernelArgs(const Function &Func) {
302  for (auto &Arg : Func.args())
303  emitKernelArg(Arg);
304 
305  emitHiddenKernelArgs(Func);
306 }
307 
308 void MetadataStreamerV2::emitKernelArg(const Argument &Arg) {
309  auto Func = Arg.getParent();
310  auto ArgNo = Arg.getArgNo();
311  const MDNode *Node;
312 
313  StringRef Name;
314  Node = Func->getMetadata("kernel_arg_name");
315  if (Node && ArgNo < Node->getNumOperands())
316  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
317  else if (Arg.hasName())
318  Name = Arg.getName();
319 
320  StringRef TypeName;
321  Node = Func->getMetadata("kernel_arg_type");
322  if (Node && ArgNo < Node->getNumOperands())
323  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
324 
325  StringRef BaseTypeName;
326  Node = Func->getMetadata("kernel_arg_base_type");
327  if (Node && ArgNo < Node->getNumOperands())
328  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
329 
330  StringRef AccQual;
331  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
332  Arg.hasNoAliasAttr()) {
333  AccQual = "read_only";
334  } else {
335  Node = Func->getMetadata("kernel_arg_access_qual");
336  if (Node && ArgNo < Node->getNumOperands())
337  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
338  }
339 
340  StringRef TypeQual;
341  Node = Func->getMetadata("kernel_arg_type_qual");
342  if (Node && ArgNo < Node->getNumOperands())
343  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
344 
345  Type *Ty = Arg.getType();
346  const DataLayout &DL = Func->getParent()->getDataLayout();
347 
348  unsigned PointeeAlign = 0;
349  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
350  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
351  PointeeAlign = Arg.getParamAlignment();
352  if (PointeeAlign == 0)
353  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
354  }
355  }
356 
357  emitKernelArg(DL, Ty, getValueKind(Arg.getType(), TypeQual, BaseTypeName),
358  PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual);
359 }
360 
361 void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
363  unsigned PointeeAlign, StringRef Name,
364  StringRef TypeName,
365  StringRef BaseTypeName,
366  StringRef AccQual, StringRef TypeQual) {
367  HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
368  auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
369 
370  Arg.mName = Name;
371  Arg.mTypeName = TypeName;
372  Arg.mSize = DL.getTypeAllocSize(Ty);
373  Arg.mAlign = DL.getABITypeAlignment(Ty);
374  Arg.mValueKind = ValueKind;
375  Arg.mValueType = getValueType(Ty, BaseTypeName);
376  Arg.mPointeeAlign = PointeeAlign;
377 
378  if (auto PtrTy = dyn_cast<PointerType>(Ty))
379  Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
380 
381  Arg.mAccQual = getAccessQualifier(AccQual);
382 
383  // TODO: Emit Arg.mActualAccQual.
384 
385  SmallVector<StringRef, 1> SplitTypeQuals;
386  TypeQual.split(SplitTypeQuals, " ", -1, false);
387  for (StringRef Key : SplitTypeQuals) {
388  auto P = StringSwitch<bool*>(Key)
389  .Case("const", &Arg.mIsConst)
390  .Case("restrict", &Arg.mIsRestrict)
391  .Case("volatile", &Arg.mIsVolatile)
392  .Case("pipe", &Arg.mIsPipe)
393  .Default(nullptr);
394  if (P)
395  *P = true;
396  }
397 }
398 
399 void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func) {
400  int HiddenArgNumBytes =
401  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
402 
403  if (!HiddenArgNumBytes)
404  return;
405 
406  auto &DL = Func.getParent()->getDataLayout();
407  auto Int64Ty = Type::getInt64Ty(Func.getContext());
408 
409  if (HiddenArgNumBytes >= 8)
410  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetX);
411  if (HiddenArgNumBytes >= 16)
412  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetY);
413  if (HiddenArgNumBytes >= 24)
414  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetZ);
415 
416  auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
418 
419  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
420  // "none" argument.
421  if (HiddenArgNumBytes >= 32) {
422  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
423  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenPrintfBuffer);
424  else
425  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
426  }
427 
428  // Emit "default queue" and "completion action" arguments if enqueue kernel is
429  // used, otherwise emit dummy "none" arguments.
430  if (HiddenArgNumBytes >= 48) {
431  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
432  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenDefaultQueue);
433  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenCompletionAction);
434  } else {
435  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
436  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
437  }
438  }
439 
440  // Emit the pointer argument for multi-grid object.
441  if (HiddenArgNumBytes >= 56)
442  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenMultiGridSyncArg);
443 }
444 
446  return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
447 }
448 
450  emitVersion();
451  emitPrintf(Mod);
452 }
453 
455  std::string HSAMetadataString;
456  if (toString(HSAMetadata, HSAMetadataString))
457  return;
458 
459  if (DumpHSAMetadata)
460  dump(HSAMetadataString);
461  if (VerifyHSAMetadata)
462  verify(HSAMetadataString);
463 }
464 
466  const SIProgramInfo &ProgramInfo) {
467  auto &Func = MF.getFunction();
468  if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
469  return;
470 
471  auto CodeProps = getHSACodeProps(MF, ProgramInfo);
472  auto DebugProps = getHSADebugProps(MF, ProgramInfo);
473 
474  HSAMetadata.mKernels.push_back(Kernel::Metadata());
475  auto &Kernel = HSAMetadata.mKernels.back();
476 
477  Kernel.mName = Func.getName();
478  Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str();
479  emitKernelLanguage(Func);
480  emitKernelAttrs(Func);
481  emitKernelArgs(Func);
482  HSAMetadata.mKernels.back().mCodeProps = CodeProps;
483  HSAMetadata.mKernels.back().mDebugProps = DebugProps;
484 }
485 
486 //===----------------------------------------------------------------------===//
487 // HSAMetadataStreamerV3
488 //===----------------------------------------------------------------------===//
489 
490 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
491  errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
492 }
493 
494 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
495  errs() << "AMDGPU HSA Metadata Parser Test: ";
496 
497  msgpack::Document FromHSAMetadataString;
498 
499  if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
500  errs() << "FAIL\n";
501  return;
502  }
503 
504  std::string ToHSAMetadataString;
505  raw_string_ostream StrOS(ToHSAMetadataString);
506  FromHSAMetadataString.toYAML(StrOS);
507 
508  errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
509  if (HSAMetadataString != ToHSAMetadataString) {
510  errs() << "Original input: " << HSAMetadataString << '\n'
511  << "Produced output: " << StrOS.str() << '\n';
512  }
513 }
514 
516 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
518  .Case("read_only", StringRef("read_only"))
519  .Case("write_only", StringRef("write_only"))
520  .Case("read_write", StringRef("read_write"))
521  .Default(None);
522 }
523 
525 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
526  switch (AddressSpace) {
528  return StringRef("private");
530  return StringRef("global");
532  return StringRef("constant");
534  return StringRef("local");
536  return StringRef("generic");
538  return StringRef("region");
539  default:
540  return None;
541  }
542 }
543 
544 StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
545  StringRef BaseTypeName) const {
546  if (TypeQual.find("pipe") != StringRef::npos)
547  return "pipe";
548 
549  return StringSwitch<StringRef>(BaseTypeName)
550  .Case("image1d_t", "image")
551  .Case("image1d_array_t", "image")
552  .Case("image1d_buffer_t", "image")
553  .Case("image2d_t", "image")
554  .Case("image2d_array_t", "image")
555  .Case("image2d_array_depth_t", "image")
556  .Case("image2d_array_msaa_t", "image")
557  .Case("image2d_array_msaa_depth_t", "image")
558  .Case("image2d_depth_t", "image")
559  .Case("image2d_msaa_t", "image")
560  .Case("image2d_msaa_depth_t", "image")
561  .Case("image3d_t", "image")
562  .Case("sampler_t", "sampler")
563  .Case("queue_t", "queue")
564  .Default(isa<PointerType>(Ty)
566  ? "dynamic_shared_pointer"
567  : "global_buffer")
568  : "by_value");
569 }
570 
571 StringRef MetadataStreamerV3::getValueType(Type *Ty, StringRef TypeName) const {
572  switch (Ty->getTypeID()) {
573  case Type::IntegerTyID: {
574  auto Signed = !TypeName.startswith("u");
575  switch (Ty->getIntegerBitWidth()) {
576  case 8:
577  return Signed ? "i8" : "u8";
578  case 16:
579  return Signed ? "i16" : "u16";
580  case 32:
581  return Signed ? "i32" : "u32";
582  case 64:
583  return Signed ? "i64" : "u64";
584  default:
585  return "struct";
586  }
587  }
588  case Type::HalfTyID:
589  return "f16";
590  case Type::FloatTyID:
591  return "f32";
592  case Type::DoubleTyID:
593  return "f64";
594  case Type::PointerTyID:
595  return getValueType(Ty->getPointerElementType(), TypeName);
596  case Type::VectorTyID:
597  return getValueType(Ty->getVectorElementType(), TypeName);
598  default:
599  return "struct";
600  }
601 }
602 
603 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
604  switch (Ty->getTypeID()) {
605  case Type::IntegerTyID: {
606  if (!Signed)
607  return (Twine('u') + getTypeName(Ty, true)).str();
608 
609  auto BitWidth = Ty->getIntegerBitWidth();
610  switch (BitWidth) {
611  case 8:
612  return "char";
613  case 16:
614  return "short";
615  case 32:
616  return "int";
617  case 64:
618  return "long";
619  default:
620  return (Twine('i') + Twine(BitWidth)).str();
621  }
622  }
623  case Type::HalfTyID:
624  return "half";
625  case Type::FloatTyID:
626  return "float";
627  case Type::DoubleTyID:
628  return "double";
629  case Type::VectorTyID: {
630  auto VecTy = cast<VectorType>(Ty);
631  auto ElTy = VecTy->getElementType();
632  auto NumElements = VecTy->getVectorNumElements();
633  return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
634  }
635  default:
636  return "unknown";
637  }
638 }
639 
641 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
642  auto Dims = HSAMetadataDoc->getArrayNode();
643  if (Node->getNumOperands() != 3)
644  return Dims;
645 
646  for (auto &Op : Node->operands())
647  Dims.push_back(Dims.getDocument()->getNode(
648  uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
649  return Dims;
650 }
651 
652 void MetadataStreamerV3::emitVersion() {
653  auto Version = HSAMetadataDoc->getArrayNode();
654  Version.push_back(Version.getDocument()->getNode(VersionMajor));
655  Version.push_back(Version.getDocument()->getNode(VersionMinor));
656  getRootMetadata("amdhsa.version") = Version;
657 }
658 
659 void MetadataStreamerV3::emitPrintf(const Module &Mod) {
660  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
661  if (!Node)
662  return;
663 
664  auto Printf = HSAMetadataDoc->getArrayNode();
665  for (auto Op : Node->operands())
666  if (Op->getNumOperands())
667  Printf.push_back(Printf.getDocument()->getNode(
668  cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
669  getRootMetadata("amdhsa.printf") = Printf;
670 }
671 
672 void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
673  msgpack::MapDocNode Kern) {
674  // TODO: What about other languages?
675  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
676  if (!Node || !Node->getNumOperands())
677  return;
678  auto Op0 = Node->getOperand(0);
679  if (Op0->getNumOperands() <= 1)
680  return;
681 
682  Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
683  auto LanguageVersion = Kern.getDocument()->getArrayNode();
684  LanguageVersion.push_back(Kern.getDocument()->getNode(
685  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
686  LanguageVersion.push_back(Kern.getDocument()->getNode(
687  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
688  Kern[".language_version"] = LanguageVersion;
689 }
690 
691 void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
692  msgpack::MapDocNode Kern) {
693 
694  if (auto Node = Func.getMetadata("reqd_work_group_size"))
695  Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
696  if (auto Node = Func.getMetadata("work_group_size_hint"))
697  Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
698  if (auto Node = Func.getMetadata("vec_type_hint")) {
699  Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
700  getTypeName(
701  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
702  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
703  /*Copy=*/true);
704  }
705  if (Func.hasFnAttribute("runtime-handle")) {
706  Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
707  Func.getFnAttribute("runtime-handle").getValueAsString().str(),
708  /*Copy=*/true);
709  }
710 }
711 
712 void MetadataStreamerV3::emitKernelArgs(const Function &Func,
713  msgpack::MapDocNode Kern) {
714  unsigned Offset = 0;
715  auto Args = HSAMetadataDoc->getArrayNode();
716  for (auto &Arg : Func.args())
717  emitKernelArg(Arg, Offset, Args);
718 
719  emitHiddenKernelArgs(Func, Offset, Args);
720 
721  Kern[".args"] = Args;
722 }
723 
724 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
726  auto Func = Arg.getParent();
727  auto ArgNo = Arg.getArgNo();
728  const MDNode *Node;
729 
730  StringRef Name;
731  Node = Func->getMetadata("kernel_arg_name");
732  if (Node && ArgNo < Node->getNumOperands())
733  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
734  else if (Arg.hasName())
735  Name = Arg.getName();
736 
738  Node = Func->getMetadata("kernel_arg_type");
739  if (Node && ArgNo < Node->getNumOperands())
740  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
741 
742  StringRef BaseTypeName;
743  Node = Func->getMetadata("kernel_arg_base_type");
744  if (Node && ArgNo < Node->getNumOperands())
745  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
746 
748  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
749  Arg.hasNoAliasAttr()) {
750  AccQual = "read_only";
751  } else {
752  Node = Func->getMetadata("kernel_arg_access_qual");
753  if (Node && ArgNo < Node->getNumOperands())
754  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
755  }
756 
757  StringRef TypeQual;
758  Node = Func->getMetadata("kernel_arg_type_qual");
759  if (Node && ArgNo < Node->getNumOperands())
760  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
761 
762  Type *Ty = Arg.getType();
763  const DataLayout &DL = Func->getParent()->getDataLayout();
764 
765  unsigned PointeeAlign = 0;
766  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
767  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
768  PointeeAlign = Arg.getParamAlignment();
769  if (PointeeAlign == 0)
770  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
771  }
772  }
773 
774  emitKernelArg(Func->getParent()->getDataLayout(), Arg.getType(),
775  getValueKind(Arg.getType(), TypeQual, BaseTypeName), Offset,
776  Args, PointeeAlign, Name, TypeName, BaseTypeName, AccQual,
777  TypeQual);
778 }
779 
780 void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
781  StringRef ValueKind, unsigned &Offset,
783  unsigned PointeeAlign, StringRef Name,
784  StringRef TypeName,
785  StringRef BaseTypeName,
786  StringRef AccQual, StringRef TypeQual) {
787  auto Arg = Args.getDocument()->getMapNode();
788 
789  if (!Name.empty())
790  Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
791  if (!TypeName.empty())
792  Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
793  auto Size = DL.getTypeAllocSize(Ty);
794  auto Align = DL.getABITypeAlignment(Ty);
795  Arg[".size"] = Arg.getDocument()->getNode(Size);
796  Offset = alignTo(Offset, Align);
797  Arg[".offset"] = Arg.getDocument()->getNode(Offset);
798  Offset += Size;
799  Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
800  Arg[".value_type"] =
801  Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
802  if (PointeeAlign)
803  Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign);
804 
805  if (auto PtrTy = dyn_cast<PointerType>(Ty))
806  if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
807  Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
808 
809  if (auto AQ = getAccessQualifier(AccQual))
810  Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
811 
812  // TODO: Emit Arg[".actual_access"].
813 
814  SmallVector<StringRef, 1> SplitTypeQuals;
815  TypeQual.split(SplitTypeQuals, " ", -1, false);
816  for (StringRef Key : SplitTypeQuals) {
817  if (Key == "const")
818  Arg[".is_const"] = Arg.getDocument()->getNode(true);
819  else if (Key == "restrict")
820  Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
821  else if (Key == "volatile")
822  Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
823  else if (Key == "pipe")
824  Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
825  }
826 
827  Args.push_back(Arg);
828 }
829 
830 void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
831  unsigned &Offset,
832  msgpack::ArrayDocNode Args) {
833  int HiddenArgNumBytes =
834  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
835 
836  if (!HiddenArgNumBytes)
837  return;
838 
839  auto &DL = Func.getParent()->getDataLayout();
840  auto Int64Ty = Type::getInt64Ty(Func.getContext());
841 
842  if (HiddenArgNumBytes >= 8)
843  emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, Args);
844  if (HiddenArgNumBytes >= 16)
845  emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, Args);
846  if (HiddenArgNumBytes >= 24)
847  emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, Args);
848 
849  auto Int8PtrTy =
851 
852  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
853  // "none" argument.
854  if (HiddenArgNumBytes >= 32) {
855  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
856  emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, Args);
857  else
858  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
859  }
860 
861  // Emit "default queue" and "completion action" arguments if enqueue kernel is
862  // used, otherwise emit dummy "none" arguments.
863  if (HiddenArgNumBytes >= 48) {
864  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
865  emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
866  emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
867  } else {
868  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
869  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
870  }
871  }
872 
873  // Emit the pointer argument for multi-grid object.
874  if (HiddenArgNumBytes >= 56)
875  emitKernelArg(DL, Int8PtrTy, "hidden_multigrid_sync_arg", Offset, Args);
876 }
877 
879 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
880  const SIProgramInfo &ProgramInfo) const {
881  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
883  const Function &F = MF.getFunction();
884 
885  auto Kern = HSAMetadataDoc->getMapNode();
886 
887  Align MaxKernArgAlign;
888  Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
889  STM.getKernArgSegmentSize(F, MaxKernArgAlign));
890  Kern[".group_segment_fixed_size"] =
891  Kern.getDocument()->getNode(ProgramInfo.LDSSize);
892  Kern[".private_segment_fixed_size"] =
893  Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
894  Kern[".kernarg_segment_align"] =
895  Kern.getDocument()->getNode(std::max(Align(4), MaxKernArgAlign).value());
896  Kern[".wavefront_size"] =
897  Kern.getDocument()->getNode(STM.getWavefrontSize());
898  Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
899  Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
900  Kern[".max_flat_workgroup_size"] =
901  Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
902  Kern[".sgpr_spill_count"] =
903  Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
904  Kern[".vgpr_spill_count"] =
905  Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
906 
907  return Kern;
908 }
909 
911  return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
912 }
913 
915  emitVersion();
916  emitPrintf(Mod);
917  getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
918 }
919 
921  std::string HSAMetadataString;
922  raw_string_ostream StrOS(HSAMetadataString);
923  HSAMetadataDoc->toYAML(StrOS);
924 
925  if (DumpHSAMetadata)
926  dump(StrOS.str());
927  if (VerifyHSAMetadata)
928  verify(StrOS.str());
929 }
930 
932  const SIProgramInfo &ProgramInfo) {
933  auto &Func = MF.getFunction();
934  auto Kern = getHSAKernelProps(MF, ProgramInfo);
935 
938 
939  auto Kernels =
940  getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
941 
942  {
943  Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
944  Kern[".symbol"] = Kern.getDocument()->getNode(
945  (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
946  emitKernelLanguage(Func, Kern);
947  emitKernelAttrs(Func, Kern);
948  emitKernelArgs(Func, Kern);
949  }
950 
951  Kernels.push_back(Kern);
952 }
953 
954 } // end namespace HSAMD
955 } // end namespace AMDGPU
956 } // end namespace llvm
Type * getVectorElementType() const
Definition: Type.h:376
const NoneType None
Definition: None.h:23
void emitKernel(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) override
A parsed version of the target data layout string in and methods for querying it. ...
Definition: DataLayout.h:111
constexpr char Align[]
Key for Kernel::Arg::Metadata::mAlign.
Type
MessagePack types as defined in the standard, with the exception of Integer being divided into a sign...
Definition: MsgPackReader.h:48
raw_ostream & errs()
This returns a reference to a raw_ostream for standard error.
This class represents an incoming formal argument to a Function.
Definition: Argument.h:29
LLVM_NODISCARD std::string str() const
str - Get the contents as an std::string.
Definition: StringRef.h:232
AMDGPU specific subclass of TargetSubtarget.
bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override
This class represents lattice values for constants.
Definition: AllocatorList.h:23
A Module instance is used to store all the information related to an LLVM module. ...
Definition: Module.h:66
2: 32-bit floating point type
Definition: Type.h:59
amdgpu Simplify well known AMD library false FunctionCallee Value const Twine & Name
Document * getDocument() const
constexpr char PointeeAlign[]
Key for Kernel::Arg::Metadata::mPointeeAlign.
void toYAML(raw_ostream &OS)
Convert MsgPack Document to YAML text.
LLVM_NODISCARD bool startswith(StringRef Prefix) const
Check if this string starts with the given Prefix.
Definition: StringRef.h:270
bool hasFnAttribute(Attribute::AttrKind Kind) const
Return true if the function has the attribute.
Definition: Function.h:323
Metadata node.
Definition: Metadata.h:863
F(f)
const MDOperand & getOperand(unsigned I) const
Definition: Metadata.h:1068
unsigned getPointerAddressSpace() const
Get the address space of this pointer or pointer vector type.
Definition: DerivedTypes.h:635
unsigned getKernArgSegmentSize(const Function &F, Align &MaxAlign) const
1: 16-bit floating point type
Definition: Type.h:58
static IntegerType * getInt64Ty(LLVMContext &C)
Definition: Type.cpp:181
15: Pointers
Definition: Type.h:75
Address space for region memory. (GDS)
Definition: AMDGPU.h:271
StringSwitch & Case(StringLiteral S, T Value)
Definition: StringSwitch.h:67
Track resource usage for kernels / entry functions.
Definition: SIProgramInfo.h:21
A DocNode that is an array.
Address space for global memory (RAT0, VTX0).
Definition: AMDGPU.h:270
Type * getPointerElementType() const
Definition: Type.h:381
const DataLayout & getDataLayout() const
Get the data layout for the module&#39;s target platform.
Definition: Module.cpp:369
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
Definition: Twine.h:80
TypeID getTypeID() const
Return the type id for the type.
Definition: Type.h:138
constexpr char Printf[]
Key for HSA::Metadata::mPrintf.
Defines struct to track resource usage for kernels and entry functions.
MDNode * getMetadata(unsigned KindID) const
Get the current metadata attachments for the given kind, if any.
Definition: Metadata.cpp:1440
LLVM_NODISCARD R Default(T Value)
Definition: StringSwitch.h:181
bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override
LLVM_NODISCARD bool empty() const
empty - Check if the string is empty.
Definition: StringRef.h:140
bool fromYAML(StringRef S)
Read YAML text into the MsgPack document. Returns false on failure.
std::error_code fromString(std::string String, Metadata &HSAMetadata)
Converts String to HSAMetadata.
Key
PAL metadata keys.
Type * getType() const
All values are typed, get the type of this value.
Definition: Value.h:246
static cl::opt< bool > VerifyHSAMetadata("amdgpu-verify-hsa-metadata", cl::desc("Verify AMDGPU HSA Metadata"))
DocNode getNode()
Create a nil node associated with this Document.
op_range operands() const
Definition: Metadata.h:1066
bool onlyReadsMemory() const
Return true if this argument has the readonly or readnone attribute.
Definition: Function.cpp:169
NamedMDNode * getNamedMetadata(const Twine &Name) const
Return the first NamedMDNode in the module with the specified name.
Definition: Module.cpp:250
uint64_t value() const
This is a hole in the type system and should not be abused.
Definition: Alignment.h:86
constexpr char Attrs[]
Key for Kernel::Metadata::mAttrs.
AMDGPU HSA Metadata Streamer.
ArrayDocNode getArrayNode()
Create an empty Array node associated with this Document.
bool hasNoAliasAttr() const
Return true if this argument has the noalias attribute.
Definition: Function.cpp:138
11: Arbitrary bit width integers
Definition: Type.h:71
Address space for private memory.
Definition: AMDGPU.h:275
#define P(N)
A switch()-like statement whose cases are string literals.
Definition: StringSwitch.h:42
bool hasName() const
Definition: Value.h:252
const TargetSubtargetInfo & getSubtarget() const
getSubtarget - Return the subtarget for which this machine code is being compiled.
The instances of the Type class are immutable: once they are created, they are never changed...
Definition: Type.h:46
static cl::opt< bool > DumpHSAMetadata("amdgpu-dump-hsa-metadata", cl::desc("Dump AMDGPU HSA Metadata"))
Address space for local memory.
Definition: AMDGPU.h:274
In-memory representation of kernel metadata.
Address space for flat memory.
Definition: AMDGPU.h:269
ValueKind
Value kinds.
This file contains the declarations for the subclasses of Constant, which represent the different fla...
bool isPointerTy() const
True if this is an instance of PointerType.
Definition: Type.h:224
std::vector< uint32_t > mVersion
HSA metadata version. Required.
TypeSize getTypeAllocSize(Type *Ty) const
Returns the offset in bytes between successive objects of the specified type, including alignment pad...
Definition: DataLayout.h:486
LLVM_NODISCARD size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition: StringRef.h:299
Ty * getInfo()
getInfo - Keep track of various per-function pieces of information for backends that would like to do...
constexpr char TypeName[]
Key for Kernel::Arg::Metadata::mTypeName.
amdgpu Simplify well known AMD library false FunctionCallee Value * Arg
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function. ...
Definition: Function.cpp:205
static wasm::ValType getType(const TargetRegisterClass *RC)
static PointerType * getInt8PtrTy(LLVMContext &C, unsigned AS=0)
Definition: Type.cpp:224
MapDocNode getMapNode()
Create an empty Map node associated with this Document.
std::string & str()
Flushes the stream contents to the target string and returns the string&#39;s reference.
Definition: raw_ostream.h:519
This struct is a compact representation of a valid (non-zero power of two) alignment.
Definition: Alignment.h:40
AccessQualifier
Access qualifiers.
std::vector< std::string > mPrintf
Printf metadata. Optional.
unsigned getWavefrontSize() const
Align max(MaybeAlign Lhs, Align Rhs)
Definition: Alignment.h:390
16: SIMD &#39;packed&#39; format, or other vector type
Definition: Type.h:76
CallingConv::ID getCallingConv() const
getCallingConv()/setCallingConv(CC) - These method get and set the calling convention of this functio...
Definition: Function.h:212
unsigned getParamAlignment() const
If this is a byval or inalloca argument, return its alignment.
Definition: Function.cpp:111
Module.h This file contains the declarations for the Module class.
LLVM_NODISCARD std::pair< StringRef, StringRef > split(char Separator) const
Split into two substrings around the first occurrence of a separator character.
Definition: StringRef.h:710
AddressSpace
Definition: NVPTXBaseInfo.h:21
unsigned getABITypeAlignment(Type *Ty) const
Returns the minimum ABI-required alignment for the specified type.
Definition: DataLayout.cpp:755
constexpr char CodeProps[]
Key for Kernel::Metadata::mCodeProps.
ValueType
Value types.
const Function & getFunction() const
Return the LLVM function that this machine code represents.
constexpr char LanguageVersion[]
Key for Kernel::Metadata::mLanguageVersion.
SPIR_KERNEL - Calling convention for SPIR kernel functions.
Definition: CallingConv.h:141
constexpr uint32_t VersionMinor
HSA metadata minor version.
unsigned getArgNo() const
Return the index of this formal argument in its containing function.
Definition: Argument.h:47
Address space for constant memory (VTX2).
Definition: AMDGPU.h:273
Simple in-memory representation of a document of msgpack objects with ability to find and create arra...
A DocNode that is a map.
constexpr char Kernels[]
Key for HSA::Metadata::mKernels.
void emitKernel(const MachineFunction &MF, const SIProgramInfo &ProgramInfo) override
std::error_code toString(Metadata HSAMetadata, std::string &String)
Converts HSAMetadata to String.
This class keeps track of the SPI_SP_INPUT_ADDR config register, which tells the hardware which inter...
uint64_t alignTo(uint64_t Size, Align A)
Returns a multiple of A needed to store Size bytes.
Definition: Alignment.h:163
const Function * getParent() const
Definition: Argument.h:41
static const size_t npos
Definition: StringRef.h:50
unsigned getIntegerBitWidth() const
Definition: DerivedTypes.h:102
StringRef getValueAsString() const
Return the attribute&#39;s value as a string.
Definition: Attributes.cpp:220
constexpr char DebugProps[]
Key for Kernel::Metadata::mDebugProps.
AddressSpaceQualifier
Address space qualifiers.
StringRef getName() const
Return a constant reference to the value&#39;s name.
Definition: Value.cpp:214
int getIntegerAttribute(const Function &F, StringRef Name, int Default)
constexpr char AccQual[]
Key for Kernel::Arg::Metadata::mAccQual.
uint32_t Size
Definition: Profile.cpp:46
3: 64-bit floating point type
Definition: Type.h:60
virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict)=0
Emit HSA Metadata.
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
A raw_ostream that writes to an std::string.
Definition: raw_ostream.h:503
Module * getParent()
Get the module that this global value is contained inside of...
Definition: GlobalValue.h:575
Attribute getFnAttribute(Attribute::AttrKind Kind) const
Return the attribute for the given attribute kind.
Definition: Function.h:333
StringRef - Represent a constant reference to a string, i.e.
Definition: StringRef.h:48
std::vector< Kernel::Metadata > mKernels
Kernels metadata. Required.
unsigned getNumOperands() const
Return number of MDNode operands.
Definition: Metadata.h:1074
Calling convention for AMDGPU code object kernels.
Definition: CallingConv.h:205
const uint64_t Version
Definition: InstrProf.h:980
constexpr char Args[]
Key for Kernel::Metadata::mArgs.
constexpr uint32_t VersionMajor
HSA metadata major version.
iterator_range< arg_iterator > args()
Definition: Function.h:724