[AMDGPU/Metadata] Rename HSAMD::MetadataStreamer classes

[raghavmedicherla]

Summary

Renamed all HSAMD::MetadataStreamer classes to improve readability of the code.

[scott.linder]

LGTM, thank you!

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

TargetRegistry::RegisterAsmPrinter(getTheGCNTarget(),
                                      createAMDGPUAsmPrinterPass);
 }
 
 AMDGPUAsmPrinter::AMDGPUAsmPrinter(TargetMachine &TM,
                                    std::unique_ptr<MCStreamer> Streamer)
     : AsmPrinter(TM, std::move(Streamer)) {
   if (TM.getTargetTriple().getOS() == Triple::AMDHSA) {
     if (isHsaAbiVersion2(getGlobalSTI())) {
-      HSAMetadataStream.reset(new HSAMD::MetadataStreamerV2());
+      HSAMetadataStream.reset(new HSAMD::MetadataStreamerYamlV2());
     } else if (isHsaAbiVersion3(getGlobalSTI())) {
-      HSAMetadataStream.reset(new HSAMD::MetadataStreamerV3());
+      HSAMetadataStream.reset(new HSAMD::MetadataStreamerMsgPackV3());
     } else if (isHsaAbiVersion5(getGlobalSTI())) {
-      HSAMetadataStream.reset(new HSAMD::MetadataStreamerV5());
+      HSAMetadataStream.reset(new HSAMD::MetadataStreamerMsgPackV5());
     } else {
-      HSAMetadataStream.reset(new HSAMD::MetadataStreamerV4());
+      HSAMetadataStream.reset(new HSAMD::MetadataStreamerMsgPackV4());
     }
   }
 }
 
 StringRef AMDGPUAsmPrinter::getPassName() const {
   return "AMDGPU Assembly Printer";
 }
 

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h

virtual void emitKernel(const MachineFunction &MF,
                           const SIProgramInfo &ProgramInfo) = 0;
 
 protected:
   virtual void emitVersion() = 0;
   virtual void emitHiddenKernelArgs(const MachineFunction &MF, unsigned &Offset,
                                     msgpack::ArrayDocNode Args) = 0;
 };
 
