Index: llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp
@@ -90,13 +90,13 @@
     : 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());
     }
   }
 }
Index: llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
+++ llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h
@@ -60,8 +60,7 @@
                                     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>();
@@ -116,8 +115,8 @@
   }
 
 public:
-  MetadataStreamerV3() = default;
-  ~MetadataStreamerV3() = default;
+  MetadataStreamerMsgpackV3() = default;
+  ~MetadataStreamerMsgpackV3() = default;
 
   bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override;
 
@@ -130,34 +129,31 @@
                   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;
 
@@ -213,8 +209,8 @@
   }
 
 public:
-  MetadataStreamerV2() = default;
-  ~MetadataStreamerV2() = default;
+  MetadataStreamerYamlV2() = default;
+  ~MetadataStreamerYamlV2() = default;
 
   bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override;
 
Index: llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
@@ -49,13 +49,13 @@
 namespace HSAMD {
 
 //===----------------------------------------------------------------------===//
-// HSAMetadataStreamerV2
+// HSAMetadataStreamerYamlV2
 //===----------------------------------------------------------------------===//
-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;
@@ -79,7 +79,7 @@
 }
 
 AccessQualifier
-MetadataStreamerV2::getAccessQualifier(StringRef AccQual) const {
+MetadataStreamerYamlV2::getAccessQualifier(StringRef AccQual) const {
   if (AccQual.empty())
     return AccessQualifier::Unknown;
 
@@ -91,8 +91,7 @@
 }
 
 AddressSpaceQualifier
-MetadataStreamerV2::getAddressSpaceQualifier(
-    unsigned AddressSpace) const {
+MetadataStreamerYamlV2::getAddressSpaceQualifier(unsigned AddressSpace) const {
   switch (AddressSpace) {
   case AMDGPUAS::PRIVATE_ADDRESS:
     return AddressSpaceQualifier::Private;
@@ -111,8 +110,8 @@
   }
 }
 
-ValueKind MetadataStreamerV2::getValueKind(Type *Ty, StringRef TypeQual,
-                                           StringRef BaseTypeName) const {
+ValueKind MetadataStreamerYamlV2::getValueKind(Type *Ty, StringRef TypeQual,
+                                               StringRef BaseTypeName) const {
   if (TypeQual.contains("pipe"))
     return ValueKind::Pipe;
 
@@ -139,7 +138,7 @@
                       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)
@@ -177,7 +176,7 @@
 }
 
 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;
@@ -187,9 +186,8 @@
   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;
@@ -218,20 +216,19 @@
   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");
@@ -244,7 +241,7 @@
           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?
@@ -262,7 +259,7 @@
       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"))
@@ -280,15 +277,15 @@
   }
 }
 
-void MetadataStreamerV2::emitKernelArgs(const Function &Func,
-                                        const GCNSubtarget &ST) {
+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;
@@ -344,12 +341,10 @@
                 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();
 
@@ -381,8 +376,8 @@
   }
 }
 
-void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
-                                              const GCNSubtarget &ST) {
+void MetadataStreamerYamlV2::emitHiddenKernelArgs(const Function &Func,
+                                                  const GCNSubtarget &ST) {
   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
   if (!HiddenArgNumBytes)
     return;
@@ -433,17 +428,17 @@
   }
 }
 
-bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
+bool MetadataStreamerYamlV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
   return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
 }
 
-void MetadataStreamerV2::begin(const Module &Mod,
-                               const IsaInfo::AMDGPUTargetID &TargetID) {
+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;
@@ -454,8 +449,8 @@
     verify(HSAMetadataString);
 }
 
-void MetadataStreamerV2::emitKernel(const MachineFunction &MF,
-                                    const SIProgramInfo &ProgramInfo) {
+void MetadataStreamerYamlV2::emitKernel(const MachineFunction &MF,
+                                        const SIProgramInfo &ProgramInfo) {
   auto &Func = MF.getFunction();
   if (Func.getCallingConv() != CallingConv::AMDGPU_KERNEL)
     return;
@@ -480,11 +475,11 @@
 // 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;
@@ -506,7 +501,7 @@
 }
 
 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"))
@@ -514,8 +509,8 @@
       .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");
@@ -534,8 +529,9 @@
   }
 }
 
