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  unsigned MaxKernArgAlign;
222  HSACodeProps.mKernargSegmentSize = STM.getKernArgSegmentSize(F,
223  MaxKernArgAlign);
224  HSACodeProps.mGroupSegmentFixedSize = ProgramInfo.LDSSize;
225  HSACodeProps.mPrivateSegmentFixedSize = ProgramInfo.ScratchSize;
226  HSACodeProps.mKernargSegmentAlign = std::max(MaxKernArgAlign, 4u);
227  HSACodeProps.mWavefrontSize = STM.getWavefrontSize();
228  HSACodeProps.mNumSGPRs = ProgramInfo.NumSGPR;
229  HSACodeProps.mNumVGPRs = ProgramInfo.NumVGPR;
230  HSACodeProps.mMaxFlatWorkGroupSize = MFI.getMaxFlatWorkGroupSize();
231  HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack;
232  HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled();
233  HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
234  HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
235 
236  return HSACodeProps;
237 }
238 
239 Kernel::DebugProps::Metadata
240 MetadataStreamerV2::getHSADebugProps(const MachineFunction &MF,
241  const SIProgramInfo &ProgramInfo) const {
242  return HSAMD::Kernel::DebugProps::Metadata();
243 }
244 
245 void MetadataStreamerV2::emitVersion() {
246  auto &Version = HSAMetadata.mVersion;
247 
248  Version.push_back(VersionMajor);
249  Version.push_back(VersionMinor);
250 }
251 
252 void MetadataStreamerV2::emitPrintf(const Module &Mod) {
253  auto &Printf = HSAMetadata.mPrintf;
254 
255  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
256  if (!Node)
257  return;
258 
259  for (auto Op : Node->operands())
260  if (Op->getNumOperands())
261  Printf.push_back(cast<MDString>(Op->getOperand(0))->getString());
262 }
263 
264 void MetadataStreamerV2::emitKernelLanguage(const Function &Func) {
265  auto &Kernel = HSAMetadata.mKernels.back();
266 
267  // TODO: What about other languages?
268  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
269  if (!Node || !Node->getNumOperands())
270  return;
271  auto Op0 = Node->getOperand(0);
272  if (Op0->getNumOperands() <= 1)
273  return;
274 
275  Kernel.mLanguage = "OpenCL C";
276  Kernel.mLanguageVersion.push_back(
277  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
278  Kernel.mLanguageVersion.push_back(
279  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
280 }
281 
282 void MetadataStreamerV2::emitKernelAttrs(const Function &Func) {
283  auto &Attrs = HSAMetadata.mKernels.back().mAttrs;
284 
285  if (auto Node = Func.getMetadata("reqd_work_group_size"))
286  Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
287  if (auto Node = Func.getMetadata("work_group_size_hint"))
288  Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
289  if (auto Node = Func.getMetadata("vec_type_hint")) {
290  Attrs.mVecTypeHint = getTypeName(
291  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
292  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
293  }
294  if (Func.hasFnAttribute("runtime-handle")) {
295  Attrs.mRuntimeHandle =
296  Func.getFnAttribute("runtime-handle").getValueAsString().str();
297  }
298 }
299 
300 void MetadataStreamerV2::emitKernelArgs(const Function &Func) {
301  for (auto &Arg : Func.args())
302  emitKernelArg(Arg);
303 
304  emitHiddenKernelArgs(Func);
305 }
306 
307 void MetadataStreamerV2::emitKernelArg(const Argument &Arg) {
308  auto Func = Arg.getParent();
309  auto ArgNo = Arg.getArgNo();
310  const MDNode *Node;
311 
312  StringRef Name;
313  Node = Func->getMetadata("kernel_arg_name");
314  if (Node && ArgNo < Node->getNumOperands())
315  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
316  else if (Arg.hasName())
317  Name = Arg.getName();
318 
319  StringRef TypeName;
320  Node = Func->getMetadata("kernel_arg_type");
321  if (Node && ArgNo < Node->getNumOperands())
322  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
323 
324  StringRef BaseTypeName;
325  Node = Func->getMetadata("kernel_arg_base_type");
326  if (Node && ArgNo < Node->getNumOperands())
327  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
328 
329  StringRef AccQual;
330  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
331  Arg.hasNoAliasAttr()) {
332  AccQual = "read_only";
333  } else {
334  Node = Func->getMetadata("kernel_arg_access_qual");
335  if (Node && ArgNo < Node->getNumOperands())
336  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
337  }
338 
339  StringRef TypeQual;
340  Node = Func->getMetadata("kernel_arg_type_qual");
341  if (Node && ArgNo < Node->getNumOperands())
342  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
343 
344  Type *Ty = Arg.getType();
345  const DataLayout &DL = Func->getParent()->getDataLayout();
346 
347  unsigned PointeeAlign = 0;
348  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
349  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
350  PointeeAlign = Arg.getParamAlignment();
351  if (PointeeAlign == 0)
352  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
353  }
354  }
355 
356  emitKernelArg(DL, Ty, getValueKind(Arg.getType(), TypeQual, BaseTypeName),
357  PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual);
358 }
359 
360 void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
362  unsigned PointeeAlign, StringRef Name,
363  StringRef TypeName,
364  StringRef BaseTypeName,
365  StringRef AccQual, StringRef TypeQual) {
366  HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
367  auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
368 
369  Arg.mName = Name;
370  Arg.mTypeName = TypeName;
371  Arg.mSize = DL.getTypeAllocSize(Ty);
372  Arg.mAlign = DL.getABITypeAlignment(Ty);
373  Arg.mValueKind = ValueKind;
374  Arg.mValueType = getValueType(Ty, BaseTypeName);
375  Arg.mPointeeAlign = PointeeAlign;
376 
377  if (auto PtrTy = dyn_cast<PointerType>(Ty))
378  Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
379 
380  Arg.mAccQual = getAccessQualifier(AccQual);
381 
382  // TODO: Emit Arg.mActualAccQual.
383 
384  SmallVector<StringRef, 1> SplitTypeQuals;
385  TypeQual.split(SplitTypeQuals, " ", -1, false);
386  for (StringRef Key : SplitTypeQuals) {
387  auto P = StringSwitch<bool*>(Key)
388  .Case("const", &Arg.mIsConst)
389  .Case("restrict", &Arg.mIsRestrict)
390  .Case("volatile", &Arg.mIsVolatile)
391  .Case("pipe", &Arg.mIsPipe)
392  .Default(nullptr);
393  if (P)
394  *P = true;
395  }
396 }
397 
398 void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func) {
399  int HiddenArgNumBytes =
400  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
401 
402  if (!HiddenArgNumBytes)
403  return;
404 
405  auto &DL = Func.getParent()->getDataLayout();
406  auto Int64Ty = Type::getInt64Ty(Func.getContext());
407 
408  if (HiddenArgNumBytes >= 8)
409  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetX);
410  if (HiddenArgNumBytes >= 16)
411  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetY);
412  if (HiddenArgNumBytes >= 24)
413  emitKernelArg(DL, Int64Ty, ValueKind::HiddenGlobalOffsetZ);
414 
415  auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
417 
418  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
419  // "none" argument.
420  if (HiddenArgNumBytes >= 32) {
421  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
422  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenPrintfBuffer);
423  else
424  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
425  }
426 
427  // Emit "default queue" and "completion action" arguments if enqueue kernel is
428  // used, otherwise emit dummy "none" arguments.
429  if (HiddenArgNumBytes >= 48) {
430  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
431  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenDefaultQueue);
432  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenCompletionAction);
433  } else {
434  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
435  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
436  }
437  }
438 
439  // Emit the pointer argument for multi-grid object.
440  if (HiddenArgNumBytes >= 56)
441  emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenMultiGridSyncArg);
442 }
443 
445  return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
446 }
447 
449  emitVersion();
450  emitPrintf(Mod);
451 }
452 
454  std::string HSAMetadataString;
455  if (toString(HSAMetadata, HSAMetadataString))
456  return;
457 
458  if (DumpHSAMetadata)
459  dump(HSAMetadataString);
460  if (VerifyHSAMetadata)
461  verify(HSAMetadataString);
462 }
463 
465  const SIProgramInfo &ProgramInfo) {
466  auto &Func = MF.getFunction();
467  if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
468  return;
469 
470  auto CodeProps = getHSACodeProps(MF, ProgramInfo);
471  auto DebugProps = getHSADebugProps(MF, ProgramInfo);
472 
473  HSAMetadata.mKernels.push_back(Kernel::Metadata());
474  auto &Kernel = HSAMetadata.mKernels.back();
475 
476  Kernel.mName = Func.getName();
477  Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str();
478  emitKernelLanguage(Func);
479  emitKernelAttrs(Func);
480  emitKernelArgs(Func);
481  HSAMetadata.mKernels.back().mCodeProps = CodeProps;
482  HSAMetadata.mKernels.back().mDebugProps = DebugProps;
483 }
484 
485 //===----------------------------------------------------------------------===//
486 // HSAMetadataStreamerV3
487 //===----------------------------------------------------------------------===//
488 
489 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
490  errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
491 }
492 
493 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
494  errs() << "AMDGPU HSA Metadata Parser Test: ";
495 
496  msgpack::Document FromHSAMetadataString;
497 
498  if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
499  errs() << "FAIL\n";
500  return;
501  }
502 
503  std::string ToHSAMetadataString;
504  raw_string_ostream StrOS(ToHSAMetadataString);
505  FromHSAMetadataString.toYAML(StrOS);
506 
507  errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
508  if (HSAMetadataString != ToHSAMetadataString) {
509  errs() << "Original input: " << HSAMetadataString << '\n'
510  << "Produced output: " << StrOS.str() << '\n';
511  }
512 }
513 
515 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
517  .Case("read_only", StringRef("read_only"))
518  .Case("write_only", StringRef("write_only"))
519  .Case("read_write", StringRef("read_write"))
520  .Default(None);
521 }
522 
524 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
525  switch (AddressSpace) {
527  return StringRef("private");
529  return StringRef("global");
531  return StringRef("constant");
533  return StringRef("local");
535  return StringRef("generic");
537  return StringRef("region");
538  default:
539  return None;
540  }
541 }
542 
543 StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
544  StringRef BaseTypeName) const {
545  if (TypeQual.find("pipe") != StringRef::npos)
546  return "pipe";
547 
548  return StringSwitch<StringRef>(BaseTypeName)
549  .Case("image1d_t", "image")
550  .Case("image1d_array_t", "image")
551  .Case("image1d_buffer_t", "image")
552  .Case("image2d_t", "image")
553  .Case("image2d_array_t", "image")
554  .Case("image2d_array_depth_t", "image")
555  .Case("image2d_array_msaa_t", "image")
556  .Case("image2d_array_msaa_depth_t", "image")
557  .Case("image2d_depth_t", "image")
558  .Case("image2d_msaa_t", "image")
559  .Case("image2d_msaa_depth_t", "image")
560  .Case("image3d_t", "image")
561  .Case("sampler_t", "sampler")
562  .Case("queue_t", "queue")
563  .Default(isa<PointerType>(Ty)
565  ? "dynamic_shared_pointer"
566  : "global_buffer")
567  : "by_value");
568 }
569 
570 StringRef MetadataStreamerV3::getValueType(Type *Ty, StringRef TypeName) const {
571  switch (Ty->getTypeID()) {
572  case Type::IntegerTyID: {
573  auto Signed = !TypeName.startswith("u");
574  switch (Ty->getIntegerBitWidth()) {
575  case 8:
576  return Signed ? "i8" : "u8";
577  case 16:
578  return Signed ? "i16" : "u16";
579  case 32:
580  return Signed ? "i32" : "u32";
581  case 64:
582  return Signed ? "i64" : "u64";
583  default:
584  return "struct";
585  }
586  }
587  case Type::HalfTyID:
588  return "f16";
589  case Type::FloatTyID:
590  return "f32";
591  case Type::DoubleTyID:
592  return "f64";
593  case Type::PointerTyID:
594  return getValueType(Ty->getPointerElementType(), TypeName);
595  case Type::VectorTyID:
596  return getValueType(Ty->getVectorElementType(), TypeName);
597  default:
598  return "struct";
599  }
600 }
601 
602 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
603  switch (Ty->getTypeID()) {
604  case Type::IntegerTyID: {
605  if (!Signed)
606  return (Twine('u') + getTypeName(Ty, true)).str();
607 
608  auto BitWidth = Ty->getIntegerBitWidth();
609  switch (BitWidth) {
610  case 8:
611  return "char";
612  case 16:
613  return "short";
614  case 32:
615  return "int";
616  case 64:
617  return "long";
618  default:
619  return (Twine('i') + Twine(BitWidth)).str();
620  }
621  }
622  case Type::HalfTyID:
623  return "half";
624  case Type::FloatTyID:
625  return "float";
626  case Type::DoubleTyID:
627  return "double";
628  case Type::VectorTyID: {
629  auto VecTy = cast<VectorType>(Ty);
630  auto ElTy = VecTy->getElementType();
631  auto NumElements = VecTy->getVectorNumElements();
632  return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
633  }
634  default:
635  return "unknown";
636  }
637 }
638 
640 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
641  auto Dims = HSAMetadataDoc->getArrayNode();
642  if (Node->getNumOperands() != 3)
643  return Dims;
644 
645  for (auto &Op : Node->operands())
646  Dims.push_back(Dims.getDocument()->getNode(
647  uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
648  return Dims;
649 }
650 
651 void MetadataStreamerV3::emitVersion() {
652  auto Version = HSAMetadataDoc->getArrayNode();
653  Version.push_back(Version.getDocument()->getNode(VersionMajor));
654  Version.push_back(Version.getDocument()->getNode(VersionMinor));
655  getRootMetadata("amdhsa.version") = Version;
656 }
657 
658 void MetadataStreamerV3::emitPrintf(const Module &Mod) {
659  auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
660  if (!Node)
661  return;
662 
663  auto Printf = HSAMetadataDoc->getArrayNode();
664  for (auto Op : Node->operands())
665  if (Op->getNumOperands())
666  Printf.push_back(Printf.getDocument()->getNode(
667  cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
668  getRootMetadata("amdhsa.printf") = Printf;
669 }
670 
671 void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
672  msgpack::MapDocNode Kern) {
673  // TODO: What about other languages?
674  auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
675  if (!Node || !Node->getNumOperands())
676  return;
677  auto Op0 = Node->getOperand(0);
678  if (Op0->getNumOperands() <= 1)
679  return;
680 
681  Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
682  auto LanguageVersion = Kern.getDocument()->getArrayNode();
683  LanguageVersion.push_back(Kern.getDocument()->getNode(
684  mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
685  LanguageVersion.push_back(Kern.getDocument()->getNode(
686  mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
687  Kern[".language_version"] = LanguageVersion;
688 }
689 
690 void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
691  msgpack::MapDocNode Kern) {
692 
693  if (auto Node = Func.getMetadata("reqd_work_group_size"))
694  Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
695  if (auto Node = Func.getMetadata("work_group_size_hint"))
696  Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
697  if (auto Node = Func.getMetadata("vec_type_hint")) {
698  Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
699  getTypeName(
700  cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
701  mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
702  /*Copy=*/true);
703  }
704  if (Func.hasFnAttribute("runtime-handle")) {
705  Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
706  Func.getFnAttribute("runtime-handle").getValueAsString().str(),
707  /*Copy=*/true);
708  }
709 }
710 
711 void MetadataStreamerV3::emitKernelArgs(const Function &Func,
712  msgpack::MapDocNode Kern) {
713  unsigned Offset = 0;
714  auto Args = HSAMetadataDoc->getArrayNode();
715  for (auto &Arg : Func.args())
716  emitKernelArg(Arg, Offset, Args);
717 
718  emitHiddenKernelArgs(Func, Offset, Args);
719 
720  Kern[".args"] = Args;
721 }
722 
723 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
725  auto Func = Arg.getParent();
726  auto ArgNo = Arg.getArgNo();
727  const MDNode *Node;
728 
729  StringRef Name;
730  Node = Func->getMetadata("kernel_arg_name");
731  if (Node && ArgNo < Node->getNumOperands())
732  Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
733  else if (Arg.hasName())
734  Name = Arg.getName();
735 
737  Node = Func->getMetadata("kernel_arg_type");
738  if (Node && ArgNo < Node->getNumOperands())
739  TypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
740 
741  StringRef BaseTypeName;
742  Node = Func->getMetadata("kernel_arg_base_type");
743  if (Node && ArgNo < Node->getNumOperands())
744  BaseTypeName = cast<MDString>(Node->getOperand(ArgNo))->getString();
745 
747  if (Arg.getType()->isPointerTy() && Arg.onlyReadsMemory() &&
748  Arg.hasNoAliasAttr()) {
749  AccQual = "read_only";
750  } else {
751  Node = Func->getMetadata("kernel_arg_access_qual");
752  if (Node && ArgNo < Node->getNumOperands())
753  AccQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
754  }
755 
756  StringRef TypeQual;
757  Node = Func->getMetadata("kernel_arg_type_qual");
758  if (Node && ArgNo < Node->getNumOperands())
759  TypeQual = cast<MDString>(Node->getOperand(ArgNo))->getString();
760 
761  Type *Ty = Arg.getType();
762  const DataLayout &DL = Func->getParent()->getDataLayout();
763 
764  unsigned PointeeAlign = 0;
765  if (auto PtrTy = dyn_cast<PointerType>(Ty)) {
766  if (PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
767  PointeeAlign = Arg.getParamAlignment();
768  if (PointeeAlign == 0)
769  PointeeAlign = DL.getABITypeAlignment(PtrTy->getElementType());
770  }
771  }
772 
773  emitKernelArg(Func->getParent()->getDataLayout(), Arg.getType(),
774  getValueKind(Arg.getType(), TypeQual, BaseTypeName), Offset,
775  Args, PointeeAlign, Name, TypeName, BaseTypeName, AccQual,
776  TypeQual);
777 }
778 
779 void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
780  StringRef ValueKind, unsigned &Offset,
782  unsigned PointeeAlign, StringRef Name,
783  StringRef TypeName,
784  StringRef BaseTypeName,
785  StringRef AccQual, StringRef TypeQual) {
786  auto Arg = Args.getDocument()->getMapNode();
787 
788  if (!Name.empty())
789  Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
790  if (!TypeName.empty())
791  Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
792  auto Size = DL.getTypeAllocSize(Ty);
793  auto Align = DL.getABITypeAlignment(Ty);
794  Arg[".size"] = Arg.getDocument()->getNode(Size);
795  Offset = alignTo(Offset, Align);
796  Arg[".offset"] = Arg.getDocument()->getNode(Offset);
797  Offset += Size;
798  Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
799  Arg[".value_type"] =
800  Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
801  if (PointeeAlign)
802  Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign);
803 
804  if (auto PtrTy = dyn_cast<PointerType>(Ty))
805  if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
806  Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
807 
808  if (auto AQ = getAccessQualifier(AccQual))
809  Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
810 
811  // TODO: Emit Arg[".actual_access"].
812 
813  SmallVector<StringRef, 1> SplitTypeQuals;
814  TypeQual.split(SplitTypeQuals, " ", -1, false);
815  for (StringRef Key : SplitTypeQuals) {
816  if (Key == "const")
817  Arg[".is_const"] = Arg.getDocument()->getNode(true);
818  else if (Key == "restrict")
819  Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
820  else if (Key == "volatile")
821  Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
822  else if (Key == "pipe")
823  Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
824  }
825 
826  Args.push_back(Arg);
827 }
828 
829 void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
830  unsigned &Offset,
831  msgpack::ArrayDocNode Args) {
832  int HiddenArgNumBytes =
833  getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
834 
835  if (!HiddenArgNumBytes)
836  return;
837 
838  auto &DL = Func.getParent()->getDataLayout();
839  auto Int64Ty = Type::getInt64Ty(Func.getContext());
840 
841  if (HiddenArgNumBytes >= 8)
842  emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, Args);
843  if (HiddenArgNumBytes >= 16)
844  emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, Args);
845  if (HiddenArgNumBytes >= 24)
846  emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, Args);
847 
848  auto Int8PtrTy =
850 
851  // Emit "printf buffer" argument if printf is used, otherwise emit dummy
852  // "none" argument.
853  if (HiddenArgNumBytes >= 32) {
854  if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
855  emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, Args);
856  else
857  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
858  }
859 
860  // Emit "default queue" and "completion action" arguments if enqueue kernel is
861  // used, otherwise emit dummy "none" arguments.
862  if (HiddenArgNumBytes >= 48) {
863  if (Func.hasFnAttribute("calls-enqueue-kernel")) {
864  emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
865  emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
866  } else {
867  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
868  emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
869  }
870  }
871 
872  // Emit the pointer argument for multi-grid object.
873  if (HiddenArgNumBytes >= 56)
874  emitKernelArg(DL, Int8PtrTy, "hidden_multigrid_sync_arg", Offset, Args);
875 }
876 
878 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
879  const SIProgramInfo &ProgramInfo) const {
880  const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
882  const Function &F = MF.getFunction();
883 
884  auto Kern = HSAMetadataDoc->getMapNode();
885 
886  unsigned MaxKernArgAlign;
887  Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
888  STM.getKernArgSegmentSize(F, MaxKernArgAlign));
889  Kern[".group_segment_fixed_size"] =
890  Kern.getDocument()->getNode(ProgramInfo.LDSSize);
891  Kern[".private_segment_fixed_size"] =
892  Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
893  Kern[".kernarg_segment_align"] =
894  Kern.getDocument()->getNode(std::max(uint32_t(4), MaxKernArgAlign));
895  Kern[".wavefront_size"] =
896  Kern.getDocument()->getNode(STM.getWavefrontSize());
897  Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
898  Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
899  Kern[".max_flat_workgroup_size"] =
900  Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
901  Kern[".sgpr_spill_count"] =
902  Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
903  Kern[".vgpr_spill_count"] =
904  Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
905 
906  return Kern;
907 }
908 
910  return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
911 }
912 
914  emitVersion();
915  emitPrintf(Mod);
916  getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
917 }
918 
920  std::string HSAMetadataString;
921  raw_string_ostream StrOS(HSAMetadataString);
922  HSAMetadataDoc->toYAML(StrOS);
923 
924  if (DumpHSAMetadata)
925  dump(StrOS.str());
926  if (VerifyHSAMetadata)
927  verify(StrOS.str());
928 }
929 
931  const SIProgramInfo &ProgramInfo) {
932  auto &Func = MF.getFunction();
933  auto Kern = getHSAKernelProps(MF, ProgramInfo);
934 
937 
938  auto Kernels =
939  getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
940 
941  {
942  Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
943  Kern[".symbol"] = Kern.getDocument()->getNode(
944  (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
945  emitKernelLanguage(Func, Kern);
946  emitKernelAttrs(Func, Kern);
947  emitKernelArgs(Func, Kern);
948  }
949 
950  Kernels.push_back(Kern);
951 }
952 
953 } // end namespace HSAMD
954 } // end namespace AMDGPU
955 } // end namespace llvm
Type * getVectorElementType() const
Definition: Type.h:371
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
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.
GCNRegPressure max(const GCNRegPressure &P1, const GCNRegPressure &P2)
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:218
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:65
Address space for region memory. (GDS)
Definition: AMDGPU.h:271
2: 32-bit floating point type
Definition: Type.h:58
amdgpu Simplify well known AMD library false FunctionCallee Value const Twine & Name
Address space for local memory.
Definition: AMDGPU.h:274
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:256
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:580
1: 16-bit floating point type
Definition: Type.h:57
static IntegerType * getInt64Ty(LLVMContext &C)
Definition: Type.cpp:176
15: Pointers
Definition: Type.h:74
Address space for constant memory (VTX2).
Definition: AMDGPU.h:273
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.
Type * getPointerElementType() const
Definition: Type.h:376
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:137
constexpr char Printf[]
Key for HSA::Metadata::mPrintf.
Defines struct to track resource usage for kernels and entry functions.
Address space for global memory (RAT0, VTX0).
Definition: AMDGPU.h:270
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:126
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:245
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
Address space for flat memory.
Definition: AMDGPU.h:269
NamedMDNode * getNamedMetadata(const Twine &Name) const
Return the first NamedMDNode in the module with the specified name.
Definition: Module.cpp:250
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:70
#define P(N)
A switch()-like statement whose cases are string literals.
Definition: StringSwitch.h:42
bool hasName() const
Definition: Value.h:251
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:45
static cl::opt< bool > DumpHSAMetadata("amdgpu-dump-hsa-metadata", cl::desc("Dump AMDGPU HSA Metadata"))
In-memory representation of kernel metadata.
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:223
std::vector< uint32_t > mVersion
HSA metadata version. Required.
LLVM_NODISCARD size_t find(char C, size_t From=0) const
Search for the first character C in the string.
Definition: StringRef.h:285
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
Address space for private memory.
Definition: AMDGPU.h:275
static wasm::ValType getType(const TargetRegisterClass *RC)
static PointerType * getInt8PtrTy(LLVMContext &C, unsigned AS=0)
Definition: Type.cpp:219
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
unsigned getKernArgSegmentSize(const Function &F, unsigned &MaxAlign) const
AccessQualifier
Access qualifiers.
std::vector< std::string > mPrintf
Printf metadata. Optional.
unsigned getWavefrontSize() const
Calling convention for AMDGPU code object kernels.
Definition: CallingConv.h:200
16: SIMD &#39;packed&#39; format, or other vector type
Definition: Type.h:75
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:696
SPIR_KERNEL - Calling convention for SPIR kernel functions.
Definition: CallingConv.h:136
AddressSpace
Definition: NVPTXBaseInfo.h:21
unsigned getABITypeAlignment(Type *Ty) const
Returns the minimum ABI-required alignment for the specified type.
Definition: DataLayout.cpp:752
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.
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
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
uint64_t getTypeAllocSize(Type *Ty) const
Returns the offset in bytes between successive objects of the specified type, including alignment pad...
Definition: DataLayout.h:470
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:126
const Function * getParent() const
Definition: Argument.h:41
static const size_t npos
Definition: StringRef.h:50
unsigned getIntegerBitWidth() const
Definition: DerivedTypes.h:97
StringRef getValueAsString() const
Return the attribute&#39;s value as a string.
Definition: Attributes.cpp:223
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:59
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
const uint64_t Version
Definition: InstrProf.h:984
constexpr char Args[]
Key for Kernel::Metadata::mArgs.
constexpr uint32_t VersionMajor
HSA metadata major version.
iterator_range< arg_iterator > args()
Definition: Function.h:719