-// TODO: Rename MetadataStreamerV3 -> MetadataStreamerMsgPackV3.
-class MetadataStreamerV3 : public MetadataStreamer {
+class MetadataStreamerMsgPackV3 : public MetadataStreamer {
 protected:
   std::unique_ptr<msgpack::Document> HSAMetadataDoc =
       std::make_unique<msgpack::Document>();
 
   void dump(StringRef HSAMetadataString) const;
 
   void verify(StringRef HSAMetadataString) const;
 
msgpack::DocNode &getRootMetadata(StringRef Key) {
     return HSAMetadataDoc->getRoot().getMap(/*Convert=*/true)[Key];
   }
 
   msgpack::DocNode &getHSAMetadataRoot() {
     return HSAMetadataDoc->getRoot();
   }
 
 public:
-  MetadataStreamerV3() = default;
-  ~MetadataStreamerV3() = default;
+  MetadataStreamerMsgPackV3() = default;
+  ~MetadataStreamerMsgPackV3() = default;
 
   bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override;
 
   void begin(const Module &Mod,
              const IsaInfo::AMDGPUTargetID &TargetID) override;
 
   void end() override;
 
   void emitKernel(const MachineFunction &MF,
                   const SIProgramInfo &ProgramInfo) override;
 };
 
-// TODO: Rename MetadataStreamerV4 -> MetadataStreamerMsgPackV4.
-class MetadataStreamerV4 : public MetadataStreamerV3 {
+class MetadataStreamerMsgPackV4 : public MetadataStreamerMsgPackV3 {
 protected:
   void emitVersion() override;
   void emitTargetID(const IsaInfo::AMDGPUTargetID &TargetID);
 
 public:
-  MetadataStreamerV4() = default;
-  ~MetadataStreamerV4() = default;
+  MetadataStreamerMsgPackV4() = default;
+  ~MetadataStreamerMsgPackV4() = default;
 
   void begin(const Module &Mod,
              const IsaInfo::AMDGPUTargetID &TargetID) override;
 };
 
-// TODO: Rename MetadataStreamerV5 -> MetadataStreamerMsgPackV5.
-class MetadataStreamerV5 final : public MetadataStreamerV4 {
+class MetadataStreamerMsgPackV5 final : public MetadataStreamerMsgPackV4 {
 protected:
   void emitVersion() override;
   void emitHiddenKernelArgs(const MachineFunction &MF, unsigned &Offset,
                             msgpack::ArrayDocNode Args) override;
 
 public:
-  MetadataStreamerV5() = default;
-  ~MetadataStreamerV5() = default;
+  MetadataStreamerMsgPackV5() = default;
+  ~MetadataStreamerMsgPackV5() = default;
 };
 
 // TODO: Rename MetadataStreamerV2 -> MetadataStreamerYamlV2.
-class MetadataStreamerV2 final : public MetadataStreamer {
+class MetadataStreamerYamlV2 final : public MetadataStreamer {
 private:
   Metadata HSAMetadata;
 
   void dump(StringRef HSAMetadataString) const;
 
   void verify(StringRef HSAMetadataString) const;
 
   AccessQualifier getAccessQualifier(StringRef AccQual) const;

 protected:
   void emitVersion() override;
   void emitHiddenKernelArgs(const MachineFunction &MF, unsigned &Offset,
                             msgpack::ArrayDocNode Args) override {
     llvm_unreachable("Dummy override should not be invoked!");
   }
 
 public:
-  MetadataStreamerV2() = default;
-  ~MetadataStreamerV2() = default;
+  MetadataStreamerYamlV2() = default;
+  ~MetadataStreamerYamlV2() = default;
 
   bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override;
 
   void begin(const Module &Mod,
              const IsaInfo::AMDGPUTargetID &TargetID) override;
 
   void end() override;
 

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

static cl::opt<bool> VerifyHSAMetadata(
     cl::desc("Verify AMDGPU HSA Metadata"));
 
 namespace AMDGPU {
 namespace HSAMD {
 
 //===----------------------------------------------------------------------===//
 // HSAMetadataStreamerV2
 //===----------------------------------------------------------------------===//
-void MetadataStreamerV2::dump(StringRef HSAMetadataString) const {
+void MetadataStreamerYamlV2::dump(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
 }
 
-void MetadataStreamerV2::verify(StringRef HSAMetadataString) const {
+void MetadataStreamerYamlV2::verify(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata Parser Test: ";
 
   HSAMD::Metadata FromHSAMetadataString;
   if (fromString(HSAMetadataString, FromHSAMetadataString)) {
     errs() << "FAIL\n";
     return;
   }
 
   std::string ToHSAMetadataString;
   if (toString(FromHSAMetadataString, ToHSAMetadataString)) {
     errs() << "FAIL\n";
     return;
   }
 
   errs() << (HSAMetadataString == ToHSAMetadataString ? "PASS" : "FAIL")
          << '\n';
   if (HSAMetadataString != ToHSAMetadataString) {
     errs() << "Original input: " << HSAMetadataString << '\n'
            << "Produced output: " << ToHSAMetadataString << '\n';
   }
 }
 
 AccessQualifier
-MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
+MetadataStreamerYamlV2::getAccessQualifier(StringRef AccQual) const {
   if (AccQual.empty())
     return AccessQualifier::Unknown;
 
   return StringSwitch<AccessQualifier>(AccQual)
              .Case("read_only",  AccessQualifier::ReadOnly)
              .Case("write_only", AccessQualifier::WriteOnly)
              .Case("read_write", AccessQualifier::ReadWrite)
              .Default(AccessQualifier::Default);
 }
 
 AddressSpaceQualifier
-MetadataStreamerV2::getAddressSpaceQualifier(
-    unsigned AddressSpace) const {
+MetadataStreamerYamlV2::getAddressSpaceQualifier(unsigned AddressSpace) const {
   switch (AddressSpace) {
   case AMDGPUAS::PRIVATE_ADDRESS:
     return AddressSpaceQualifier::Private;
   case AMDGPUAS::GLOBAL_ADDRESS:
     return AddressSpaceQualifier::Global;
   case AMDGPUAS::CONSTANT_ADDRESS:
     return AddressSpaceQualifier::Constant;
   case AMDGPUAS::LOCAL_ADDRESS:
     return AddressSpaceQualifier::Local;
   case AMDGPUAS::FLAT_ADDRESS:
     return AddressSpaceQualifier::Generic;
   case AMDGPUAS::REGION_ADDRESS:
     return AddressSpaceQualifier::Region;
   default:
     return AddressSpaceQualifier::Unknown;
   }
 }
 
-ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual,
+ValueKind MetadataStreamerYamlV2::getValueKind(Type *Ty, StringRef TypeQual,
                                                StringRef BaseTypeName) const {
   if (TypeQual.contains("pipe"))
     return ValueKind::Pipe;
 
   return StringSwitch<ValueKind>(BaseTypeName)
              .Case("image1d_t", ValueKind::Image)
              .Case("image1d_array_t", ValueKind::Image)
              .Case("image1d_buffer_t", ValueKind::Image)
              .Case("image2d_t", ValueKind::Image)
return StringSwitch<ValueKind>(BaseTypeName)
              .Default(isa<PointerType>(Ty) ?
                           (Ty->getPointerAddressSpace() ==
                            AMDGPUAS::LOCAL_ADDRESS ?
                            ValueKind::DynamicSharedPointer :
                            ValueKind::GlobalBuffer) :
                       ValueKind::ByValue);
 }
 
-std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const {
+std::string MetadataStreamerYamlV2::getTypeName(Type *Ty, bool Signed) const {
   switch (Ty->getTypeID()) {
   case Type::IntegerTyID: {
     if (!Signed)
       return (Twine('u') + getTypeName(Ty, true)).str();
 
     auto BitWidth = Ty->getIntegerBitWidth();
     switch (BitWidth) {
     case 8:
case Type::FixedVectorTyID: {
     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
   }
   default:
     return "unknown";
   }
 }
 
 std::vector<uint32_t>
-MetadataStreamerV2::getWorkGroupDimensions(MDNode *Node) const {
+MetadataStreamerYamlV2::getWorkGroupDimensions(MDNode *Node) const {
   std::vector<uint32_t> Dims;
   if (Node->getNumOperands() != 3)
     return Dims;
 
   for (auto &Op : Node->operands())
     Dims.push_back(mdconst::extract<ConstantInt>(Op)->getZExtValue());
   return Dims;
 }
 
-Kernel::CodeProps::Metadata
-MetadataStreamerV2::getHSACodeProps(const MachineFunction &MF,
-                                    const SIProgramInfo &ProgramInfo) const {
+Kernel::CodeProps::Metadata MetadataStreamerYamlV2::getHSACodeProps(
+    const MachineFunction &MF, const SIProgramInfo &ProgramInfo) const {
   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
   HSAMD::Kernel::CodeProps::Metadata HSACodeProps;
   const Function &F = MF.getFunction();
 
   assert(F.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
          F.getCallingConv() == CallingConv::SPIR_KERNEL);
 
Kernel::CodeProps::Metadata MetadataStreamerYamlV2::getHSACodeProps(
   HSACodeProps.mIsDynamicCallStack = ProgramInfo.DynamicCallStack;
   HSACodeProps.mIsXNACKEnabled = STM.isXNACKEnabled();
   HSACodeProps.mNumSpilledSGPRs = MFI.getNumSpilledSGPRs();
   HSACodeProps.mNumSpilledVGPRs = MFI.getNumSpilledVGPRs();
 
   return HSACodeProps;
 }
 
-Kernel::DebugProps::Metadata
-MetadataStreamerV2::getHSADebugProps(const MachineFunction &MF,
-                                     const SIProgramInfo &ProgramInfo) const {
+Kernel::DebugProps::Metadata MetadataStreamerYamlV2::getHSADebugProps(
+    const MachineFunction &MF, const SIProgramInfo &ProgramInfo) const {
   return HSAMD::Kernel::DebugProps::Metadata();
 }
 
-void MetadataStreamerV2::emitVersion() {
+void MetadataStreamerYamlV2::emitVersion() {
   auto &Version = HSAMetadata.mVersion;
 
   Version.push_back(VersionMajorV2);
   Version.push_back(VersionMinorV2);
 }
 
-void MetadataStreamerV2::emitPrintf(const Module &Mod) {
+void MetadataStreamerYamlV2::emitPrintf(const Module &Mod) {
   auto &Printf = HSAMetadata.mPrintf;
 
   auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
   if (!Node)
     return;
 
   for (auto *Op : Node->operands())
     if (Op->getNumOperands())
       Printf.push_back(
           std::string(cast<MDString>(Op->getOperand(0))->getString()));
 }
 
-void MetadataStreamerV2::emitKernelLanguage(const Function &Func) {
+void MetadataStreamerYamlV2::emitKernelLanguage(const Function &Func) {
   auto &Kernel = HSAMetadata.mKernels.back();
 
   // TODO: What about other languages?
   auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
   if (!Node || !Node->getNumOperands())
     return;
   auto Op0 = Node->getOperand(0);
   if (Op0->getNumOperands() <= 1)
     return;
 
   Kernel.mLanguage = "OpenCL C";
   Kernel.mLanguageVersion.push_back(
       mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue());
   Kernel.mLanguageVersion.push_back(
       mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue());
 }
 
-void MetadataStreamerV2::emitKernelAttrs(const Function &Func) {
+void MetadataStreamerYamlV2::emitKernelAttrs(const Function &Func) {
   auto &Attrs = HSAMetadata.mKernels.back().mAttrs;
 
   if (auto Node = Func.getMetadata("reqd_work_group_size"))
     Attrs.mReqdWorkGroupSize = getWorkGroupDimensions(Node);
   if (auto Node = Func.getMetadata("work_group_size_hint"))
     Attrs.mWorkGroupSizeHint = getWorkGroupDimensions(Node);
   if (auto Node = Func.getMetadata("vec_type_hint")) {
     Attrs.mVecTypeHint = getTypeName(
         cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
         mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue());
   }
   if (Func.hasFnAttribute("runtime-handle")) {
     Attrs.mRuntimeHandle =
         Func.getFnAttribute("runtime-handle").getValueAsString().str();
   }
 }
 
-void MetadataStreamerV2::emitKernelArgs(const Function &Func,
+void MetadataStreamerYamlV2::emitKernelArgs(const Function &Func,
                                             const GCNSubtarget &ST) {
   for (auto &Arg : Func.args())
     emitKernelArg(Arg);
 
   emitHiddenKernelArgs(Func, ST);
 }
 
-void MetadataStreamerV2::emitKernelArg(const Argument &Arg) {
+void MetadataStreamerYamlV2::emitKernelArg(const Argument &Arg) {
   auto Func = Arg.getParent();
   auto ArgNo = Arg.getArgNo();
   const MDNode *Node;
 
   StringRef Name;
   Node = Func->getMetadata("kernel_arg_name");
   if (Node && ArgNo < Node->getNumOperands())
     Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
void MetadataStreamerYamlV2::emitKernelArg(const Argument &Arg) {
   Align ArgAlign;
   std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
 
   emitKernelArg(DL, ArgTy, ArgAlign,
                 getValueKind(ArgTy, TypeQual, BaseTypeName), PointeeAlign, Name,
                 TypeName, BaseTypeName, AccQual, TypeQual);
 }
 
-void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
-                                       Align Alignment, ValueKind ValueKind,
-                                       MaybeAlign PointeeAlign, StringRef Name,
-                                       StringRef TypeName,
-                                       StringRef BaseTypeName,
-                                       StringRef AccQual, StringRef TypeQual) {
+void MetadataStreamerYamlV2::emitKernelArg(
+    const DataLayout &DL, Type *Ty, Align Alignment, ValueKind ValueKind,
+    MaybeAlign PointeeAlign, StringRef Name, StringRef TypeName,
+    StringRef BaseTypeName, StringRef AccQual, StringRef TypeQual) {
   HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
   auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
 
   Arg.mName = std::string(Name);
   Arg.mTypeName = std::string(TypeName);
   Arg.mSize = DL.getTypeAllocSize(Ty);
   Arg.mAlign = Alignment.value();
   Arg.mValueKind = ValueKind;
auto P = StringSwitch<bool*>(Key)
                  .Case("volatile", &Arg.mIsVolatile)
                  .Case("pipe",     &Arg.mIsPipe)
                  .Default(nullptr);
     if (P)
       *P = true;
   }
 }
 
-void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
+void MetadataStreamerYamlV2::emitHiddenKernelArgs(const Function &Func,
                                                   const GCNSubtarget &ST) {
   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
   if (!HiddenArgNumBytes)
     return;
 
   auto &DL = Func.getParent()->getDataLayout();
   auto Int64Ty = Type::getInt64Ty(Func.getContext());
 
   if (HiddenArgNumBytes >= 8)
void MetadataStreamerYamlV2::emitHiddenKernelArgs(const Function &Func,
   if (HiddenArgNumBytes >= 56) {
     if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg"))
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenMultiGridSyncArg);
     else
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
   }
 }
 
-bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
+bool MetadataStreamerYamlV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
   return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
 }
 
-void MetadataStreamerV2::begin(const Module &Mod,
+void MetadataStreamerYamlV2::begin(const Module &Mod,
                                    const IsaInfo::AMDGPUTargetID &TargetID) {
   emitVersion();
   emitPrintf(Mod);
 }
 
-void MetadataStreamerV2::end() {
+void MetadataStreamerYamlV2::end() {
   std::string HSAMetadataString;
   if (toString(HSAMetadata, HSAMetadataString))
     return;
 
   if (DumpHSAMetadata)
     dump(HSAMetadataString);
   if (VerifyHSAMetadata)
     verify(HSAMetadataString);
 }
 
-void MetadataStreamerV2::emitKernel(const MachineFunction &MF,
+void MetadataStreamerYamlV2::emitKernel(const MachineFunction &MF,
                                         const SIProgramInfo &ProgramInfo) {
   auto &Func = MF.getFunction();
   if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
     return;
 
   auto CodeProps = getHSACodeProps(MF, ProgramInfo);
   auto DebugProps = getHSADebugProps(MF, ProgramInfo);
 
   HSAMetadata.mKernels.push_back(Kernel::Metadata());
   auto &Kernel = HSAMetadata.mKernels.back();
 
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
   Kernel.mName = std::string(Func.getName());
   Kernel.mSymbolName = (Twine(Func.getName()) + Twine("@kd")).str();
   emitKernelLanguage(Func);
   emitKernelAttrs(Func);
   emitKernelArgs(Func, ST);
   HSAMetadata.mKernels.back().mCodeProps = CodeProps;
   HSAMetadata.mKernels.back().mDebugProps = DebugProps;
 }
 
 //===----------------------------------------------------------------------===//
 // HSAMetadataStreamerV3
 //===----------------------------------------------------------------------===//
 
-void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
+void MetadataStreamerMsgPackV3::dump(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
 }
 
-void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
+void MetadataStreamerMsgPackV3::verify(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata Parser Test: ";
 
   msgpack::Document FromHSAMetadataString;
 
   if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
     errs() << "FAIL\n";
     return;
   }
 
   std::string ToHSAMetadataString;
   raw_string_ostream StrOS(ToHSAMetadataString);
   FromHSAMetadataString.toYAML(StrOS);
 
   errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
   if (HSAMetadataString != ToHSAMetadataString) {
     errs() << "Original input: " << HSAMetadataString << '\n'
            << "Produced output: " << StrOS.str() << '\n';
   }
 }
 
 Optional<StringRef>
-MetadataStreamerV3::getAccessQualifier(StringRef AccQual) const {
+MetadataStreamerMsgPackV3::getAccessQualifier(StringRef AccQual) const {
   return StringSwitch<Optional<StringRef>>(AccQual)
       .Case("read_only", StringRef("read_only"))
       .Case("write_only", StringRef("write_only"))
       .Case("read_write", StringRef("read_write"))
       .Default(None);
 }
 
-Optional<StringRef>
-MetadataStreamerV3::getAddressSpaceQualifier(unsigned AddressSpace) const {
+Optional<StringRef> MetadataStreamerMsgPackV3::getAddressSpaceQualifier(
+    unsigned AddressSpace) const {
   switch (AddressSpace) {
   case AMDGPUAS::PRIVATE_ADDRESS:
     return StringRef("private");
   case AMDGPUAS::GLOBAL_ADDRESS:
     return StringRef("global");
   case AMDGPUAS::CONSTANT_ADDRESS:
     return StringRef("constant");
   case AMDGPUAS::LOCAL_ADDRESS:
     return StringRef("local");
   case AMDGPUAS::FLAT_ADDRESS:
     return StringRef("generic");
   case AMDGPUAS::REGION_ADDRESS:
     return StringRef("region");
   default:
     return None;
   }
 }
 
-StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
-                                           StringRef BaseTypeName) const {
+StringRef
+MetadataStreamerMsgPackV3::getValueKind(Type *Ty, StringRef TypeQual,
   if (TypeQual.contains("pipe"))
     return "pipe";
 
   return StringSwitch<StringRef>(BaseTypeName)
       .Case("image1d_t", "image")
       .Case("image1d_array_t", "image")
       .Case("image1d_buffer_t", "image")
       .Case("image2d_t", "image")
return StringSwitch<StringRef>(BaseTypeName)
       .Case("queue_t", "queue")
       .Default(isa<PointerType>(Ty)
                    ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
                           ? "dynamic_shared_pointer"
                           : "global_buffer")
                    : "by_value");
 }
 
-std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
+std::string MetadataStreamerMsgPackV3::getTypeName(Type *Ty,
+                                                   bool Signed) const {
   switch (Ty->getTypeID()) {
   case Type::IntegerTyID: {
     if (!Signed)
       return (Twine('u') + getTypeName(Ty, true)).str();
 
     auto BitWidth = Ty->getIntegerBitWidth();
     switch (BitWidth) {
     case 8:
case Type::FixedVectorTyID: {
     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
   }
   default:
     return "unknown";
   }
 }
 
 msgpack::ArrayDocNode
-MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
+MetadataStreamerMsgPackV3::getWorkGroupDimensions(MDNode *Node) const {
   auto Dims = HSAMetadataDoc->getArrayNode();
   if (Node->getNumOperands() != 3)
     return Dims;
 
   for (auto &Op : Node->operands())
     Dims.push_back(Dims.getDocument()->getNode(
         uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
   return Dims;
 }
 
-void MetadataStreamerV3::emitVersion() {
+void MetadataStreamerMsgPackV3::emitVersion() {
   auto Version = HSAMetadataDoc->getArrayNode();
   Version.push_back(Version.getDocument()->getNode(VersionMajorV3));
   Version.push_back(Version.getDocument()->getNode(VersionMinorV3));
   getRootMetadata("amdhsa.version") = Version;
 }
 
-void MetadataStreamerV3::emitPrintf(const Module &Mod) {
+void MetadataStreamerMsgPackV3::emitPrintf(const Module &Mod) {
   auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
   if (!Node)
     return;
 
   auto Printf = HSAMetadataDoc->getArrayNode();
   for (auto *Op : Node->operands())
     if (Op->getNumOperands())
       Printf.push_back(Printf.getDocument()->getNode(
           cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
   getRootMetadata("amdhsa.printf") = Printf;
 }
 
-void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
+void MetadataStreamerMsgPackV3::emitKernelLanguage(const Function &Func,
                                                    msgpack::MapDocNode Kern) {
   // TODO: What about other languages?
   auto Node = Func.getParent()->getNamedMetadata("opencl.ocl.version");
   if (!Node || !Node->getNumOperands())
     return;
   auto Op0 = Node->getOperand(0);
   if (Op0->getNumOperands() <= 1)
     return;
 
   Kern[".language"] = Kern.getDocument()->getNode("OpenCL C");
   auto LanguageVersion = Kern.getDocument()->getArrayNode();
   LanguageVersion.push_back(Kern.getDocument()->getNode(
       mdconst::extract<ConstantInt>(Op0->getOperand(0))->getZExtValue()));
   LanguageVersion.push_back(Kern.getDocument()->getNode(
       mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
   Kern[".language_version"] = LanguageVersion;
 }
 
-void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
+void MetadataStreamerMsgPackV3::emitKernelAttrs(const Function &Func,
                                                 msgpack::MapDocNode Kern) {
 
   if (auto Node = Func.getMetadata("reqd_work_group_size"))
     Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
   if (auto Node = Func.getMetadata("work_group_size_hint"))
     Kern[".workgroup_size_hint"] = getWorkGroupDimensions(Node);
   if (auto Node = Func.getMetadata("vec_type_hint")) {
     Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
         getTypeName(
             cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
             mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
         /*Copy=*/true);
   }
   if (Func.hasFnAttribute("runtime-handle")) {
     Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
         Func.getFnAttribute("runtime-handle").getValueAsString().str(),
         /*Copy=*/true);
   }
   if (Func.hasFnAttribute("device-init"))
     Kern[".kind"] = Kern.getDocument()->getNode("init");
   else if (Func.hasFnAttribute("device-fini"))
     Kern[".kind"] = Kern.getDocument()->getNode("fini");
 }
 
-void MetadataStreamerV3::emitKernelArgs(const MachineFunction &MF,
+void MetadataStreamerMsgPackV3::emitKernelArgs(const MachineFunction &MF,
                                                msgpack::MapDocNode Kern) {
   auto &Func = MF.getFunction();
   unsigned Offset = 0;
   auto Args = HSAMetadataDoc->getArrayNode();
   for (auto &Arg : Func.args())
     emitKernelArg(Arg, Offset, Args);
 
   emitHiddenKernelArgs(MF, Offset, Args);
 
   Kern[".args"] = Args;
 }
 
-void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
+void MetadataStreamerMsgPackV3::emitKernelArg(const Argument &Arg,
+                                              unsigned &Offset,
                                               msgpack::ArrayDocNode Args) {
   auto Func = Arg.getParent();
   auto ArgNo = Arg.getArgNo();
   const MDNode *Node;
 
   StringRef Name;
   Node = Func->getMetadata("kernel_arg_name");
   if (Node && ArgNo < Node->getNumOperands())
     Name = cast<MDString>(Node->getOperand(ArgNo))->getString();
void MetadataStreamerMsgPackV3::emitKernelArg(const Argument &Arg,
   Align ArgAlign;
   std::tie(ArgTy, ArgAlign) = getArgumentTypeAlign(Arg, DL);
 
   emitKernelArg(DL, ArgTy, ArgAlign,
                 getValueKind(ArgTy, TypeQual, BaseTypeName), Offset, Args,
                 PointeeAlign, Name, TypeName, BaseTypeName, AccQual, TypeQual);
 }
 
-void MetadataStreamerV3::emitKernelArg(
+void MetadataStreamerMsgPackV3::emitKernelArg(
     const DataLayout &DL, Type *Ty, Align Alignment, StringRef ValueKind,
     unsigned &Offset, msgpack::ArrayDocNode Args, MaybeAlign PointeeAlign,
     StringRef Name, StringRef TypeName, StringRef BaseTypeName,
     StringRef AccQual, StringRef TypeQual) {
   auto Arg = Args.getDocument()->getMapNode();
 
   if (!Name.empty())
     Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
else if (Key == "volatile")
       Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
     else if (Key == "pipe")
       Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
   }
 
   Args.push_back(Arg);
 }
 
-void MetadataStreamerV3::emitHiddenKernelArgs(const MachineFunction &MF,
-                                              unsigned &Offset,
-                                              msgpack::ArrayDocNode Args) {
+void MetadataStreamerMsgPackV3::emitHiddenKernelArgs(
+    const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
   auto &Func = MF.getFunction();
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
 
   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
   if (!HiddenArgNumBytes)
     return;
 
   const Module *M = Func.getParent();
if (!Func.hasFnAttribute("amdgpu-no-multigrid-sync-arg")) {
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset,
                     Args);
     } else {
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
     }
   }
 }
 
-msgpack::MapDocNode
-MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
-                                      const SIProgramInfo &ProgramInfo) const {
+msgpack::MapDocNode MetadataStreamerMsgPackV3::getHSAKernelProps(
+    const MachineFunction &MF, const SIProgramInfo &ProgramInfo) const {
   const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
   const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
   const Function &F = MF.getFunction();
 
   auto Kern = HSAMetadataDoc->getMapNode();
 
   Align MaxKernArgAlign;
   Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
msgpack::MapDocNode MetadataStreamerMsgPackV3::getHSAKernelProps(
   Kern[".sgpr_spill_count"] =
       Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
   Kern[".vgpr_spill_count"] =
       Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
 
   return Kern;
 }
 
-bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
+bool MetadataStreamerMsgPackV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
   return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
 }
 
-void MetadataStreamerV3::begin(const Module &Mod,
+void MetadataStreamerMsgPackV3::begin(const Module &Mod,
                                       const IsaInfo::AMDGPUTargetID &TargetID) {
   emitVersion();
   emitPrintf(Mod);
   getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
 }
 
-void MetadataStreamerV3::end() {
+void MetadataStreamerMsgPackV3::end() {
   std::string HSAMetadataString;
   raw_string_ostream StrOS(HSAMetadataString);
   HSAMetadataDoc->toYAML(StrOS);
 
   if (DumpHSAMetadata)
     dump(StrOS.str());
   if (VerifyHSAMetadata)
     verify(StrOS.str());
 }
 
-void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
+void MetadataStreamerMsgPackV3::emitKernel(const MachineFunction &MF,
                                            const SIProgramInfo &ProgramInfo) {
   auto &Func = MF.getFunction();
   auto Kern = getHSAKernelProps(MF, ProgramInfo);
 
   assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
          Func.getCallingConv() == CallingConv::SPIR_KERNEL);
 
   auto Kernels =
       getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
void MetadataStreamerMsgPackV3::emitKernel(const MachineFunction &MF,
 
   Kernels.push_back(Kern);
 }
 
 //===----------------------------------------------------------------------===//
 // HSAMetadataStreamerV4
 //===----------------------------------------------------------------------===//
 
-void MetadataStreamerV4::emitVersion() {
+void MetadataStreamerMsgPackV4::emitVersion() {
   auto Version = HSAMetadataDoc->getArrayNode();
   Version.push_back(Version.getDocument()->getNode(VersionMajorV4));
   Version.push_back(Version.getDocument()->getNode(VersionMinorV4));
   getRootMetadata("amdhsa.version") = Version;
 }
 
-void MetadataStreamerV4::emitTargetID(const IsaInfo::AMDGPUTargetID &TargetID) {
+void MetadataStreamerMsgPackV4::emitTargetID(
+    const IsaInfo::AMDGPUTargetID &TargetID) {
   getRootMetadata("amdhsa.target") =
       HSAMetadataDoc->getNode(TargetID.toString(), /*Copy=*/true);
 }
 
-void MetadataStreamerV4::begin(const Module &Mod,
+void MetadataStreamerMsgPackV4::begin(const Module &Mod,
                                       const IsaInfo::AMDGPUTargetID &TargetID) {
   emitVersion();
   emitTargetID(TargetID);
   emitPrintf(Mod);
   getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
 }
 
 //===----------------------------------------------------------------------===//
 // HSAMetadataStreamerV5
 //===----------------------------------------------------------------------===//
 
-void MetadataStreamerV5::emitVersion() {
+void MetadataStreamerMsgPackV5::emitVersion() {
   auto Version = HSAMetadataDoc->getArrayNode();
   Version.push_back(Version.getDocument()->getNode(VersionMajorV5));
   Version.push_back(Version.getDocument()->getNode(VersionMinorV5));
   getRootMetadata("amdhsa.version") = Version;
 }
 
-void MetadataStreamerV5::emitHiddenKernelArgs(const MachineFunction &MF,
-                                              unsigned &Offset,
-                                              msgpack::ArrayDocNode Args) {
+void MetadataStreamerMsgPackV5::emitHiddenKernelArgs(
+    const MachineFunction &MF, unsigned &Offset, msgpack::ArrayDocNode Args) {
   auto &Func = MF.getFunction();
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
 
   // No implicit kernel argument is used.
   if (ST.getImplicitArgNumBytes(Func) == 0)
     return;
 
   const Module *M = Func.getParent();