[AMDGPU] Switched HSA metadata to use MsgPackDocument

[tpr]

Summary

MsgPackDocument is the lighter-weight replacement for MsgPackTypes. This
commit switches AMDGPU HSA metadata processing to use MsgPackDocument
instead of MsgPackTypes.

Change-Id: I0751668013abe8c87db01db1170831a76079b3a6

[scott.linder]

Thank you for updating this! I only have one small suggestion, but otherwise it LGTM.

lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h

55

Small nit, but could this be llvm::make_unique<msgpack::Document>()?

[tpr]

V2: Addressed review comment.

[scott.linder]

LGTM

include/llvm/BinaryFormat/AMDGPUMetadataVerifier.h

 /// well-typed metadata and untyped metadata. When verifying in the non-strict
 /// mode, untyped metadata is coerced into the correct type if possible.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_BINARYFORMAT_AMDGPUMETADATAVERIFIER_H
 #define LLVM_BINARYFORMAT_AMDGPUMETADATAVERIFIER_H
 
-#include "llvm/BinaryFormat/MsgPackTypes.h"
+#include "llvm/BinaryFormat/MsgPackDocument.h"
 
 namespace llvm {
 namespace AMDGPU {
 namespace HSAMD {
 namespace V3 {
 
 /// Verifier for AMDGPU HSA metadata.
 ///
 /// Operates in two modes:
 ///
 /// In strict mode, metadata must already be well-typed.
 ///
 /// In non-strict mode, metadata is coerced into expected types when possible.
 class MetadataVerifier {
   bool Strict;
 
-  bool verifyScalar(msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind,
-                    function_ref<bool(msgpack::ScalarNode &)> verifyValue = {});
-  bool verifyInteger(msgpack::Node &Node);
-  bool verifyArray(msgpack::Node &Node,
-                   function_ref<bool(msgpack::Node &)> verifyNode,
+  bool verifyScalar(msgpack::DocNode &Node, msgpack::Type SKind,
+                    function_ref<bool(msgpack::DocNode &)> verifyValue = {});
+  bool verifyInteger(msgpack::DocNode &Node);
+  bool verifyArray(msgpack::DocNode &Node,
+                   function_ref<bool(msgpack::DocNode &)> verifyNode,
                    Optional<size_t> Size = None);
-  bool verifyEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required,
-                   function_ref<bool(msgpack::Node &)> verifyNode);
+  bool verifyEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
+                   function_ref<bool(msgpack::DocNode &)> verifyNode);
   bool
-  verifyScalarEntry(msgpack::MapNode &MapNode, StringRef Key, bool Required,
-                    msgpack::ScalarNode::ScalarKind SKind,
-                    function_ref<bool(msgpack::ScalarNode &)> verifyValue = {});
-  bool verifyIntegerEntry(msgpack::MapNode &MapNode, StringRef Key,
+  verifyScalarEntry(msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
+                    msgpack::Type SKind,
+                    function_ref<bool(msgpack::DocNode &)> verifyValue = {});
+  bool verifyIntegerEntry(msgpack::MapDocNode &MapNode, StringRef Key,
                           bool Required);
-  bool verifyKernelArgs(msgpack::Node &Node);
-  bool verifyKernel(msgpack::Node &Node);
+  bool verifyKernelArgs(msgpack::DocNode &Node);
+  bool verifyKernel(msgpack::DocNode &Node);
 
 public:
   /// Construct a MetadataVerifier, specifying whether it will operate in \p
   /// Strict mode.
   MetadataVerifier(bool Strict) : Strict(Strict) {}
 
   /// Verify given HSA metadata.
   ///
   /// \returns True when successful, false when metadata is invalid.
-  bool verify(msgpack::Node &HSAMetadataRoot);
+  bool verify(msgpack::DocNode &HSAMetadataRoot);
 };
 
 } // end namespace V3
 } // end namespace HSAMD
 } // end namespace AMDGPU
 } // end namespace llvm
 
 #endif // LLVM_BINARYFORMAT_AMDGPUMETADATAVERIFIER_H