-StringRef MetadataStreamerV3::getValueKind(Type *Ty, StringRef TypeQual,
-                                           StringRef BaseTypeName) const {
+StringRef
+MetadataStreamerMsgpackV3::getValueKind(Type *Ty, StringRef TypeQual,
+                                        StringRef BaseTypeName) const {
   if (TypeQual.contains("pipe"))
     return "pipe";
 
@@ -561,7 +557,8 @@
                    : "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)
@@ -599,7 +596,7 @@
 }
 
 msgpack::ArrayDocNode
-MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
+MetadataStreamerMsgpackV3::getWorkGroupDimensions(MDNode *Node) const {
   auto Dims = HSAMetadataDoc->getArrayNode();
   if (Node->getNumOperands() != 3)
     return Dims;
@@ -610,14 +607,14 @@
   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;
@@ -630,8 +627,8 @@
   getRootMetadata("amdhsa.printf") = Printf;
 }
 
-void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
-                                            msgpack::MapDocNode Kern) {
+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())
@@ -649,8 +646,8 @@
   Kern[".language_version"] = LanguageVersion;
 }
 
-void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
-                                         msgpack::MapDocNode Kern) {
+void MetadataStreamerMsgpackV3::emitKernelAttrs(const Function &Func,
+                                                msgpack::MapDocNode Kern) {
 
   if (auto Node = Func.getMetadata("reqd_work_group_size"))
     Kern[".reqd_workgroup_size"] = getWorkGroupDimensions(Node);
@@ -674,8 +671,8 @@
     Kern[".kind"] = Kern.getDocument()->getNode("fini");
 }
 
-void MetadataStreamerV3::emitKernelArgs(const MachineFunction &MF,
-                                        msgpack::MapDocNode Kern) {
+void MetadataStreamerMsgpackV3::emitKernelArgs(const MachineFunction &MF,
+                                               msgpack::MapDocNode Kern) {
   auto &Func = MF.getFunction();
   unsigned Offset = 0;
   auto Args = HSAMetadataDoc->getArrayNode();
@@ -687,8 +684,9 @@
   Kern[".args"] = Args;
 }
 
-void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
-                                       msgpack::ArrayDocNode Args) {
+void MetadataStreamerMsgpackV3::emitKernelArg(const Argument &Arg,
+                                              unsigned &Offset,
+                                              msgpack::ArrayDocNode Args) {
   auto Func = Arg.getParent();
   auto ArgNo = Arg.getArgNo();
   const MDNode *Node;
@@ -746,7 +744,7 @@
                 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,
@@ -794,9 +792,8 @@
   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>();
 
@@ -862,9 +859,8 @@
   }
 }
 
-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();
@@ -904,18 +900,18 @@
   return Kern;
 }
 
-bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
+bool MetadataStreamerMsgpackV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
   return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
 }
 
-void MetadataStreamerV3::begin(const Module &Mod,
-                               const IsaInfo::AMDGPUTargetID &TargetID) {
+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);
@@ -926,8 +922,8 @@
     verify(StrOS.str());
 }
 
-void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
-                                    const SIProgramInfo &ProgramInfo) {
+void MetadataStreamerMsgpackV3::emitKernel(const MachineFunction &MF,
+                                           const SIProgramInfo &ProgramInfo) {
   auto &Func = MF.getFunction();
   auto Kern = getHSAKernelProps(MF, ProgramInfo);
 
@@ -953,20 +949,21 @@
 // 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,
-                               const IsaInfo::AMDGPUTargetID &TargetID) {
+void MetadataStreamerMsgpackV4::begin(const Module &Mod,
+                                      const IsaInfo::AMDGPUTargetID &TargetID) {
   emitVersion();
   emitTargetID(TargetID);
   emitPrintf(Mod);
@@ -977,16 +974,15 @@
 // 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>();