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