lib/BinaryFormat/AMDGPUMetadataVerifier.cpp

 #include "llvm/Support/AMDGPUMetadata.h"
 
 namespace llvm {
 namespace AMDGPU {
 namespace HSAMD {
 namespace V3 {
 
 bool MetadataVerifier::verifyScalar(
-    msgpack::Node &Node, msgpack::ScalarNode::ScalarKind SKind,
-    function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
-  auto ScalarPtr = dyn_cast<msgpack::ScalarNode>(&Node);
-  if (!ScalarPtr)
-    return false;
-  auto &Scalar = *ScalarPtr;
-  // Do not output extraneous tags for types we know from the spec.
-  Scalar.IgnoreTag = true;
-  if (Scalar.getScalarKind() != SKind) {
+    msgpack::DocNode &Node, msgpack::Type SKind,
+    function_ref<bool(msgpack::DocNode &)> verifyValue) {
+  if (!Node.isScalar())
+    return false;
+  if (Node.getKind() != SKind) {
     if (Strict)
       return false;
     // If we are not strict, we interpret string values as "implicitly typed"
     // and attempt to coerce them to the expected type here.
-    if (Scalar.getScalarKind() != msgpack::ScalarNode::SK_String)
+    if (Node.getKind() != msgpack::Type::String)
       return false;
-    std::string StringValue = Scalar.getString();
-    Scalar.setScalarKind(SKind);
-    if (Scalar.inputYAML(StringValue) != StringRef())
+    StringRef StringValue = Node.getString();
+    Node.fromString(StringValue);
+    if (Node.getKind() != SKind)
       return false;
   }
   if (verifyValue)
-    return verifyValue(Scalar);
+    return verifyValue(Node);
   return true;
 }
 
-bool MetadataVerifier::verifyInteger(msgpack::Node &Node) {
-  if (!verifyScalar(Node, msgpack::ScalarNode::SK_UInt))
-    if (!verifyScalar(Node, msgpack::ScalarNode::SK_Int))
+bool MetadataVerifier::verifyInteger(msgpack::DocNode &Node) {
+  if (!verifyScalar(Node, msgpack::Type::UInt))
+    if (!verifyScalar(Node, msgpack::Type::Int))
       return false;
   return true;
 }
 
 bool MetadataVerifier::verifyArray(
-    msgpack::Node &Node, function_ref<bool(msgpack::Node &)> verifyNode,
+    msgpack::DocNode &Node, function_ref<bool(msgpack::DocNode &)> verifyNode,
     Optional<size_t> Size) {
-  auto ArrayPtr = dyn_cast<msgpack::ArrayNode>(&Node);
-  if (!ArrayPtr)
+  if (!Node.isArray())
     return false;
-  auto &Array = *ArrayPtr;
+  auto &Array = Node.getArray();
   if (Size && Array.size() != *Size)
     return false;
   for (auto &Item : Array)
-    if (!verifyNode(*Item.get()))
+    if (!verifyNode(Item))
       return false;
 
   return true;
 }
 
 bool MetadataVerifier::verifyEntry(
-    msgpack::MapNode &MapNode, StringRef Key, bool Required,
-    function_ref<bool(msgpack::Node &)> verifyNode) {
+    msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
+    function_ref<bool(msgpack::DocNode &)> verifyNode) {
   auto Entry = MapNode.find(Key);
   if (Entry == MapNode.end())
     return !Required;
-  return verifyNode(*Entry->second.get());
+  return verifyNode(Entry->second);
 }
 
 bool MetadataVerifier::verifyScalarEntry(
-    msgpack::MapNode &MapNode, StringRef Key, bool Required,
-    msgpack::ScalarNode::ScalarKind SKind,
-    function_ref<bool(msgpack::ScalarNode &)> verifyValue) {
-  return verifyEntry(MapNode, Key, Required, [=](msgpack::Node &Node) {
+    msgpack::MapDocNode &MapNode, StringRef Key, bool Required,
+    msgpack::Type SKind,
+    function_ref<bool(msgpack::DocNode &)> verifyValue) {
+  return verifyEntry(MapNode, Key, Required, [=](msgpack::DocNode &Node) {
     return verifyScalar(Node, SKind, verifyValue);
   });
 }
 
-bool MetadataVerifier::verifyIntegerEntry(msgpack::MapNode &MapNode,
+bool MetadataVerifier::verifyIntegerEntry(msgpack::MapDocNode &MapNode,
                                           StringRef Key, bool Required) {
-  return verifyEntry(MapNode, Key, Required, [this](msgpack::Node &Node) {
+  return verifyEntry(MapNode, Key, Required, [this](msgpack::DocNode &Node) {
     return verifyInteger(Node);
   });
 }
 
-bool MetadataVerifier::verifyKernelArgs(msgpack::Node &Node) {
-  auto ArgsMapPtr = dyn_cast<msgpack::MapNode>(&Node);
-  if (!ArgsMapPtr)
+bool MetadataVerifier::verifyKernelArgs(msgpack::DocNode &Node) {
+  if (!Node.isMap())
     return false;
-  auto &ArgsMap = *ArgsMapPtr;
+  auto &ArgsMap = Node.getMap();
 
   if (!verifyScalarEntry(ArgsMap, ".name", false,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".type_name", false,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyIntegerEntry(ArgsMap, ".size", true))
     return false;
   if (!verifyIntegerEntry(ArgsMap, ".offset", true))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("by_value", true)
                                .Case("global_buffer", true)
                                .Case("dynamic_shared_pointer", true)
                                .Case("sampler", true)
                                .Case("image", true)
                                .Case("pipe", true)
                                .Case("queue", true)
                                .Case("hidden_global_offset_x", true)
                                .Case("hidden_global_offset_y", true)
                                .Case("hidden_global_offset_z", true)
                                .Case("hidden_none", true)
                                .Case("hidden_printf_buffer", true)
                                .Case("hidden_default_queue", true)
                                .Case("hidden_completion_action", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".value_type", true,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("struct", true)
                                .Case("i8", true)
                                .Case("u8", true)
                                .Case("i16", true)
                                .Case("u16", true)
                                .Case("f16", true)
                                .Case("i32", true)
                                .Case("u32", true)
                                .Case("f32", true)
                                .Case("i64", true)
                                .Case("u64", true)
                                .Case("f64", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyIntegerEntry(ArgsMap, ".pointee_align", false))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".address_space", false,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("private", true)
                                .Case("global", true)
                                .Case("constant", true)
                                .Case("local", true)
                                .Case("generic", true)
                                .Case("region", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".access", false,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("read_only", true)
                                .Case("write_only", true)
                                .Case("read_write", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".actual_access", false,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("read_only", true)
                                .Case("write_only", true)
                                .Case("read_write", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".is_const", false,
-                         msgpack::ScalarNode::SK_Boolean))
+                         msgpack::Type::Boolean))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".is_restrict", false,
-                         msgpack::ScalarNode::SK_Boolean))
+                         msgpack::Type::Boolean))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".is_volatile", false,
-                         msgpack::ScalarNode::SK_Boolean))
+                         msgpack::Type::Boolean))
     return false;
   if (!verifyScalarEntry(ArgsMap, ".is_pipe", false,
-                         msgpack::ScalarNode::SK_Boolean))
+                         msgpack::Type::Boolean))
     return false;
 
   return true;
 }
 
-bool MetadataVerifier::verifyKernel(msgpack::Node &Node) {
-  auto KernelMapPtr = dyn_cast<msgpack::MapNode>(&Node);
-  if (!KernelMapPtr)
+bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
+  if (!Node.isMap())
     return false;
-  auto &KernelMap = *KernelMapPtr;
+  auto &KernelMap = Node.getMap();
 
   if (!verifyScalarEntry(KernelMap, ".name", true,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyScalarEntry(KernelMap, ".symbol", true,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyScalarEntry(KernelMap, ".language", false,
-                         msgpack::ScalarNode::SK_String,
-                         [](msgpack::ScalarNode &SNode) {
+                         msgpack::Type::String,
+                         [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("OpenCL C", true)
                                .Case("OpenCL C++", true)
                                .Case("HCC", true)
                                .Case("HIP", true)
                                .Case("OpenMP", true)
                                .Case("Assembler", true)
                                .Default(false);
                          }))
     return false;
   if (!verifyEntry(
-          KernelMap, ".language_version", false, [this](msgpack::Node &Node) {
+          KernelMap, ".language_version", false, [this](msgpack::DocNode &Node) {
             return verifyArray(
                 Node,
-                [this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
+                [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
           }))
     return false;
-  if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::Node &Node) {
-        return verifyArray(Node, [this](msgpack::Node &Node) {
+  if (!verifyEntry(KernelMap, ".args", false, [this](msgpack::DocNode &Node) {
+        return verifyArray(Node, [this](msgpack::DocNode &Node) {
           return verifyKernelArgs(Node);
         });
       }))
     return false;
   if (!verifyEntry(KernelMap, ".reqd_workgroup_size", false,
-                   [this](msgpack::Node &Node) {
+                   [this](msgpack::DocNode &Node) {
                      return verifyArray(Node,
-                                        [this](msgpack::Node &Node) {
+                                        [this](msgpack::DocNode &Node) {
                                           return verifyInteger(Node);
                                         },
                                         3);
                    }))
     return false;
   if (!verifyEntry(KernelMap, ".workgroup_size_hint", false,
-                   [this](msgpack::Node &Node) {
+                   [this](msgpack::DocNode &Node) {
                      return verifyArray(Node,
-                                        [this](msgpack::Node &Node) {
+                                        [this](msgpack::DocNode &Node) {
                                           return verifyInteger(Node);
                                         },
                                         3);
                    }))
     return false;
   if (!verifyScalarEntry(KernelMap, ".vec_type_hint", false,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyScalarEntry(KernelMap, ".device_enqueue_symbol", false,
-                         msgpack::ScalarNode::SK_String))
+                         msgpack::Type::String))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_size", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".group_segment_fixed_size", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".private_segment_fixed_size", true))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".kernarg_segment_align", true))
bool MetadataVerifier::verifyKernel(msgpack::DocNode &Node) {
   if (!verifyIntegerEntry(KernelMap, ".sgpr_spill_count", false))
     return false;
   if (!verifyIntegerEntry(KernelMap, ".vgpr_spill_count", false))
     return false;
 
   return true;
 }
 
-bool MetadataVerifier::verify(msgpack::Node &HSAMetadataRoot) {
-  auto RootMapPtr = dyn_cast<msgpack::MapNode>(&HSAMetadataRoot);
-  if (!RootMapPtr)
+bool MetadataVerifier::verify(msgpack::DocNode &HSAMetadataRoot) {
+  if (!HSAMetadataRoot.isMap())
     return false;
-  auto &RootMap = *RootMapPtr;
+  auto &RootMap = HSAMetadataRoot.getMap();
 
   if (!verifyEntry(
-          RootMap, "amdhsa.version", true, [this](msgpack::Node &Node) {
+          RootMap, "amdhsa.version", true, [this](msgpack::DocNode &Node) {
             return verifyArray(
                 Node,
-                [this](msgpack::Node &Node) { return verifyInteger(Node); }, 2);
+                [this](msgpack::DocNode &Node) { return verifyInteger(Node); }, 2);
           }))
     return false;
   if (!verifyEntry(
-          RootMap, "amdhsa.printf", false, [this](msgpack::Node &Node) {
-            return verifyArray(Node, [this](msgpack::Node &Node) {
-              return verifyScalar(Node, msgpack::ScalarNode::SK_String);
+          RootMap, "amdhsa.printf", false, [this](msgpack::DocNode &Node) {
+            return verifyArray(Node, [this](msgpack::DocNode &Node) {
+              return verifyScalar(Node, msgpack::Type::String);
             });
           }))
     return false;
   if (!verifyEntry(RootMap, "amdhsa.kernels", true,
-                   [this](msgpack::Node &Node) {
-                     return verifyArray(Node, [this](msgpack::Node &Node) {
+                   [this](msgpack::DocNode &Node) {
+                     return verifyArray(Node, [this](msgpack::DocNode &Node) {
                        return verifyKernel(Node);
                      });
                    }))
     return false;
 
   return true;
 }
 
 } // end namespace V3
 } // end namespace HSAMD
 } // end namespace AMDGPU
 } // end namespace llvm

lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h

 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUHSAMETADATASTREAMER_H
 #define LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUHSAMETADATASTREAMER_H
 
 #include "AMDGPU.h"
 #include "AMDKernelCodeT.h"
 #include "llvm/ADT/StringRef.h"
-#include "llvm/BinaryFormat/MsgPackTypes.h"
+#include "llvm/BinaryFormat/MsgPackDocument.h"
 #include "llvm/Support/AMDGPUMetadata.h"
 
 namespace llvm {
 
 class AMDGPUTargetStreamer;
 class Argument;
 class DataLayout;
 class Function;
public:
   virtual void end() = 0;
 
   virtual void emitKernel(const MachineFunction &MF,
                           const SIProgramInfo &ProgramInfo) = 0;
 };
 
 class MetadataStreamerV3 final : public MetadataStreamer {
 private:
-  std::shared_ptr<msgpack::Node> HSAMetadataRoot =
-      std::make_shared<msgpack::MapNode>();
+  std::unique_ptr<msgpack::Document> HSAMetadataDoc =
+      llvm::make_unique<msgpack::Document>();
 
   void dump(StringRef HSAMetadataString) const;
 
   void verify(StringRef HSAMetadataString) const;
 
   Optional<StringRef> getAccessQualifier(StringRef AccQual) const;
 
   Optional<StringRef> getAddressSpaceQualifier(unsigned AddressSpace) const;
 
   StringRef getValueKind(Type *Ty, StringRef TypeQual,
                          StringRef BaseTypeName) const;
 
   StringRef getValueType(Type *Ty, StringRef TypeName) const;
 
   std::string getTypeName(Type *Ty, bool Signed) const;
 
-  std::shared_ptr<msgpack::ArrayNode>
-  getWorkGroupDimensions(MDNode *Node) const;
+  msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const;
 
-  std::shared_ptr<msgpack::MapNode>
-  getHSAKernelProps(const MachineFunction &MF,
+  msgpack::MapDocNode getHSAKernelProps(const MachineFunction &MF,
                                         const SIProgramInfo &ProgramInfo) const;
 
   void emitVersion();
 
   void emitPrintf(const Module &Mod);
 
-  void emitKernelLanguage(const Function &Func, msgpack::MapNode &Kern);
+  void emitKernelLanguage(const Function &Func, msgpack::MapDocNode Kern);
 
-  void emitKernelAttrs(const Function &Func, msgpack::MapNode &Kern);
+  void emitKernelAttrs(const Function &Func, msgpack::MapDocNode Kern);
 
-  void emitKernelArgs(const Function &Func, msgpack::MapNode &Kern);
+  void emitKernelArgs(const Function &Func, msgpack::MapDocNode Kern);
 
   void emitKernelArg(const Argument &Arg, unsigned &Offset,
-                     msgpack::ArrayNode &Args);
+                     msgpack::ArrayDocNode Args);
 
   void emitKernelArg(const DataLayout &DL, Type *Ty, StringRef ValueKind,
-                     unsigned &Offset, msgpack::ArrayNode &Args,
+                     unsigned &Offset, msgpack::ArrayDocNode Args,
                      unsigned PointeeAlign = 0, StringRef Name = "",
                      StringRef TypeName = "", StringRef BaseTypeName = "",
                      StringRef AccQual = "", StringRef TypeQual = "");
 
   void emitHiddenKernelArgs(const Function &Func, unsigned &Offset,
-                            msgpack::ArrayNode &Args);
+                            msgpack::ArrayDocNode Args);
 
-  std::shared_ptr<msgpack::Node> &getRootMetadata(StringRef Key) {
-    return (*cast<msgpack::MapNode>(HSAMetadataRoot.get()))[Key];
+  msgpack::DocNode &getRootMetadata(StringRef Key) {
+    return HSAMetadataDoc->getRoot().getMap(/*Convert=*/true)[Key];
   }
 
-  std::shared_ptr<msgpack::Node> &getHSAMetadataRoot() {
-    return HSAMetadataRoot;
+  msgpack::DocNode &getHSAMetadataRoot() {
+    return HSAMetadataDoc->getRoot();
   }
 
 public:
   MetadataStreamerV3() = default;
   ~MetadataStreamerV3() = default;
 
   bool emitTo(AMDGPUTargetStreamer &TargetStreamer) override;
 

lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

 
 void MetadataStreamerV3::dump(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata:\n" << HSAMetadataString << '\n';
 }
 
 void MetadataStreamerV3::verify(StringRef HSAMetadataString) const {
   errs() << "AMDGPU HSA Metadata Parser Test: ";
 
-  std::shared_ptr<msgpack::Node> FromHSAMetadataString =
-      std::make_shared<msgpack::MapNode>();
+  msgpack::Document FromHSAMetadataString;
 
-  yaml::Input YIn(HSAMetadataString);
-  YIn >> FromHSAMetadataString;
-  if (YIn.error()) {
+  if (!FromHSAMetadataString.fromYAML(HSAMetadataString)) {
     errs() << "FAIL\n";
     return;
   }
 
   std::string ToHSAMetadataString;
   raw_string_ostream StrOS(ToHSAMetadataString);
-  yaml::Output YOut(StrOS);
-  YOut << FromHSAMetadataString;
+  FromHSAMetadataString.toYAML(StrOS);
 
   errs() << (HSAMetadataString == StrOS.str() ? "PASS" : "FAIL") << '\n';
   if (HSAMetadataString != ToHSAMetadataString) {
     errs() << "Original input: " << HSAMetadataString << '\n'
            << "Produced output: " << StrOS.str() << '\n';
   }
 }
 
case Type::VectorTyID: {
     auto NumElements = VecTy->getVectorNumElements();
     return (Twine(getTypeName(ElTy, Signed)) + Twine(NumElements)).str();
   }
   default:
     return "unknown";
   }
 }
 
-std::shared_ptr<msgpack::ArrayNode>
+msgpack::ArrayDocNode
 MetadataStreamerV3::getWorkGroupDimensions(MDNode *Node) const {
-  auto Dims = std::make_shared<msgpack::ArrayNode>();
+  auto Dims = HSAMetadataDoc->getArrayNode();
   if (Node->getNumOperands() != 3)
     return Dims;
 
   for (auto &Op : Node->operands())
-    Dims->push_back(std::make_shared<msgpack::ScalarNode>(
-        mdconst::extract<ConstantInt>(Op)->getZExtValue()));
+    Dims.push_back(Dims.getDocument()->getNode(
+        uint64_t(mdconst::extract<ConstantInt>(Op)->getZExtValue())));
   return Dims;
 }
 
 void MetadataStreamerV3::emitVersion() {
-  auto Version = std::make_shared<msgpack::ArrayNode>();
-  Version->push_back(std::make_shared<msgpack::ScalarNode>(V3::VersionMajor));
-  Version->push_back(std::make_shared<msgpack::ScalarNode>(V3::VersionMinor));
-  getRootMetadata("amdhsa.version") = std::move(Version);
+  auto Version = HSAMetadataDoc->getArrayNode();
+  Version.push_back(Version.getDocument()->getNode(VersionMajor));
+  Version.push_back(Version.getDocument()->getNode(VersionMinor));
+  getRootMetadata("amdhsa.version") = Version;
 }
 
 void MetadataStreamerV3::emitPrintf(const Module &Mod) {
   auto Node = Mod.getNamedMetadata("llvm.printf.fmts");
   if (!Node)
     return;
 
-  auto Printf = std::make_shared<msgpack::ArrayNode>();
+  auto Printf = HSAMetadataDoc->getArrayNode();
   for (auto Op : Node->operands())
     if (Op->getNumOperands())
-      Printf->push_back(std::make_shared<msgpack::ScalarNode>(
-          cast<MDString>(Op->getOperand(0))->getString()));
-  getRootMetadata("amdhsa.printf") = std::move(Printf);
+      Printf.push_back(Printf.getDocument()->getNode(
+          cast<MDString>(Op->getOperand(0))->getString(), /*Copy=*/true));
+  getRootMetadata("amdhsa.printf") = Printf;
 }
 
 void MetadataStreamerV3::emitKernelLanguage(const Function &Func,
-                                            msgpack::MapNode &Kern) {
+                                            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"] = std::make_shared<msgpack::ScalarNode>("OpenCL C");
-  auto LanguageVersion = std::make_shared<msgpack::ArrayNode>();
-  LanguageVersion->push_back(std::make_shared<msgpack::ScalarNode>(
+  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(std::make_shared<msgpack::ScalarNode>(
+  LanguageVersion.push_back(Kern.getDocument()->getNode(
       mdconst::extract<ConstantInt>(Op0->getOperand(1))->getZExtValue()));
-  Kern[".language_version"] = std::move(LanguageVersion);
+  Kern[".language_version"] = LanguageVersion;
 }
 
 void MetadataStreamerV3::emitKernelAttrs(const Function &Func,
-                                         msgpack::MapNode &Kern) {
+                                         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"] = std::make_shared<msgpack::ScalarNode>(getTypeName(
+    Kern[".vec_type_hint"] = Kern.getDocument()->getNode(
+        getTypeName(
             cast<ValueAsMetadata>(Node->getOperand(0))->getType(),
-        mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()));
+            mdconst::extract<ConstantInt>(Node->getOperand(1))->getZExtValue()),
+        /*Copy=*/true);
   }
   if (Func.hasFnAttribute("runtime-handle")) {
-    Kern[".device_enqueue_symbol"] = std::make_shared<msgpack::ScalarNode>(
-        Func.getFnAttribute("runtime-handle").getValueAsString().str());
+    Kern[".device_enqueue_symbol"] = Kern.getDocument()->getNode(
+        Func.getFnAttribute("runtime-handle").getValueAsString().str(),
+        /*Copy=*/true);
   }
 }
 
 void MetadataStreamerV3::emitKernelArgs(const Function &Func,
-                                        msgpack::MapNode &Kern) {
+                                        msgpack::MapDocNode Kern) {
   unsigned Offset = 0;
-  auto Args = std::make_shared<msgpack::ArrayNode>();
+  auto Args = HSAMetadataDoc->getArrayNode();
   for (auto &Arg : Func.args())
-    emitKernelArg(Arg, Offset, *Args);
+    emitKernelArg(Arg, Offset, Args);
 
-  emitHiddenKernelArgs(Func, Offset, *Args);
+  emitHiddenKernelArgs(Func, Offset, Args);
 
   // TODO: What about other languages?
   if (Func.getParent()->getNamedMetadata("opencl.ocl.version")) {
     auto &DL = Func.getParent()->getDataLayout();
     auto Int64Ty = Type::getInt64Ty(Func.getContext());
 
-    emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, *Args);
-    emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, *Args);
-    emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, *Args);
+    emitKernelArg(DL, Int64Ty, "hidden_global_offset_x", Offset, Args);
+    emitKernelArg(DL, Int64Ty, "hidden_global_offset_y", Offset, Args);
+    emitKernelArg(DL, Int64Ty, "hidden_global_offset_z", Offset, Args);
 
     auto Int8PtrTy =
         Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 
     // Emit "printf buffer" argument if printf is used, otherwise emit dummy
     // "none" argument.
     if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
-      emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, *Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_printf_buffer", Offset, Args);
     else
-      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
 
     // Emit "default queue" and "completion action" arguments if enqueue kernel
     // is used, otherwise emit dummy "none" arguments.
     if (Func.hasFnAttribute("calls-enqueue-kernel")) {
-      emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, *Args);
-      emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, *Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
     } else {
-      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
-      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, *Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
+      emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
     }
   }
 
-  Kern[".args"] = std::move(Args);
+  Kern[".args"] = Args;
 }
 
 void MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
-                                       msgpack::ArrayNode &Args) {
+                                       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 MetadataStreamerV3::emitKernelArg(const Argument &Arg, unsigned &Offset,
   emitKernelArg(Func->getParent()->getDataLayout(), Arg.getType(),
                 getValueKind(Arg.getType(), TypeQual, BaseTypeName), Offset,
                 Args, PointeeAlign, Name, TypeName, BaseTypeName, AccQual,
                 TypeQual);
 }
 
 void MetadataStreamerV3::emitKernelArg(const DataLayout &DL, Type *Ty,
                                        StringRef ValueKind, unsigned &Offset,
-                                       msgpack::ArrayNode &Args,
+                                       msgpack::ArrayDocNode Args,
                                        unsigned PointeeAlign, StringRef Name,
                                        StringRef TypeName,
                                        StringRef BaseTypeName,
                                        StringRef AccQual, StringRef TypeQual) {
-  auto ArgPtr = std::make_shared<msgpack::MapNode>();
-  auto &Arg = *ArgPtr;
+  auto Arg = Args.getDocument()->getMapNode();
 
   if (!Name.empty())
-    Arg[".name"] = std::make_shared<msgpack::ScalarNode>(Name);
+    Arg[".name"] = Arg.getDocument()->getNode(Name, /*Copy=*/true);
   if (!TypeName.empty())
-    Arg[".type_name"] = std::make_shared<msgpack::ScalarNode>(TypeName);
+    Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
   auto Size = DL.getTypeAllocSize(Ty);
   auto Align = DL.getABITypeAlignment(Ty);
-  Arg[".size"] = std::make_shared<msgpack::ScalarNode>(Size);
+  Arg[".size"] = Arg.getDocument()->getNode(Size);
   Offset = alignTo(Offset, Align);
-  Arg[".offset"] = std::make_shared<msgpack::ScalarNode>(Offset);
+  Arg[".offset"] = Arg.getDocument()->getNode(Offset);
   Offset += Size;
-  Arg[".value_kind"] = std::make_shared<msgpack::ScalarNode>(ValueKind);
+  Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
   Arg[".value_type"] =
-      std::make_shared<msgpack::ScalarNode>(getValueType(Ty, BaseTypeName));
+      Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
   if (PointeeAlign)
-    Arg[".pointee_align"] = std::make_shared<msgpack::ScalarNode>(PointeeAlign);
+    Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign);
 
   if (auto PtrTy = dyn_cast<PointerType>(Ty))
     if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
-      Arg[".address_space"] = std::make_shared<msgpack::ScalarNode>(*Qualifier);
+      Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
 
   if (auto AQ = getAccessQualifier(AccQual))
-    Arg[".access"] = std::make_shared<msgpack::ScalarNode>(*AQ);
+    Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
 
   // TODO: Emit Arg[".actual_access"].
 
   SmallVector<StringRef, 1> SplitTypeQuals;
   TypeQual.split(SplitTypeQuals, " ", -1, false);
   for (StringRef Key : SplitTypeQuals) {
     if (Key == "const")
-      Arg[".is_const"] = std::make_shared<msgpack::ScalarNode>(true);
+      Arg[".is_const"] = Arg.getDocument()->getNode(true);
     else if (Key == "restrict")
-      Arg[".is_restrict"] = std::make_shared<msgpack::ScalarNode>(true);
+      Arg[".is_restrict"] = Arg.getDocument()->getNode(true);
     else if (Key == "volatile")
-      Arg[".is_volatile"] = std::make_shared<msgpack::ScalarNode>(true);
+      Arg[".is_volatile"] = Arg.getDocument()->getNode(true);
     else if (Key == "pipe")
-      Arg[".is_pipe"] = std::make_shared<msgpack::ScalarNode>(true);
+      Arg[".is_pipe"] = Arg.getDocument()->getNode(true);
   }
 
-  Args.push_back(std::move(ArgPtr));
+  Args.push_back(Arg);
 }
 
 void MetadataStreamerV3::emitHiddenKernelArgs(const Function &Func,
                                               unsigned &Offset,
-                                              msgpack::ArrayNode &Args) {
+                                              msgpack::ArrayDocNode Args) {
   int HiddenArgNumBytes =
       getIntegerAttribute(Func, "amdgpu-implicitarg-num-bytes", 0);
 
   if (!HiddenArgNumBytes)
     return;
 
   auto &DL = Func.getParent()->getDataLayout();
   auto Int64Ty = Type::getInt64Ty(Func.getContext());
if (Func.hasFnAttribute("calls-enqueue-kernel")) {
       emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
     } else {
       emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
       emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
     }
   }
 }
 
-std::shared_ptr<msgpack::MapNode>
+msgpack::MapDocNode
 MetadataStreamerV3::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 HSAKernelProps = std::make_shared<msgpack::MapNode>();
-  auto &Kern = *HSAKernelProps;
+  auto Kern = HSAMetadataDoc->getMapNode();
 
   unsigned MaxKernArgAlign;
-  Kern[".kernarg_segment_size"] = std::make_shared<msgpack::ScalarNode>(
+  Kern[".kernarg_segment_size"] = Kern.getDocument()->getNode(
       STM.getKernArgSegmentSize(F, MaxKernArgAlign));
   Kern[".group_segment_fixed_size"] =
-      std::make_shared<msgpack::ScalarNode>(ProgramInfo.LDSSize);
+      Kern.getDocument()->getNode(ProgramInfo.LDSSize);
   Kern[".private_segment_fixed_size"] =
-      std::make_shared<msgpack::ScalarNode>(ProgramInfo.ScratchSize);
+      Kern.getDocument()->getNode(ProgramInfo.ScratchSize);
   Kern[".kernarg_segment_align"] =
-      std::make_shared<msgpack::ScalarNode>(std::max(uint32_t(4), MaxKernArgAlign));
+      Kern.getDocument()->getNode(std::max(uint32_t(4), MaxKernArgAlign));
   Kern[".wavefront_size"] =
-      std::make_shared<msgpack::ScalarNode>(STM.getWavefrontSize());
-  Kern[".sgpr_count"] = std::make_shared<msgpack::ScalarNode>(ProgramInfo.NumSGPR);
-  Kern[".vgpr_count"] = std::make_shared<msgpack::ScalarNode>(ProgramInfo.NumVGPR);
+      Kern.getDocument()->getNode(STM.getWavefrontSize());
+  Kern[".sgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumSGPR);
+  Kern[".vgpr_count"] = Kern.getDocument()->getNode(ProgramInfo.NumVGPR);
   Kern[".max_flat_workgroup_size"] =
-      std::make_shared<msgpack::ScalarNode>(MFI.getMaxFlatWorkGroupSize());
+      Kern.getDocument()->getNode(MFI.getMaxFlatWorkGroupSize());
   Kern[".sgpr_spill_count"] =
-      std::make_shared<msgpack::ScalarNode>(MFI.getNumSpilledSGPRs());
+      Kern.getDocument()->getNode(MFI.getNumSpilledSGPRs());
   Kern[".vgpr_spill_count"] =
-      std::make_shared<msgpack::ScalarNode>(MFI.getNumSpilledVGPRs());
+      Kern.getDocument()->getNode(MFI.getNumSpilledVGPRs());
 
-  return HSAKernelProps;
+  return Kern;
 }
 
 bool MetadataStreamerV3::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
-  return TargetStreamer.EmitHSAMetadata(getHSAMetadataRoot(), true);
+  return TargetStreamer.EmitHSAMetadata(*HSAMetadataDoc, true);
 }
 
 void MetadataStreamerV3::begin(const Module &Mod) {
   emitVersion();
   emitPrintf(Mod);
-  getRootMetadata("amdhsa.kernels").reset(new msgpack::ArrayNode());
+  getRootMetadata("amdhsa.kernels") = HSAMetadataDoc->getArrayNode();
 }
 
 void MetadataStreamerV3::end() {
   std::string HSAMetadataString;
   raw_string_ostream StrOS(HSAMetadataString);
-  yaml::Output YOut(StrOS);
-  YOut << HSAMetadataRoot;
+  HSAMetadataDoc->toYAML(StrOS);
 
   if (DumpHSAMetadata)
     dump(StrOS.str());
   if (VerifyHSAMetadata)
     verify(StrOS.str());
 }
 
 void MetadataStreamerV3::emitKernel(const MachineFunction &MF,
                                     const SIProgramInfo &ProgramInfo) {
   auto &Func = MF.getFunction();
-  auto KernelProps = getHSAKernelProps(MF, ProgramInfo);
+  auto Kern = getHSAKernelProps(MF, ProgramInfo);
 
   assert(Func.getCallingConv() == CallingConv::AMDGPU_KERNEL ||
          Func.getCallingConv() == CallingConv::SPIR_KERNEL);
 
-  auto &KernelsNode = getRootMetadata("amdhsa.kernels");
-  auto Kernels = cast<msgpack::ArrayNode>(KernelsNode.get());
+  auto Kernels =
+      getRootMetadata("amdhsa.kernels").getArray(/*Convert=*/true);
 
   {
-    auto &Kern = *KernelProps;
-    Kern[".name"] = std::make_shared<msgpack::ScalarNode>(Func.getName());
-    Kern[".symbol"] = std::make_shared<msgpack::ScalarNode>(
-        (Twine(Func.getName()) + Twine(".kd")).str());
+    Kern[".name"] = Kern.getDocument()->getNode(Func.getName());
+    Kern[".symbol"] = Kern.getDocument()->getNode(
+        (Twine(Func.getName()) + Twine(".kd")).str(), /*Copy=*/true);
     emitKernelLanguage(Func, Kern);
     emitKernelAttrs(Func, Kern);
     emitKernelArgs(Func, Kern);
   }
 
-  Kernels->push_back(std::move(KernelProps));
+  Kernels.push_back(Kern);
 }
 
 } // end namespace HSAMD
 } // end namespace AMDGPU
 } // end namespace llvm

lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.h

 //===-- AMDGPUTargetStreamer.h - AMDGPU Target Streamer --------*- C++ -*--===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUTARGETSTREAMER_H
 #define LLVM_LIB_TARGET_AMDGPU_MCTARGETDESC_AMDGPUTARGETSTREAMER_H
 
 #include "AMDKernelCodeT.h"
-#include "llvm/BinaryFormat/MsgPackTypes.h"
+#include "llvm/BinaryFormat/MsgPackDocument.h"
 #include "llvm/MC/MCStreamer.h"
 #include "llvm/MC/MCSubtargetInfo.h"
 #include "llvm/Support/AMDGPUMetadata.h"
 #include "llvm/Support/AMDHSAKernelDescriptor.h"
 
 namespace llvm {
 #include "AMDGPUPTNote.h"
 
public:
 
   /// Emit HSA Metadata
   ///
   /// When \p Strict is true, known metadata elements must already be
   /// well-typed. When \p Strict is false, known types are inferred and
   /// the \p HSAMetadata structure is updated with the correct types.
   ///
   /// \returns True on success, false on failure.
-  virtual bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
-                               bool Strict) = 0;
+  virtual bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) = 0;
 
   /// \returns True on success, false on failure.
   virtual bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) = 0;
 
   /// \returns True on success, false on failure.
   virtual bool EmitPALMetadata(const AMDGPU::PALMD::Metadata &PALMetadata) = 0;
 
   virtual void EmitAmdhsaKernelDescriptor(
public:
   void EmitAMDKernelCodeT(const amd_kernel_code_t &Header) override;
 
   void EmitAMDGPUSymbolType(StringRef SymbolName, unsigned Type) override;
 
   /// \returns True on success, false on failure.
   bool EmitISAVersion(StringRef IsaVersionString) override;
 
   /// \returns True on success, false on failure.
-  bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
-                       bool Strict) override;
+  bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) override;
 
   /// \returns True on success, false on failure.
   bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) override;
 
   /// \returns True on success, false on failure.
   bool EmitPALMetadata(const AMDGPU::PALMD::Metadata &PALMetadata) override;
 
   void EmitAmdhsaKernelDescriptor(
public:
   void EmitAMDKernelCodeT(const amd_kernel_code_t &Header) override;
 
   void EmitAMDGPUSymbolType(StringRef SymbolName, unsigned Type) override;
 
   /// \returns True on success, false on failure.
   bool EmitISAVersion(StringRef IsaVersionString) override;
 
   /// \returns True on success, false on failure.
-  bool EmitHSAMetadata(std::shared_ptr<msgpack::Node> &HSAMetadata,
-                       bool Strict) override;
+  bool EmitHSAMetadata(msgpack::Document &HSAMetadata, bool Strict) override;
 
   /// \returns True on success, false on failure.
   bool EmitHSAMetadata(const AMDGPU::HSAMD::Metadata &HSAMetadata) override;
 
   /// \returns True on success, false on failure.
   bool EmitPALMetadata(const AMDGPU::PALMD::Metadata &PALMetadata) override;
 
   void EmitAmdhsaKernelDescriptor(
       const MCSubtargetInfo &STI, StringRef KernelName,
       const amdhsa::kernel_descriptor_t &KernelDescriptor, uint64_t NextVGPR,
       uint64_t NextSGPR, bool ReserveVCC, bool ReserveFlatScr,
       bool ReserveXNACK) override;
 };
 
 }
 #endif

lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp

 #include "AMDGPUTargetStreamer.h"
 #include "AMDGPU.h"
 #include "SIDefines.h"
 #include "Utils/AMDGPUBaseInfo.h"
 #include "Utils/AMDKernelCodeTUtils.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/BinaryFormat/AMDGPUMetadataVerifier.h"
 #include "llvm/BinaryFormat/ELF.h"
-#include "llvm/BinaryFormat/MsgPackTypes.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/Metadata.h"
 #include "llvm/IR/Module.h"
 #include "llvm/MC/MCContext.h"
 #include "llvm/MC/MCELFStreamer.h"
 #include "llvm/MC/MCObjectFileInfo.h"
 #include "llvm/MC/MCSectionELF.h"
bool AMDGPUTargetStreamer::EmitHSAMetadataV2(StringRef HSAMetadataString) {
   HSAMD::Metadata HSAMetadata;
   if (HSAMD::fromString(HSAMetadataString, HSAMetadata))
     return false;
 
   return EmitHSAMetadata(HSAMetadata);
 }
 
 bool AMDGPUTargetStreamer::EmitHSAMetadataV3(StringRef HSAMetadataString) {
-  std::shared_ptr<msgpack::Node> HSAMetadataRoot;
-  yaml::Input YIn(HSAMetadataString);
-  YIn >> HSAMetadataRoot;
-  if (YIn.error())
+  msgpack::Document HSAMetadataDoc;
+  if (!HSAMetadataDoc.fromYAML(HSAMetadataString))
     return false;
-  return EmitHSAMetadata(HSAMetadataRoot, false);
+  return EmitHSAMetadata(HSAMetadataDoc, false);
 }
 
 StringRef AMDGPUTargetStreamer::getArchNameFromElfMach(unsigned ElfMach) {
   AMDGPU::GPUKind AK;
 
   switch (ElfMach) {
   case ELF::EF_AMDGPU_MACH_R600_R600:     AK = GK_R600;    break;
   case ELF::EF_AMDGPU_MACH_R600_R630:     AK = GK_R630;    break;
bool AMDGPUTargetAsmStreamer::EmitHSAMetadata(
 
   OS << '\t' << AssemblerDirectiveBegin << '\n';
   OS << HSAMetadataString << '\n';
   OS << '\t' << AssemblerDirectiveEnd << '\n';
   return true;
 }
 
 bool AMDGPUTargetAsmStreamer::EmitHSAMetadata(
-    std::shared_ptr<msgpack::Node> &HSAMetadataRoot, bool Strict) {
+    msgpack::Document &HSAMetadataDoc, bool Strict) {
   V3::MetadataVerifier Verifier(Strict);
-  if (!Verifier.verify(*HSAMetadataRoot))
+  if (!Verifier.verify(HSAMetadataDoc.getRoot()))
     return false;
 
   std::string HSAMetadataString;
   raw_string_ostream StrOS(HSAMetadataString);
-  yaml::Output YOut(StrOS);
-  YOut << HSAMetadataRoot;
+  HSAMetadataDoc.toYAML(StrOS);
 
   OS << '\t' << V3::AssemblerDirectiveBegin << '\n';
   OS << StrOS.str() << '\n';
   OS << '\t' << V3::AssemblerDirectiveEnd << '\n';
   return true;
 }
 
 bool AMDGPUTargetAsmStreamer::EmitPALMetadata(
EmitNote(ElfNote::NoteNameV2, DescSZ, ELF::NT_AMD_AMDGPU_ISA,
            [&](MCELFStreamer &OS) {
              OS.EmitLabel(DescBegin);
              OS.EmitBytes(IsaVersionString);
              OS.EmitLabel(DescEnd);
            });
   return true;
 }
 
-bool AMDGPUTargetELFStreamer::EmitHSAMetadata(
-    std::shared_ptr<msgpack::Node> &HSAMetadataRoot, bool Strict) {
+bool AMDGPUTargetELFStreamer::EmitHSAMetadata(msgpack::Document &HSAMetadataDoc,
+                                              bool Strict) {
   V3::MetadataVerifier Verifier(Strict);
-  if (!Verifier.verify(*HSAMetadataRoot))
+  if (!Verifier.verify(HSAMetadataDoc.getRoot()))
     return false;
 
   std::string HSAMetadataString;
-  raw_string_ostream StrOS(HSAMetadataString);
-  msgpack::Writer MPWriter(StrOS);
-  HSAMetadataRoot->write(MPWriter);
+  HSAMetadataDoc.writeToBlob(HSAMetadataString);
 
   // Create two labels to mark the beginning and end of the desc field
   // and a MCExpr to calculate the size of the desc field.
   auto &Context = getContext();
   auto *DescBegin = Context.createTempSymbol();
   auto *DescEnd = Context.createTempSymbol();
   auto *DescSZ = MCBinaryExpr::createSub(
       MCSymbolRefExpr::create(DescEnd, Context),
       MCSymbolRefExpr::create(DescBegin, Context), Context);
 
   EmitNote(ElfNote::NoteNameV3, DescSZ, ELF::NT_AMDGPU_METADATA,
            [&](MCELFStreamer &OS) {
              OS.EmitLabel(DescBegin);
-             OS.EmitBytes(StrOS.str());
+             OS.EmitBytes(HSAMetadataString);
              OS.EmitLabel(DescEnd);
            });
   return true;
 }
 
 bool AMDGPUTargetELFStreamer::EmitHSAMetadata(
     const AMDGPU::HSAMD::Metadata &HSAMetadata) {
   std::string HSAMetadataString;

test/CodeGen/AMDGPU/hsa-metadata-deduce-ro-arg-v3.ll

 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck %s
 
-; CHECK:        .symbol:          test_ro_arg.kd
-; CHECK:        .name:            test_ro_arg
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:       'float*'
-; CHECK-NEXT:       .value_kind:      global_buffer
-; CHECK-NEXT:       .name:            in
-; CHECK-NEXT:       .access:          read_only
-; CHECK-NEXT:       .offset:          0
-; CHECK-NEXT:       .is_const:        true
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .is_restrict:     true
-; CHECK-NEXT:       .value_type:      f32
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .type_name:       'float*'
-; CHECK-NEXT:       .value_kind:      global_buffer
-; CHECK-NEXT:       .name:            out
-; CHECK-NEXT:       .offset:          8
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      f32
-; CHECK-NEXT:       .address_space:   global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .access:         read_only
+; CHECK-NEXT:         .address_space:  global
+; CHECK-NEXT:         .is_const:       true
+; CHECK-NEXT:         .is_restrict:    true
+; CHECK-NEXT:         .name:           in
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'float*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     f32
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           out
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'float*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     f32
+; CHECK:          .name:           test_ro_arg
+; CHECK:          .symbol:         test_ro_arg.kd
 
 define amdgpu_kernel void @test_ro_arg(float addrspace(1)* noalias readonly %in, float addrspace(1)* %out)
     !kernel_arg_addr_space !0 !kernel_arg_access_qual !1 !kernel_arg_type !2
     !kernel_arg_base_type !2 !kernel_arg_type_qual !3 {
   ret void
 }
 
 !0 = !{i32 1, i32 1}
 !1 = !{!"none", !"none"}
 !2 = !{!"float*", !"float*"}
 !3 = !{!"const restrict", !""}

test/CodeGen/AMDGPU/hsa-metadata-enqueu-kernel-v3.ll

 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
 
-; CHECK: ---
-; CHECK:  amdhsa.kernels:
-; CHECK:        .symbol:          test_non_enqueue_kernel_caller.kd
-; CHECK:        .name:            test_non_enqueue_kernel_caller
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:     char
-; CHECK-NEXT:       .value_kind:    by_value
-; CHECK-NEXT:       .offset:        0
-; CHECK-NEXT:       .size:          1
-; CHECK-NEXT:       .value_type:    i8
-; CHECK-NEXT:       .name:          a
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_x
-; CHECK-NEXT:       .offset:        8
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_y
-; CHECK-NEXT:       .offset:        16
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_z
-; CHECK-NEXT:       .offset:        24
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
+; CHECK:              ---
+; CHECK:      amdhsa.kernels:  
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           1
+; CHECK-NEXT:         .type_name:      char
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
 ; CHECK-NOT:        .value_kind:    hidden_default_queue
 ; CHECK-NOT:        .value_kind:    hidden_completion_action
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_non_enqueue_kernel_caller
+; CHECK:          .symbol:         test_non_enqueue_kernel_caller.kd
 define amdgpu_kernel void @test_non_enqueue_kernel_caller(i8 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_enqueue_kernel_caller.kd
-; CHECK:        .name:            test_enqueue_kernel_caller
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:     char
-; CHECK-NEXT:       .value_kind:    by_value
-; CHECK-NEXT:       .offset:        0
-; CHECK-NEXT:       .size:          1
-; CHECK-NEXT:       .value_type:    i8
-; CHECK-NEXT:       .name:          a
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_x
-; CHECK-NEXT:       .offset:        8
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_y
-; CHECK-NEXT:       .offset:        16
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
-; CHECK-NEXT:     - .value_kind:    hidden_global_offset_z
-; CHECK-NEXT:       .offset:        24
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i64
-; CHECK-NEXT:     - .value_kind:    hidden_none
-; CHECK-NEXT:       .offset:        32
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i8
-; CHECK-NEXT:       .address_space: global
-; CHECK-NEXT:     - .value_kind:    hidden_default_queue
-; CHECK-NEXT:       .offset:        40
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i8
-; CHECK-NEXT:       .address_space: global
-; CHECK-NEXT:     - .value_kind:    hidden_completion_action
-; CHECK-NEXT:       .offset:        48
-; CHECK-NEXT:       .size:          8
-; CHECK-NEXT:       .value_type:    i8
-; CHECK-NEXT:       .address_space: global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           1
+; CHECK-NEXT:         .type_name:      char
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_default_queue
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_completion_action
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_enqueue_kernel_caller
+; CHECK:          .symbol:         test_enqueue_kernel_caller.kd
 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #0
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
   ret void
 }
 
 ; CHECK:  amdhsa.version:
 ; CHECK-NEXT: - 1

test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll

 %opencl.image3d_t = type opaque
 %opencl.queue_t = type opaque
 %opencl.pipe_t = type opaque
 %struct.B = type { i32 addrspace(1)*}
 %opencl.clk_event_t = type opaque
 
 @__test_block_invoke_kernel_runtime_handle = external addrspace(1) externally_initialized constant i8 addrspace(1)*
 
-; CHECK: ---
-; CHECK:  amdhsa.kernels:
-; CHECK:        .symbol:          test_char.kd
-; CHECK:        .name:            test_char
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      char
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           1
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NOT:        .value_kind:     hidden_default_queue
-; CHECK-NOT:        .value_kind:     hidden_completion_action
+; CHECK:              ---
+; CHECK-NEXT: amdhsa.kernels:  
+; CHECK-NEXT:   - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           1
+; CHECK-NEXT:         .type_name:      char
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NOT:          .value_kind:     hidden_default_queue
+; CHECK-NOT:          .value_kind:     hidden_completion_action
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_char
+; CHECK:          .symbol:         test_char.kd
 define amdgpu_kernel void @test_char(i8 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
     !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_ushort2.kd
-; CHECK:        .name:            test_ushort2
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      ushort2
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     u16
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      ushort2
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     u16
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_ushort2
+; CHECK:          .symbol:         test_ushort2.kd
 define amdgpu_kernel void @test_ushort2(<2 x i16> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
     !kernel_arg_base_type !10 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_int3.kd
-; CHECK:        .name:            test_int3
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int3
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           16
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           16
+; CHECK-NEXT:         .type_name:      int3
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_int3
+; CHECK:          .symbol:         test_int3.kd
 define amdgpu_kernel void @test_int3(<3 x i32> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
     !kernel_arg_base_type !11 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_ulong4.kd
-; CHECK:        .name:            test_ulong4
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      ulong4
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           32
-; CHECK-NEXT:       .value_type:     u64
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         56
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           32
+; CHECK-NEXT:         .type_name:      ulong4
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     u64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         56
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_ulong4
+; CHECK:          .symbol:         test_ulong4.kd
 define amdgpu_kernel void @test_ulong4(<4 x i64> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
     !kernel_arg_base_type !12 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_half8.kd
-; CHECK:        .name:            test_half8
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      half8
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           16
-; CHECK-NEXT:       .value_type:     f16
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           16
+; CHECK-NEXT:         .type_name:      half8
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     f16
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_half8
+; CHECK:          .symbol:         test_half8.kd
 define amdgpu_kernel void @test_half8(<8 x half> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
     !kernel_arg_base_type !13 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_float16.kd
-; CHECK:        .name:            test_float16
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      float16
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           64
-; CHECK-NEXT:       .value_type:     f32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         64
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         72
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         80
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         88
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           64
+; CHECK-NEXT:         .type_name:      float16
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     f32
+; CHECK-NEXT:       - .offset:         64
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         72
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         80
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         88
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_float16
+; CHECK:          .symbol:         test_float16.kd
 define amdgpu_kernel void @test_float16(<16 x float> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
     !kernel_arg_base_type !14 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_double16.kd
-; CHECK:        .name:            test_double16
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      double16
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           128
-; CHECK-NEXT:       .value_type:     f64
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         128
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         136
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         144
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         152
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           128
+; CHECK-NEXT:         .type_name:      double16
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     f64
+; CHECK-NEXT:       - .offset:         128
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         136
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         144
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         152
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_double16
+; CHECK:          .symbol:         test_double16.kd
 define amdgpu_kernel void @test_double16(<16 x double> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
     !kernel_arg_base_type !15 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_pointer.kd
-; CHECK:        .name:            test_pointer
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_pointer
+; CHECK:          .symbol:         test_pointer.kd
 define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16
     !kernel_arg_base_type !16 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_image.kd
-; CHECK:        .name:            test_image
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      image2d_t
-; CHECK-NEXT:       .value_kind:     image
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      image2d_t
+; CHECK-NEXT:         .value_kind:     image
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_image
+; CHECK:          .symbol:         test_image.kd
 define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
     !kernel_arg_base_type !17 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_sampler.kd
-; CHECK:        .name:            test_sampler
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      sampler_t
-; CHECK-NEXT:       .value_kind:     sampler
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      sampler_t
+; CHECK-NEXT:         .value_kind:     sampler
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_sampler
+; CHECK:          .symbol:         test_sampler.kd
 define amdgpu_kernel void @test_sampler(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18
     !kernel_arg_base_type !18 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_queue.kd
-; CHECK:        .name:            test_queue
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      queue_t
-; CHECK-NEXT:       .value_kind:     queue
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      queue_t
+; CHECK-NEXT:         .value_kind:     queue
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_queue
+; CHECK:          .symbol:         test_queue.kd
 define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
     !kernel_arg_base_type !19 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_struct.kd
-; CHECK:        .name:            test_struct
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      struct A
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space: private
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  private
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      struct A
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_struct
+; CHECK:          .symbol:         test_struct.kd
 define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
     !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_i128.kd
-; CHECK:        .name:            test_i128
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      i128
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           16
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           16
+; CHECK-NEXT:         .type_name:      i128
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_i128
+; CHECK:          .symbol:         test_i128.kd
 define amdgpu_kernel void @test_i128(i128 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
     !kernel_arg_base_type !21 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_multi_arg.kd
-; CHECK:        .name:            test_multi_arg
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .type_name:      short2
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         4
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i16
-; CHECK-NEXT:       .name:           b
-; CHECK-NEXT:     - .type_name:      char3
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .name:           c
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .name:           b
+; CHECK-NEXT:         .offset:         4
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      short2
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i16
+; CHECK-NEXT:       - .name:           c
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      char3
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_multi_arg
+; CHECK:          .symbol:         test_multi_arg.kd
 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c)
     !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
     !kernel_arg_base_type !24 !kernel_arg_type_qual !25 {
   ret void
 }
 
-; CHECK:        .symbol:          test_addr_space.kd
-; CHECK:        .name:            test_addr_space
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           g
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           c
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space: constant
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           l
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .pointee_align:  4
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           g
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .address_space:  constant
+; CHECK-NEXT:         .name:           c
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           l
+; CHECK-NEXT:         .offset:         16
+; CHECK-NEXT:         .pointee_align:  4
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_addr_space
+; CHECK:          .symbol:         test_addr_space.kd
 define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
                                            i32 addrspace(4)* %c,
                                            i32 addrspace(3)* %l)
     !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51
     !kernel_arg_base_type !51 !kernel_arg_type_qual !25 {
   ret void
 }
 
-; CHECK:        .symbol:          test_type_qual.kd
-; CHECK:        .name:            test_type_qual
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .is_volatile:    true
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           b
-; CHECK-NEXT:       .is_const:       true
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .is_restrict:    true
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     pipe
-; CHECK-NEXT:       .name:           c
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .is_pipe:        true
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .is_volatile:    true
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .is_const:       true
+; CHECK-NEXT:         .is_restrict:    true
+; CHECK-NEXT:         .name:           b
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .is_pipe:        true
+; CHECK-NEXT:         .name:           c
+; CHECK-NEXT:         .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     pipe
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_type_qual
+; CHECK:          .symbol:         test_type_qual.kd
 define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a,
                                           i32 addrspace(1)* %b,
                                           %opencl.pipe_t addrspace(1)* %c)
     !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51
     !kernel_arg_base_type !51 !kernel_arg_type_qual !70 {
   ret void
 }
 
-; CHECK:        .symbol:          test_access_qual.kd
-; CHECK:        .name:            test_access_qual
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      image1d_t
-; CHECK-NEXT:       .value_kind:     image
-; CHECK-NEXT:       .name:           ro
-; CHECK-NEXT:       .access:         read_only
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      image2d_t
-; CHECK-NEXT:       .value_kind:     image
-; CHECK-NEXT:       .name:           wo
-; CHECK-NEXT:       .access:         write_only
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      image3d_t
-; CHECK-NEXT:       .value_kind:     image
-; CHECK-NEXT:       .name:           rw
-; CHECK-NEXT:       .access:         read_write
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .access:         read_only
+; CHECK-NEXT:         .address_space:  global
+; CHECK-NEXT:         .name:           ro
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      image1d_t
+; CHECK-NEXT:         .value_kind:     image
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .access:         write_only
+; CHECK-NEXT:         .address_space:  global
+; CHECK-NEXT:         .name:           wo
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      image2d_t
+; CHECK-NEXT:         .value_kind:     image
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .access:         read_write
+; CHECK-NEXT:         .address_space:  global
+; CHECK-NEXT:         .name:           rw
+; CHECK-NEXT:         .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      image3d_t
+; CHECK-NEXT:         .value_kind:     image
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_access_qual
+; CHECK:          .symbol:         test_access_qual.kd
 define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro,
                                             %opencl.image2d_t addrspace(1)* %wo,
                                             %opencl.image3d_t addrspace(1)* %rw)
     !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62
     !kernel_arg_base_type !62 !kernel_arg_type_qual !25 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_half.kd
-; CHECK:        .name:            test_vec_type_hint_half
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   half
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_half
+; CHECK:          .symbol:         test_vec_type_hint_half.kd
+; CHECK:          .vec_type_hint:  half
 define amdgpu_kernel void @test_vec_type_hint_half(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_float.kd
-; CHECK:        .name:            test_vec_type_hint_float
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   float
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_float
+; CHECK:          .symbol:         test_vec_type_hint_float.kd
+; CHECK:          .vec_type_hint:  float
 define amdgpu_kernel void @test_vec_type_hint_float(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_double.kd
-; CHECK:        .name:            test_vec_type_hint_double
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   double
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_double
+; CHECK:          .symbol:         test_vec_type_hint_double.kd
+; CHECK:          .vec_type_hint:  double
 define amdgpu_kernel void @test_vec_type_hint_double(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_char.kd
-; CHECK:        .name:            test_vec_type_hint_char
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   char
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_char
+; CHECK:          .symbol:         test_vec_type_hint_char.kd
+; CHECK:          .vec_type_hint:  char
 define amdgpu_kernel void @test_vec_type_hint_char(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_short.kd
-; CHECK:        .name:            test_vec_type_hint_short
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   short
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_short
+; CHECK:          .symbol:         test_vec_type_hint_short.kd
+; CHECK:          .vec_type_hint:  short
 define amdgpu_kernel void @test_vec_type_hint_short(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_long.kd
-; CHECK:        .name:            test_vec_type_hint_long
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   long
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_long
+; CHECK:          .symbol:         test_vec_type_hint_long.kd
+; CHECK:          .vec_type_hint:  long
 define amdgpu_kernel void @test_vec_type_hint_long(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 {
   ret void
 }
 
-; CHECK:        .symbol:          test_vec_type_hint_unknown.kd
-; CHECK:        .name:            test_vec_type_hint_unknown
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      int
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK:        .vec_type_hint:   unknown
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_vec_type_hint_unknown
+; CHECK:          .symbol:         test_vec_type_hint_unknown.kd
+; CHECK:          .vec_type_hint:  unknown
 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 {
   ret void
 }
 
-; CHECK:        .reqd_workgroup_size:
-; CHECK-NEXT:     - 1
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 4
-; CHECK:        .symbol:          test_reqd_wgs_vec_type_hint.kd
-; CHECK:        .name:            test_reqd_wgs_vec_type_hint
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:          int
-; CHECK-NEXT:       .value_kind:         by_value
-; CHECK-NEXT:       .offset:             0
-; CHECK-NEXT:       .size:               4
-; CHECK-NEXT:       .value_type:         i32
-; CHECK-NEXT:       .name:               a
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_x
-; CHECK-NEXT:       .offset:             8
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_y
-; CHECK-NEXT:       .offset:             16
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_z
-; CHECK-NEXT:       .offset:             24
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_printf_buffer
-; CHECK-NEXT:       .offset:             32
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i8
-; CHECK-NEXT:       .address_space:      global
-; CHECK:        .vec_type_hint:       int
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_reqd_wgs_vec_type_hint
+; CHECK:          .reqd_workgroup_size: 
+; CHECK-NEXT:       - 1
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 4
+; CHECK:          .symbol:         test_reqd_wgs_vec_type_hint.kd
+; CHECK:          .vec_type_hint:  int
 define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5
     !reqd_work_group_size !6 {
   ret void
 }
 
-; CHECK:        .symbol:          test_wgs_hint_vec_type_hint.kd
-; CHECK:        .workgroup_size_hint:
-; CHECK-NEXT:     - 8
-; CHECK-NEXT:     - 16
-; CHECK-NEXT:     - 32
-; CHECK:        .name:            test_wgs_hint_vec_type_hint
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:          int
-; CHECK-NEXT:       .value_kind:         by_value
-; CHECK-NEXT:       .offset:             0
-; CHECK-NEXT:       .size:               4
-; CHECK-NEXT:       .value_type:         i32
-; CHECK-NEXT:       .name:               a
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_x
-; CHECK-NEXT:       .offset:             8
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_y
-; CHECK-NEXT:       .offset:             16
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_global_offset_z
-; CHECK-NEXT:       .offset:             24
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i64
-; CHECK-NEXT:     - .value_kind:         hidden_printf_buffer
-; CHECK-NEXT:       .offset:             32
-; CHECK-NEXT:       .size:               8
-; CHECK-NEXT:       .value_type:         i8
-; CHECK-NEXT:       .address_space:      global
-; CHECK:        .vec_type_hint:       uint4
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      int
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_wgs_hint_vec_type_hint
+; CHECK:          .symbol:         test_wgs_hint_vec_type_hint.kd
+; CHECK:          .vec_type_hint:  uint4
+; CHECK:          .workgroup_size_hint: 
+; CHECK-NEXT:       - 8
+; CHECK-NEXT:       - 16
+; CHECK-NEXT:       - 32
 define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7
     !work_group_size_hint !8 {
   ret void
 }
 
-; CHECK:        .symbol:          test_arg_ptr_to_ptr.kd
-; CHECK:        .name:            test_arg_ptr_to_ptr
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'int  addrspace(5)* addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'int  addrspace(5)* addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_arg_ptr_to_ptr
+; CHECK:          .symbol:         test_arg_ptr_to_ptr.kd
 define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a)
     !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
     !kernel_arg_base_type !80 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_arg_struct_contains_ptr.kd
-; CHECK:        .name:            test_arg_struct_contains_ptr
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      struct B
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space: private
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  private
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      struct B
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_arg_struct_contains_ptr
+; CHECK:          .symbol:         test_arg_struct_contains_ptr.kd
 define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B addrspace(5)* byval %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82
     !kernel_arg_base_type !82 !kernel_arg_type_qual !4 {
  ret void
 }
 
-; CHECK:        .symbol:          test_arg_vector_of_ptr.kd
-; CHECK:        .name:            test_arg_vector_of_ptr
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           16
-; CHECK-NEXT:       .value_type:     i32
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           16
+; CHECK-NEXT:         .type_name:      'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i32
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_arg_vector_of_ptr
+; CHECK:          .symbol:         test_arg_vector_of_ptr.kd
 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a)
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
     !kernel_arg_base_type !83 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_arg_unknown_builtin_type.kd
-; CHECK:        .name:            test_arg_unknown_builtin_type
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      clk_event_t
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      clk_event_t
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_arg_unknown_builtin_type
+; CHECK:          .symbol:         test_arg_unknown_builtin_type.kd
 define amdgpu_kernel void @test_arg_unknown_builtin_type(
     %opencl.clk_event_t addrspace(1)* %a)
     !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84
     !kernel_arg_base_type !84 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_pointee_align.kd
-; CHECK:        .name:            test_pointee_align
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      'long  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     global_buffer
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .type_name:      'char  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           b
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  1
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .type_name:      'char2  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           c
-; CHECK-NEXT:       .offset:         12
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  2
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .type_name:      'char3  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           d
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  4
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .type_name:      'char4  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           e
-; CHECK-NEXT:       .offset:         20
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  4
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .type_name:      'char8  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           f
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  8
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .type_name:      'char16  addrspace(5)*'
-; CHECK-NEXT:       .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:       .name:           g
-; CHECK-NEXT:       .offset:         28
-; CHECK-NEXT:       .size:           4
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .pointee_align:  16
-; CHECK-NEXT:       .address_space: local
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         56
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .type_name:      'long  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           b
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .pointee_align:  1
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           c
+; CHECK-NEXT:         .offset:         12
+; CHECK-NEXT:         .pointee_align:  2
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char2  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           d
+; CHECK-NEXT:         .offset:         16
+; CHECK-NEXT:         .pointee_align:  4
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char3  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           e
+; CHECK-NEXT:         .offset:         20
+; CHECK-NEXT:         .pointee_align:  4
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char4  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           f
+; CHECK-NEXT:         .offset:         24
+; CHECK-NEXT:         .pointee_align:  8
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char8  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  local
+; CHECK-NEXT:         .name:           g
+; CHECK-NEXT:         .offset:         28
+; CHECK-NEXT:         .pointee_align:  16
+; CHECK-NEXT:         .size:           4
+; CHECK-NEXT:         .type_name:      'char16  addrspace(5)*'
+; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         56
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_pointee_align
+; CHECK:          .symbol:         test_pointee_align.kd
 define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
                                               i8 addrspace(3)* %b,
                                               <2 x i8> addrspace(3)* %c,
                                               <3 x i8> addrspace(3)* %d,
                                               <4 x i8> addrspace(3)* %e,
                                               <8 x i8> addrspace(3)* %f,
                                               <16 x i8> addrspace(3)* %g)
     !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93
     !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {
   ret void
 }
 
-; CHECK:        .symbol:          __test_block_invoke_kernel.kd
-; CHECK:        .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
-; CHECK:        .name:            __test_block_invoke_kernel
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      __block_literal
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           25
-; CHECK-NEXT:       .value_type:     struct
-; CHECK-NEXT:       .name:           arg
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         56
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           arg
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           25
+; CHECK-NEXT:         .type_name:      __block_literal
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     struct
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         56
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           __test_block_invoke_kernel
+; CHECK:          .symbol:         __test_block_invoke_kernel.kd
 define amdgpu_kernel void @__test_block_invoke_kernel(
     <{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #0
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
     !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK:        .symbol:          test_enqueue_kernel_caller.kd
-; CHECK:        .name:            test_enqueue_kernel_caller
-; CHECK:        .language:        OpenCL C
-; CHECK:        .language_version:
-; CHECK-NEXT:     - 2
-; CHECK-NEXT:     - 0
-; CHECK:        .args:
-; CHECK-NEXT:     - .type_name:      char
-; CHECK-NEXT:       .value_kind:     by_value
-; CHECK-NEXT:       .offset:         0
-; CHECK-NEXT:       .size:           1
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .name:           a
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       .offset:         8
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       .offset:         16
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       .offset:         24
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i64
-; CHECK-NEXT:     - .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:       .offset:         32
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_default_queue
-; CHECK-NEXT:       .offset:         40
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
-; CHECK-NEXT:     - .value_kind:     hidden_completion_action
-; CHECK-NEXT:       .offset:         48
-; CHECK-NEXT:       .size:           8
-; CHECK-NEXT:       .value_type:     i8
-; CHECK-NEXT:       .address_space:  global
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           a
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           1
+; CHECK-NEXT:         .type_name:      char
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_default_queue
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_completion_action
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test_enqueue_kernel_caller
+; CHECK:          .symbol:         test_enqueue_kernel_caller.kd
 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #1
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
     !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
   ret void
 }
 
-; CHECK: .symbol:            unknown_addrspace_kernarg.kd
-; CHECK: .name:              unknown_addrspace_kernarg
-; CHECK: .args:
-; CHECK-NEXT: .value_kind:      global_buffer
-; CHECK-NEXT: .offset:          0
-; CHECK-NEXT: .size:            8
-; CHECK-NEXT: .value_type:      i32
-; CHECK-NEXT: .name:            ptr
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .name:           ptr
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     i32
+; CHECK:          .name:           unknown_addrspace_kernarg
+; CHECK:          .symbol:         unknown_addrspace_kernarg.kd
 define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
   ret void
 }
 
-; CHECK:  amdhsa.version:
-; CHECK-NEXT: - 1
-; CHECK-NEXT: - 0
 ; CHECK:  amdhsa.printf:
 ; CHECK-NEXT: - '1:1:4:%d\n'
 ; CHECK-NEXT: - '2:1:8:%g\n'
+; CHECK:  amdhsa.version:
+; CHECK-NEXT: - 1
+; CHECK-NEXT: - 0
 
 attributes #0 = { "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
 attributes #1 = { "calls-enqueue-kernel" }
 
 !llvm.printf.fmts = !{!100, !101}
 
 !1 = !{i32 0}
 !2 = !{!"none"}

test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll

 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX700 --check-prefix=NOTES %s
 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX803 --check-prefix=NOTES %s
 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 
-; CHECK: ---
-; CHECK:  amdhsa.kernels:
-; CHECK:        .symbol:     test.kd
-; CHECK:        .name:       test
-; CHECK:        .args:
-; CHECK-NEXT:     - .value_kind:      global_buffer
-; CHECK-NEXT:       .name:            r
-; CHECK-NEXT:       .offset:          0
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      f16
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .value_kind:      global_buffer
-; CHECK-NEXT:       .name:            a
-; CHECK-NEXT:       .offset:          8
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      f16
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .value_kind:      global_buffer
-; CHECK-NEXT:       .name:            b
-; CHECK-NEXT:       .offset:          16
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      f16
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .value_kind:      hidden_global_offset_x
-; CHECK-NEXT:       .offset:          24
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i64
-; CHECK-NEXT:     - .value_kind:      hidden_global_offset_y
-; CHECK-NEXT:       .offset:          32
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i64
-; CHECK-NEXT:     - .value_kind:      hidden_global_offset_z
-; CHECK-NEXT:       .offset:          40
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i64
-; CHECK-NEXT:     - .value_kind:      hidden_none
-; CHECK-NEXT:       .offset:          48
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i8
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .value_kind:      hidden_none
-; CHECK-NEXT:       .offset:          56
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i8
-; CHECK-NEXT:       .address_space:   global
-; CHECK-NEXT:     - .value_kind:      hidden_none
-; CHECK-NEXT:       .offset:          64
-; CHECK-NEXT:       .size:            8
-; CHECK-NEXT:       .value_type:      i8
-; CHECK-NEXT:       .address_space:   global
+; CHECK:              ---
+; CHECK:      amdhsa.kernels:  
+; CHECK:        - .args:           
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           r
+; CHECK-NEXT:         .offset:         0
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     f16
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           a
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     f16
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .name:           b
+; CHECK-NEXT:         .offset:         16
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     global_buffer
+; CHECK-NEXT:         .value_type:     f16
+; CHECK-NEXT:       - .offset:         24
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         32
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         40
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         48
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         56
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         64
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_type:     i8
+; CHECK:          .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK:          .name:           test
+; CHECK:          .symbol:         test.kd
 define amdgpu_kernel void @test(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) {
 entry:
   %a.val = load half, half addrspace(1)* %a
   %b.val = load half, half addrspace(1)* %b
   %r.val = fadd half %a.val, %b.val

test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll

 %opencl.image2d_array_depth_t = type opaque
 %opencl.image2d_array_msaa_t = type opaque
 %opencl.image2d_array_msaa_depth_t = type opaque
 %opencl.image2d_depth_t = type opaque
 %opencl.image2d_msaa_t = type opaque
 %opencl.image2d_msaa_depth_t = type opaque
 %opencl.image3d_t = type opaque
 
-; CHECK: ---
-; CHECK:  amdhsa.kernels:
-; CHECK:      .symbol:     test.kd
-; CHECK:      .name:       test
-; CHECK:      .args:
-; CHECK:        - .type_name:  image1d_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       a
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image1d_array_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       b
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image1d_buffer_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       c
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       d
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_array_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       e
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_array_depth_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       f
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_array_msaa_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       g
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_array_msaa_depth_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       h
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_depth_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       i
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_msaa_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       j
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image2d_msaa_depth_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       k
-; CHECK:          .size:       8
-; CHECK:        - .type_name:  image3d_t
-; CHECK:          .value_kind: image
-; CHECK:          .name:       l
-; CHECK:          .size:       8
+; CHECK:         ---
+; CHECK: amdhsa.kernels:  
+; CHECK:   - .args:           
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           a
+; CHECK:         .offset:         0
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image1d_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           b
+; CHECK:         .offset:         8
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image1d_array_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           c
+; CHECK:         .offset:         16
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image1d_buffer_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           d
+; CHECK:         .offset:         24
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           e
+; CHECK:         .offset:         32
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_array_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           f
+; CHECK:         .offset:         40
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_array_depth_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           g
+; CHECK:         .offset:         48
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_array_msaa_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           h
+; CHECK:         .offset:         56
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_array_msaa_depth_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           i
+; CHECK:         .offset:         64
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_depth_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           j
+; CHECK:         .offset:         72
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_msaa_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           k
+; CHECK:         .offset:         80
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image2d_msaa_depth_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
+; CHECK:       - .address_space:  global
+; CHECK:         .name:           l
+; CHECK:         .offset:         88
+; CHECK:         .size:           8
+; CHECK:         .type_name:      image3d_t
+; CHECK:         .value_kind:     image
+; CHECK:         .value_type:     struct
 define amdgpu_kernel void @test(%opencl.image1d_t addrspace(1)* %a,
                                 %opencl.image1d_array_t addrspace(1)* %b,
                                 %opencl.image1d_buffer_t addrspace(1)* %c,
                                 %opencl.image2d_t addrspace(1)* %d,
                                 %opencl.image2d_array_t addrspace(1)* %e,
                                 %opencl.image2d_array_depth_t addrspace(1)* %f,
                                 %opencl.image2d_array_msaa_t addrspace(1)* %g,
                                 %opencl.image2d_array_msaa_depth_t addrspace(1)* %h,

test/CodeGen/AMDGPU/hsa-metadata-kernel-code-props-v3.ll

 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -enable-misched=0 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX700 --check-prefix=NOTES %s
 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -enable-misched=0 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX803 --check-prefix=NOTES %s
 ; RUN: llc -mattr=+code-object-v3 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -enable-misched=0 -filetype=obj -o - < %s | llvm-readobj -elf-output-style=GNU -notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 
 @var = addrspace(1) global float 0.0
 
 ; CHECK: ---
 ; CHECK:  amdhsa.kernels:
 
-; CHECK: - .max_flat_workgroup_size:    256
-; CHECK:   .kernarg_segment_size:       24
-; CHECK:   .private_segment_fixed_size: 0
-; CHECK:   .wavefront_size:             64
-; CHECK:   .symbol:     test.kd
-; CHECK:   .name:       test
-; CHECK:   .sgpr_count:                 8
-; CHECK:   .kernarg_segment_align:      8
-; CHECK:   .vgpr_count:                 6
-; CHECK:   .group_segment_fixed_size:   0
+; CHECK:   - .args:           
+; CHECK:     .group_segment_fixed_size: 0
+; CHECK:     .kernarg_segment_align: 8
+; CHECK:     .kernarg_segment_size: 24
+; CHECK:     .max_flat_workgroup_size: 256
+; CHECK:     .name:           test
+; CHECK:     .private_segment_fixed_size: 0
+; CHECK:     .sgpr_count:     8
+; CHECK:     .symbol:         test.kd
+; CHECK:     .vgpr_count:     6
+; CHECK:     .wavefront_size: 64
 define amdgpu_kernel void @test(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) {
 entry:
   %a.val = load half, half addrspace(1)* %a
   %b.val = load half, half addrspace(1)* %b
   %r.val = fadd half %a.val, %b.val
   store half %r.val, half addrspace(1)* %r
   ret void
 }
 
-; CHECK:   .symbol:     num_spilled_sgprs.kd
 ; CHECK:   .name:       num_spilled_sgprs
 ; GFX700:   .sgpr_spill_count: 40
 ; GFX803:   .sgpr_spill_count: 24
 ; GFX900:   .sgpr_spill_count: 24
+; CHECK:   .symbol:     num_spilled_sgprs.kd
 define amdgpu_kernel void @num_spilled_sgprs(
     i32 addrspace(1)* %out0, i32 addrspace(1)* %out1, [8 x i32],
     i32 addrspace(1)* %out2, i32 addrspace(1)* %out3, [8 x i32],
     i32 addrspace(1)* %out4, i32 addrspace(1)* %out5, [8 x i32],
     i32 addrspace(1)* %out6, i32 addrspace(1)* %out7, [8 x i32],
     i32 addrspace(1)* %out8, i32 addrspace(1)* %out9, [8 x i32],
     i32 addrspace(1)* %outa, i32 addrspace(1)* %outb, [8 x i32],
     i32 addrspace(1)* %outc, i32 addrspace(1)* %outd, [8 x i32],
entry:
   store i32 %inb, i32 addrspace(1)* %outb
   store i32 %inc, i32 addrspace(1)* %outc
   store i32 %ind, i32 addrspace(1)* %outd
   store i32 %ine, i32 addrspace(1)* %oute
   store i32 %inf, i32 addrspace(1)* %outf
   ret void
 }
 
-; CHECK:   .symbol:     num_spilled_vgprs.kd
 ; CHECK:   .name:       num_spilled_vgprs
+; CHECK:   .symbol:     num_spilled_vgprs.kd
 ; CHECK:   .vgpr_spill_count: 14
 define amdgpu_kernel void @num_spilled_vgprs() #1 {
   %val0 = load volatile float, float addrspace(1)* @var
   %val1 = load volatile float, float addrspace(1)* @var
   %val2 = load volatile float, float addrspace(1)* @var
   %val3 = load volatile float, float addrspace(1)* @var
   %val4 = load volatile float, float addrspace(1)* @var
   %val5 = load volatile float, float addrspace(1)* @var

test/MC/AMDGPU/hsa-metadata-kernel-args-v3.s

 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
 
-// CHECK:  .amdgpu_metadata
-// CHECK:    amdhsa.kernels:
-// CHECK-NEXT:      - .max_flat_workgroup_size: 256
-// CHECK-NEXT:        .wavefront_size: 128
-// CHECK-NEXT:        .symbol:      'test_kernel@kd'
-// CHECK-NEXT:        .kernarg_segment_size: 8
-// CHECK-NEXT:        .private_segment_fixed_size: 32
-// CHECK-NEXT:        .name:            test_kernel
-// CHECK-NEXT:        .language:        OpenCL C
-// CHECK-NEXT:        .sgpr_count: 14
-// CHECK-NEXT:        .kernarg_segment_align: 64
-// CHECK-NEXT:        .vgpr_count: 40
-// CHECK-NEXT:        .group_segment_fixed_size: 16
-// CHECK-NEXT:        .language_version:
-// CHECK-NEXT:          - 2
-// CHECK-NEXT:          - 0
-// CHECK-NEXT:        .args:
-// CHECK-NEXT:          - .type_name:      char
-// CHECK-NEXT:            .value_kind:     by_value
-// CHECK-NEXT:            .offset:         1
-// CHECK-NEXT:            .size:          1
-// CHECK-NEXT:            .value_type:     i8
-// CHECK-NEXT:          - .value_kind:     hidden_global_offset_x
-// CHECK-NEXT:            .offset:         8
-// CHECK-NEXT:            .size:          8
-// CHECK-NEXT:            .value_type:     i64
-// CHECK-NEXT:          - .value_kind:     hidden_global_offset_y
-// CHECK-NEXT:            .offset:         8
-// CHECK-NEXT:            .size:          8
-// CHECK-NEXT:            .value_type:     i64
-// CHECK-NEXT:          - .value_kind:     hidden_global_offset_z
-// CHECK-NEXT:            .offset:         8
-// CHECK-NEXT:            .size:          8
-// CHECK-NEXT:            .value_type:     i64
-// CHECK-NEXT:          - .value_kind:     hidden_printf_buffer
-// CHECK-NEXT:            .offset:         8
-// CHECK-NEXT:            .size:          8
-// CHECK-NEXT:            .value_type:     i8
-// CHECK-NEXT:            .address_space: global
-// CHECK:    amdhsa.version:
-// CHECK-NEXT:       - 1
-// CHECK-NEXT:       - 0
-// CHECK:    amdhsa.printf:
-// CHECK-NEXT:      - '1:1:4:%d\n'
-// CHECK-NEXt:      - '2:1:8:%g\n'
-// CHECK:  .end_amdgpu_metadata
+; CHECK:        .amdgpu_metadata
+; CHECK:      amdhsa.kernels:  
+; CHECK-NEXT:   - .args:           
+; CHECK-NEXT:       - .offset:         1
+; CHECK-NEXT:         .size:           1
+; CHECK-NEXT:         .type_name:      char
+; CHECK-NEXT:         .value_kind:     by_value
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
+; CHECK-NEXT:         .value_type:     i64
+; CHECK-NEXT:       - .address_space:  global
+; CHECK-NEXT:         .offset:         8
+; CHECK-NEXT:         .size:           8
+; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
+; CHECK-NEXT:         .value_type:     i8
+; CHECK-NEXT:     .group_segment_fixed_size: 16
+; CHECK-NEXT:     .kernarg_segment_align: 64
+; CHECK-NEXT:     .kernarg_segment_size: 8
+; CHECK-NEXT:     .language:       OpenCL C
+; CHECK-NEXT:     .language_version: 
+; CHECK-NEXT:       - 2
+; CHECK-NEXT:       - 0
+; CHECK-NEXT:     .max_flat_workgroup_size: 256
+; CHECK-NEXT:     .name:           test_kernel
+; CHECK-NEXT:     .private_segment_fixed_size: 32
+; CHECK-NEXT:     .sgpr_count:     14
+; CHECK-NEXT:     .symbol:         'test_kernel@kd'
+; CHECK-NEXT:     .vgpr_count:     40
+; CHECK-NEXT:     .wavefront_size: 128
+; CHECK-NEXT: amdhsa.printf:   
+; CHECK-NEXT:   - '1:1:4:%d\n'
+; CHECK-NEXT:   - '2:1:8:%g\n'
+; CHECK-NEXT: amdhsa.version:  
+; CHECK-NEXT:   - 1
+; CHECK-NEXT:   - 0
+; CHECK:        .end_amdgpu_metadata
 .amdgpu_metadata
   amdhsa.version:
     - 1
     - 0
   amdhsa.printf:
     - '1:1:4:%d\n'
     - '2:1:8:%g\n'
   amdhsa.kernels:

test/MC/AMDGPU/hsa-metadata-kernel-attrs-v3.s

 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
 
-// CHECK:  .amdgpu_metadata
-// CHECK:    amdhsa.kernels:
-// CHECK:      - .max_flat_workgroup_size: 256
-// CHECK:        .wavefront_size: 128
-// CHECK:        .symbol:      'test_kernel@kd'
-// CHECK:        .reqd_workgroup_size:
-// CHECK-NEXT:     - 1
-// CHECK-NEXT:     - 2
-// CHECK-NEXT:     - 4
-// CHECK:        .kernarg_segment_size: 8
-// CHECK:        .private_segment_fixed_size: 32
-// CHECK:        .workgroup_size_hint:
-// CHECK-NEXT:     - 8
-// CHECK-NEXT:     - 16
-// CHECK-NEXT:     - 32
-// CHECK:        .name:            test_kernel
-// CHECK:        .language:        OpenCL C
-// CHECK:        .sgpr_count:      14
-// CHECK:        .kernarg_segment_align: 64
-// CHECK:        .vgpr_count:      40
-// CHECK:        .language_version:
-// CHECK-NEXT:     - 2
-// CHECK-NEXT:     - 0
-// CHECK:        .vec_type_hint:       int
-// CHECK:    amdhsa.version:
-// CHECK-NEXT: - 1
-// CHECK-NEXT: - 0
-// CHECK:    amdhsa.printf:
-// CHECK:      - '1:1:4:%d\n'
-// CHECK:      - '2:1:8:%g\n'
-// CHECK: .end_amdgpu_metadata
+// CHECK:       .amdgpu_metadata
+// CHECK:      amdhsa.kernels:  
+// CHECK:        - .group_segment_fixed_size: 16
+// CHECK:          .kernarg_segment_align: 64
+// CHECK:          .kernarg_segment_size: 8
+// CHECK:          .language:       OpenCL C
+// CHECK:          .language_version: 
+// CHECK-NEXT:       - 2
+// CHECK-NEXT:       - 0
+// CHECK:          .max_flat_workgroup_size: 256
+// CHECK:          .name:           test_kernel
+// CHECK:          .private_segment_fixed_size: 32
+// CHECK:          .reqd_workgroup_size: 
+// CHECK-NEXT:       - 1
+// CHECK-NEXT:       - 2
+// CHECK-NEXT:       - 4
+// CHECK:          .sgpr_count:     14
+// CHECK:          .symbol:         'test_kernel@kd'
+// CHECK:          .vec_type_hint:  int
+// CHECK:          .vgpr_count:     40
+// CHECK:          .wavefront_size: 128
+// CHECK:          .workgroup_size_hint: 
+// CHECK-NEXT:       - 8
+// CHECK-NEXT:       - 16
+// CHECK-NEXT:       - 32
+// CHECK:      amdhsa.printf:   
+// CHECK:        - '1:1:4:%d\n'
+// CHECK:        - '2:1:8:%g\n'
+// CHECK:      amdhsa.version:  
+// CHECK-NEXT:   - 1
+// CHECK-NEXT:   - 0
+// CHECK:       .end_amdgpu_metadata
 .amdgpu_metadata
   amdhsa.version:
     - 1
     - 0
   amdhsa.printf:
     - '1:1:4:%d\n'
     - '2:1:8:%g\n'
   amdhsa.kernels:

test/MC/AMDGPU/hsa-metadata-kernel-code-props-v3.s

 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
 // RUN: llvm-mc -mattr=+code-object-v3 -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
 
-// CHECK:  .amdgpu_metadata
-// CHECK:    amdhsa.kernels:
-// CHECK:      - .sgpr_count: 40
-// CHECK:        .max_flat_workgroup_size:    256
-// CHECK:        .symbol: 'test_kernel@kd'
-// CHECK:        .kernarg_segment_size:      24
-// CHECK:        .group_segment_fixed_size:   24
-// CHECK:        .private_segment_fixed_size: 16
-// CHECK:        .vgpr_count: 14
-// CHECK:        .vgpr_spill_count: 1
-// CHECK:        .kernarg_segment_align:     16
-// CHECK:        .sgpr_spill_count: 1
-// CHECK:        .wavefront_size:           64
-// CHECK:        .name:       test_kernel
-// CHECK:    amdhsa.version:
-// CHECK-NEXT: - 1
-// CHECK-NEXT: - 0
+// CHECK:       .amdgpu_metadata
+// CHECK:      amdhsa.kernels:  
+// CHECK:        - .group_segment_fixed_size: 24
+// CHECK:          .kernarg_segment_align: 16
+// CHECK:          .kernarg_segment_size: 24
+// CHECK:          .max_flat_workgroup_size: 256
+// CHECK:          .name:           test_kernel
+// CHECK:          .private_segment_fixed_size: 16
+// CHECK:          .sgpr_count:     40
+// CHECK:          .sgpr_spill_count: 1
+// CHECK:          .symbol:         'test_kernel@kd'
+// CHECK:          .vgpr_count:     14
+// CHECK:          .vgpr_spill_count: 1
+// CHECK:          .wavefront_size: 64
+// CHECK:      amdhsa.version:  
+// CHECK-NEXT:   - 1
+// CHECK-NEXT:   - 0
 .amdgpu_metadata
   amdhsa.version:
     - 1
     - 0
   amdhsa.printf:
     - '1:1:4:%d\n'
     - '2:1:8:%g\n'
   amdhsa.kernels:

test/MC/AMDGPU/hsa-v3.s

- .name:       amd_kernel_code_t_minimal
       .private_segment_fixed_size: 32
       .kernarg_segment_align: 64
       .wavefront_size: 128
       .sgpr_count: 14
       .vgpr_count: 40
       .max_flat_workgroup_size: 256
 .end_amdgpu_metadata
 
-// ASM: .amdgpu_metadata
-// ASM:    amdhsa.kernels:
-// ASM:  - .sgpr_count:     14
-// ASM:    .max_flat_workgroup_size: 256
-// ASM:    .symbol:         'amd_kernel_code_t_test_all@kd'
-// ASM:    .kernarg_segment_size: 8
-// ASM:    .group_segment_fixed_size: 16
-// ASM:    .private_segment_fixed_size: 32
-// ASM:    .vgpr_count:     40
-// ASM:    .kernarg_segment_align: 64
-// ASM:    .wavefront_size: 128
-// ASM:    .name:           amd_kernel_code_t_test_all
-// ASM:  - .sgpr_count:     14
-// ASM:    .max_flat_workgroup_size: 256
-// ASM:    .symbol:         'amd_kernel_code_t_minimal@kd'
-// ASM:    .kernarg_segment_size: 8
-// ASM:    .group_segment_fixed_size: 16
-// ASM:    .private_segment_fixed_size: 32
-// ASM:    .vgpr_count:     40
-// ASM:    .kernarg_segment_align: 64
-// ASM:    .wavefront_size: 128
-// ASM:    .name:           amd_kernel_code_t_minimal
-// ASM:    amdhsa.version:
-// ASM-NEXT: - 3
-// ASM-NEXT: - 0
-// ASM: .end_amdgpu_metadata
+// ASM:       .amdgpu_metadata
+// ASM:      amdhsa.kernels:  
+// ASM:        - .group_segment_fixed_size: 16
+// ASM:          .kernarg_segment_align: 64
+// ASM:          .kernarg_segment_size: 8
+// ASM:          .max_flat_workgroup_size: 256
+// ASM:          .name:           amd_kernel_code_t_test_all
+// ASM:          .private_segment_fixed_size: 32
+// ASM:          .sgpr_count:     14
+// ASM:          .symbol:         'amd_kernel_code_t_test_all@kd'
+// ASM:          .vgpr_count:     40
+// ASM:          .wavefront_size: 128
+// ASM:        - .group_segment_fixed_size: 16
+// ASM:          .kernarg_segment_align: 64
+// ASM:          .kernarg_segment_size: 8
+// ASM:          .max_flat_workgroup_size: 256
+// ASM:          .name:           amd_kernel_code_t_minimal
+// ASM:          .private_segment_fixed_size: 32
+// ASM:          .sgpr_count:     14
+// ASM:          .symbol:         'amd_kernel_code_t_minimal@kd'
+// ASM:          .vgpr_count:     40
+// ASM:          .wavefront_size: 128
+// ASM:      amdhsa.version:  
+// ASM-NEXT:   - 3
+// ASM-NEXT:   - 0
+// ASM:       .end_amdgpu_metadata

tools/llvm-readobj/ELFDumper.cpp

struct AMDGPUNote {
   std::string Value;
 };
 
 template <typename ELFT>
 static AMDGPUNote getAMDGPUNote(uint32_t NoteType, ArrayRef<uint8_t> Desc) {
   switch (NoteType) {
   default:
     return {"", ""};
-  case ELF::NT_AMDGPU_METADATA:
+  case ELF::NT_AMDGPU_METADATA: {
     auto MsgPackString =
         StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());
-    msgpack::Reader MsgPackReader(MsgPackString);
-    auto OptMsgPackNodeOrErr = msgpack::Node::read(MsgPackReader);
-    if (errorToBool(OptMsgPackNodeOrErr.takeError()))
+    msgpack::Document MsgPackDoc;
+    if (!MsgPackDoc.readFromBlob(MsgPackString, /*Multi=*/false))
       return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
-    auto &OptMsgPackNode = *OptMsgPackNodeOrErr;
-    if (!OptMsgPackNode)
-      return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
-    auto &MsgPackNode = *OptMsgPackNode;
 
     AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
-    if (!Verifier.verify(*MsgPackNode))
+    if (!Verifier.verify(MsgPackDoc.getRoot()))
       return {"AMDGPU Metadata", "Invalid AMDGPU Metadata"};
 
     std::string HSAMetadataString;
     raw_string_ostream StrOS(HSAMetadataString);
-    yaml::Output YOut(StrOS);
-    YOut << MsgPackNode;
+    MsgPackDoc.toYAML(StrOS);
 
     return {"AMDGPU Metadata", StrOS.str()};
   }
   }
+}
 
 template <class ELFT>
 void GNUStyle<ELFT>::printNotes(const ELFFile<ELFT> *Obj) {
   auto PrintHeader = [&](const typename ELFT::Off Offset,
                          const typename ELFT::Addr Size) {
     OS << "Displaying notes found at file offset " << format_hex(Offset, 10)
        << " with length " << format_hex(Size, 10) << ":\n"
        << "  Owner                 Data size\tDescription\n";