lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

   1 //===--- AMDGPUHSAMetadataStreamer.cpp --------------------------*- C++ -*-===//
   2 //
   3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
   4 // See https://llvm.org/LICENSE.txt for license information.
   5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 /// \file
  10 /// AMDGPU HSA Metadata Streamer.
  11 ///
  12 //
  13 //===----------------------------------------------------------------------===//
  14
  15 #include "AMDGPUHSAMetadataStreamer.h"
  16 #include "AMDGPU.h"
  17 #include "AMDGPUSubtarget.h"
  18 #include "MCTargetDesc/AMDGPUTargetStreamer.h"
  19 #include "SIMachineFunctionInfo.h"
  20 #include "SIProgramInfo.h"
  21 #include "Utils/AMDGPUBaseInfo.h"
  22 #include "llvm/ADT/StringSwitch.h"
  23 #include "llvm/IR/Constants.h"
  24 #include "llvm/IR/Module.h"
  25 #include "llvm/Support/raw_ostream.h"
  26
  27 namespace llvm {
  28
  29 static cl::opt<bool> DumpHSAMetadata(
  30     "amdgpu-dump-hsa-metadata",
  31     cl::desc("Dump AMDGPU HSA Metadata"));
  32 static cl::opt<bool> VerifyHSAMetadata(
  33     "amdgpu-verify-hsa-metadata",
  34     cl::desc("Verify AMDGPU HSA Metadata"));
  35
  36 namespace AMDGPU {
  37 namespace HSAMD {
  38
  39 //===----------------------------------------------------------------------===//
  40 // HSAMetadataStreamerV2
  41 //===----------------------------------------------------------------------===//
  42 void MetadataStreamerV2::dump(StringRef HSAMetadataString) const {
  43   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
  44 }
  45
  46 void MetadataStreamerV2::verify(StringRef HSAMetadataString) const {
  47   errs() << "AMDGPU HSA Metadata Parser Test: ";
  48
  49   HSAMD::Metadata FromHSAMetadataString;
  50   if (fromString(HSAMetadataString, FromHSAMetadataString)) {
  51     errs() << "FAIL\n";
  52     return;
  53   }
  54
  55   std::string ToHSAMetadataString;
  56   if (toString(FromHSAMetadataString, ToHSAMetadataString)) {
  57     errs() << "FAIL\n";
  58     return;
  59   }
  60
  61   errs() << (HSAMetadataString == ToHSAMetadataString ? "PASS" : "FAIL")
  62          << '\n';
  63   if (HSAMetadataString != ToHSAMetadataString) {
  64     errs() << "Original input: " << HSAMetadataString << '\n'
  65            << "Produced output: " << ToHSAMetadataString << '\n';
  66   }
  67 }
  68
  69 AccessQualifier
  70 MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
  71   if (AccQual.empty())
  72     return AccessQualifier::Unknown;
  73
  74   return StringSwitch<AccessQualifier>(AccQual)
  75              .Case("read_only",  AccessQualifier::ReadOnly)
  76              .Case("write_only", AccessQualifier::WriteOnly)
  77              .Case("read_write", AccessQualifier::ReadWrite)
  78              .Default(AccessQualifier::Default);
  79 }
  80
  81 AddressSpaceQualifier
  82 MetadataStreamerV2::getAddressSpaceQualifier(
  83     unsigned AddressSpace) const {
  84   switch (AddressSpace) {
  85   case AMDGPUAS::PRIVATE_ADDRESS:
  86     return AddressSpaceQualifier::Private;
  87   case AMDGPUAS::GLOBAL_ADDRESS:
  88     return AddressSpaceQualifier::Global;
  89   case AMDGPUAS::CONSTANT_ADDRESS:
  90     return AddressSpaceQualifier::Constant;
  91   case AMDGPUAS::LOCAL_ADDRESS:
  92     return AddressSpaceQualifier::Local;
  93   case AMDGPUAS::FLAT_ADDRESS:
  94     return AddressSpaceQualifier::Generic;
  95   case AMDGPUAS::REGION_ADDRESS:
  96     return AddressSpaceQualifier::Region;
  97   default:
  98     return AddressSpaceQualifier::Unknown;
  99   }
 100 }
 101
 102 ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual,
 103                                            StringRef BaseTypeName) const {
 104   if (TypeQual.find("pipe") != StringRef::npos)
 105     return ValueKind::Pipe;
 106
 107   return StringSwitch<ValueKind>(BaseTypeName)
 108              .Case("image1d_t", ValueKind::Image)
 109              .Case("image1d_array_t", ValueKind::Image)
 110              .Case("image1d_buffer_t", ValueKind::Image)
 111              .Case("image2d_t", ValueKind::Image)
 112              .Case("image2d_array_t", ValueKind::Image)
 113              .Case("image2d_array_depth_t", ValueKind::Image)
 114              .Case("image2d_array_msaa_t", ValueKind::Image)
 115              .Case("image2d_array_msaa_depth_t", ValueKind::Image)
 116              .Case("image2d_depth_t", ValueKind::Image)
 117              .Case("image2d_msaa_t", ValueKind::Image)
 118              .Case("image2d_msaa_depth_t", ValueKind::Image)
 119              .Case("image3d_t", ValueKind::Image)
 120              .Case("sampler_t", ValueKind::Sampler)
 121              .Case("queue_t", ValueKind::Queue)
 122              .Default(isa<PointerType>(Ty) ?
 123                           (Ty->getPointerAddressSpace() ==
 124                            AMDGPUAS::LOCAL_ADDRESS ?
 125                            ValueKind::DynamicSharedPointer :
 126                            ValueKind::GlobalBuffer) :
 127                       ValueKind::ByValue);
 128 }
 129
 130 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:
 136       return Signed ? ValueType::I8 : ValueType::U8;
 137     case 16:
 138       return Signed ? ValueType::I16 : ValueType::U16;
 139     case 32:
 140       return Signed ? ValueType::I32 : ValueType::U32;
 141     case 64:
 142       return Signed ? ValueType::I64 : ValueType::U64;
 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,
 361                                        ValueKind ValueKind,
 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(),
 416                                       AMDGPUAS::GLOBAL_ADDRESS);
 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
 444 bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
 445   return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
 446 }
 447
 448 void MetadataStreamerV2::begin(const Module &Mod) {
 449   emitVersion();
 450   emitPrintf(Mod);
 451 }
 452
 453 void MetadataStreamerV2::end() {
 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
 464 void MetadataStreamerV2::emitKernel(const MachineFunction &MF,
 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
 514 Optional<StringRef>
 515 MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
 516   return StringSwitch<Optional<StringRef>>(AccQual)
 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
 523 Optional<StringRef>
 524 MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
 525   switch (AddressSpace) {
 526   case AMDGPUAS::PRIVATE_ADDRESS:
 527     return StringRef("private");
 528   case AMDGPUAS::GLOBAL_ADDRESS:
 529     return StringRef("global");
 530   case AMDGPUAS::CONSTANT_ADDRESS:
 531     return StringRef("constant");
 532   case AMDGPUAS::LOCAL_ADDRESS:
 533     return StringRef("local");
 534   case AMDGPUAS::FLAT_ADDRESS:
 535     return StringRef("generic");
 536   case AMDGPUAS::REGION_ADDRESS:
 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)
 564                    ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
 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
 639 msgpack::ArrayDocNode
 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,
 724                                        msgpack::ArrayDocNode Args) {
 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
 736   StringRef TypeName;
 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
 746   StringRef AccQual;
 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,
 781                                        msgpack::ArrayDocNode Args,
 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 =
 849       Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 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
 877 msgpack::MapDocNode
 878 MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
 879                                       const SIProgramInfo &ProgramInfo) const {
 880   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
 881   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
 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
 909 bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
 910   return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
 911 }
 912
 913 void MetadataStreamerV3::begin(const Module &Mod) {
 914   emitVersion();
 915   emitPrintf(Mod);
 916   getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
 917 }
 918
 919 void MetadataStreamerV3::end() {
 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
 930 void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
 931                                     const SIProgramInfo &ProgramInfo) {
 932   auto &Func = MF.getFunction();
 933   auto Kern = getHSAKernelProps(MF, ProgramInfo);
 934
 935   assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
 936          Func.getCallingConv() == CallingConv::SPIR_KERNEL);
 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