AMDGPU: Remove .value_type from kernel metadata

[arsenm]

Summary

This doesn't appear used for anything, and is emitted incorrectly
based on the description. This also depends on the IR type, and
pointee element type.

[arsenm]

Patch committed to remove the use from rocclr

[kzhuravl]

Do we also want to remove it from v2 metadata?

[arsenm]

In D82818#2141973, @kzhuravl wrote:

Do we also want to remove it from v2 metadata?

Probably, but I looked briefly and didn't actually see the direct equivalent

[b-sumner]

In D82818#2141981, @arsenm wrote:

In D82818#2141973, @kzhuravl wrote:

Do we also want to remove it from v2 metadata?

Probably, but I looked briefly and didn't actually see the direct equivalent

It's "ValueType" in the v2 metadata.

[arsenm]

Remove v2, and also accept parsing

[scott.linder]

llvm/docs/AMDGPUUsage.rst

2321

Should we delete this as well? Or at least mark it as non-Required?

[arsenm]

Update documentation

[kzhuravl]

llvm/docs/AMDGPUUsage.rst

2321

It looks like (according to phab diff) you removed space between "ValueType" and 'string', which will break the table.

2339

compatibility

2820–2842

Should same text apply as above? "Unused and deprecated. This should no longer..."

llvm/include/llvm/Support/AMDGPUMetadata.h

82–83

compatibility

llvm/lib/Support/AMDGPUMetadata.cpp

115

same as above

[arsenm]

Documentation fixes

[kzhuravl]

lgtm

[arsenm]

31f4e43f3f391e5c5034580f972e0acc78f99b63

llvm/docs/AMDGPUUsage.rst

"ValueKind"       string         Required  Kernel argument kind that
                                                   the ancestor tree for determining
                                                   when the parent kernel has finished.
 
                                                 "HiddenMultiGridSyncArg"
                                                   A global address space pointer for
                                                   multi-grid synchronization is
                                                   passed in the kernarg.
 
      "ValueType"       string         Required  Kernel argument value type. Only
                                                 present if "ValueKind" is
                                                 "ByValue". For vector data
                                                 types, the value is for the
                                                 element type. Values include:
 
                                                 - "Struct"
                                                 - "I8"
                                                 - "U8"
                                                 - "I16"
                                                 - "U16"
                                                 - "F16"
                                                 - "I32"
                                                 - "U32"
                                                 - "F32"
                                                 - "I64"
                                                 - "U64"
                                                 - "F64"
 
                                                 .. TODO::
                                                    How can it be determined if a
                                                    vector type, and what size
                                                    vector?
      "PointeeAlign"    integer                  Alignment in bytes of pointee
                                                 type for pointer type kernel
                                                 argument. Must be a power
                                                 of 2. Only present if
".value_kind"          string         Required  Kernel argument kind that
                                                        the ancestor tree for determining
                                                        when the parent kernel has finished.
 
                                                      "hidden_multigrid_sync_arg"
                                                        A global address space pointer for
                                                        multi-grid synchronization is
                                                        passed in the kernarg.
 
-     ".value_type"          string         Required  Kernel argument value type. Only
-                                                     present if ".value_kind" is
-                                                     "by_value". For vector data
-                                                     types, the value is for the
-                                                     element type. Values include:
-
-                                                     - "struct"
-                                                     - "i8"
-                                                     - "u8"
-                                                     - "i16"
-                                                     - "u16"
-                                                     - "f16"
-                                                     - "i32"
-                                                     - "u32"
-                                                     - "f32"
-                                                     - "i64"
-                                                     - "u64"
-                                                     - "f64"
-
-                                                     .. TODO::
-                                                        How can it be determined if a
-                                                        vector type, and what size
-                                                        vector?
      ".pointee_align"       integer                  Alignment in bytes of pointee
                                                      type for pointer type kernel
                                                      argument. Must be a power
                                                      of 2. Only present if
                                                      ".value_kind" is
                                                      "dynamic_shared_pointer".
      ".address_space"       string                   Kernel argument address space
                                                      qualifier. Only present if

llvm/include/llvm/Support/AMDGPUMetadata.h

enum class ValueKind : uint8_t {
   HiddenPrintfBuffer     = 11,
   HiddenDefaultQueue     = 12,
   HiddenCompletionAction = 13,
   HiddenMultiGridSyncArg = 14,
   HiddenHostcallBuffer   = 15,
   Unknown                = 0xff
 };
 
-/// Value types.
+/// Value types. This is deprecated and only remains for compatability parsing
+/// of old metadata.
 enum class ValueType : uint8_t {
   Struct  = 0,
   I8      = 1,
   U8      = 2,
   I16     = 3,
   U16     = 4,
   F16     = 5,
   I32     = 6,

 /// Key for Kernel::Arg::Metadata::mSize.
 constexpr char Size[] = "Size";
 /// Key for Kernel::Arg::Metadata::mOffset.
 constexpr char Offset[] = "Offset";
 /// Key for Kernel::Arg::Metadata::mAlign.
 constexpr char Align[] = "Align";
 /// Key for Kernel::Arg::Metadata::mValueKind.
 constexpr char ValueKind[] = "ValueKind";
-/// Key for Kernel::Arg::Metadata::mValueType.
+/// Key for Kernel::Arg::Metadata::mValueType. (deprecated)
 constexpr char ValueType[] = "ValueType";
 /// Key for Kernel::Arg::Metadata::mPointeeAlign.
 constexpr char PointeeAlign[] = "PointeeAlign";
 /// Key for Kernel::Arg::Metadata::mAddrSpaceQual.
 constexpr char AddrSpaceQual[] = "AddrSpaceQual";
 /// Key for Kernel::Arg::Metadata::mAccQual.
 constexpr char AccQual[] = "AccQual";
 /// Key for Kernel::Arg::Metadata::mActualAccQual.
struct Metadata final {
   /// Size in bytes. Required.
   uint32_t mSize = 0;
   /// Offset in bytes. Required for code object v3, unused for code object v2.
   uint32_t mOffset = 0;
   /// Alignment in bytes. Required.
   uint32_t mAlign = 0;
   /// Value kind. Required.
   ValueKind mValueKind = ValueKind::Unknown;
-  /// Value type. Required.
-  ValueType mValueType = ValueType::Unknown;
   /// Pointee alignment in bytes. Optional.
   uint32_t mPointeeAlign = 0;
   /// Address space qualifier. Optional.
   AddressSpaceQualifier mAddrSpaceQual = AddressSpaceQualifier::Unknown;
   /// Access qualifier. Optional.
   AccessQualifier mAccQual = AccessQualifier::Unknown;
   /// Actual access qualifier. Optional.
   AccessQualifier mActualAccQual = AccessQualifier::Unknown;

llvm/lib/BinaryFormat/AMDGPUMetadataVerifier.cpp

if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
                                .Case("hidden_printf_buffer", true)
                                .Case("hidden_hostcall_buffer", true)
                                .Case("hidden_default_queue", true)
                                .Case("hidden_completion_action", true)
                                .Case("hidden_multigrid_sync_arg", true)
                                .Default(false);
                          }))
     return false;
-  if (!verifyScalarEntry(ArgsMap, ".value_type", true,
-                         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::Type::String,
                          [](msgpack::DocNode &SNode) {
                            return StringSwitch<bool>(SNode.getString())
                                .Case("private", true)
                                .Case("global", true)

llvm/lib/Support/AMDGPUMetadata.cpp

 template <>
 struct MappingTraits<Kernel::Arg::Metadata> {
   static void mapping(IO &YIO, Kernel::Arg::Metadata &MD) {
     YIO.mapOptional(Kernel::Arg::Key::Name, MD.mName, std::string());
     YIO.mapOptional(Kernel::Arg::Key::TypeName, MD.mTypeName, std::string());
     YIO.mapRequired(Kernel::Arg::Key::Size, MD.mSize);
     YIO.mapRequired(Kernel::Arg::Key::Align, MD.mAlign);
     YIO.mapRequired(Kernel::Arg::Key::ValueKind, MD.mValueKind);
-    YIO.mapRequired(Kernel::Arg::Key::ValueType, MD.mValueType);
+
+    // Removed. Accepted for parsing compatability, but not emitted.
+    Optional<ValueType> Unused;
+    YIO.mapOptional(Kernel::Arg::Key::ValueType, Unused);
+
     YIO.mapOptional(Kernel::Arg::Key::PointeeAlign, MD.mPointeeAlign,
                     uint32_t(0));
     YIO.mapOptional(Kernel::Arg::Key::AddrSpaceQual, MD.mAddrSpaceQual,
                     AddressSpaceQualifier::Unknown);
     YIO.mapOptional(Kernel::Arg::Key::AccQual, MD.mAccQual,
                     AccessQualifier::Unknown);
     YIO.mapOptional(Kernel::Arg::Key::ActualAccQual, MD.mActualAccQual,
                     AccessQualifier::Unknown);

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.h

private:
 
   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;
 
   msgpack::ArrayDocNode getWorkGroupDimensions(MDNode *Node) const;
 
   msgpack::MapDocNode getHSAKernelProps(const MachineFunction &MF,
                                         const SIProgramInfo &ProgramInfo) const;
 
   void emitVersion();
private:
 
   AccessQualifier getAccessQualifier(StringRef AccQual) const;
 
   AddressSpaceQualifier getAddressSpaceQualifier(unsigned AddressSpace) const;
 
   ValueKind getValueKind(Type *Ty, StringRef TypeQual,
                          StringRef BaseTypeName) const;
 
-  ValueType getValueType(Type *Ty, StringRef TypeName) const;
-
   std::string getTypeName(Type *Ty, bool Signed) const;
 
   std::vector<uint32_t> getWorkGroupDimensions(MDNode *Node) const;
 
   Kernel::CodeProps::Metadata getHSACodeProps(
       const MachineFunction &MF,
       const SIProgramInfo &ProgramInfo) const;
   Kernel::DebugProps::Metadata getHSADebugProps(

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

return StringSwitch<ValueKind>(BaseTypeName)
              .Default(isa<PointerType>(Ty) ?
                           (Ty->getPointerAddressSpace() ==
                            AMDGPUAS::LOCAL_ADDRESS ?
                            ValueKind::DynamicSharedPointer :
                            ValueKind::GlobalBuffer) :
                       ValueKind::ByValue);
 }
 
-ValueType MetadataStreamerV2::getValueType(Type *Ty, StringRef TypeName) const {
-  switch (Ty->getTypeID()) {
-  case Type::IntegerTyID: {
-    auto Signed = !TypeName.startswith("u");
-    switch (Ty->getIntegerBitWidth()) {
-    case 8:
-      return Signed ? ValueType::I8 : ValueType::U8;
-    case 16:
-      return Signed ? ValueType::I16 : ValueType::U16;
-    case 32:
-      return Signed ? ValueType::I32 : ValueType::U32;
-    case 64:
-      return Signed ? ValueType::I64 : ValueType::U64;
-    default:
-      return ValueType::Struct;
-    }
-  }
-  case Type::HalfTyID:
-    return ValueType::F16;
-  case Type::FloatTyID:
-    return ValueType::F32;
-  case Type::DoubleTyID:
-    return ValueType::F64;
-  case Type::PointerTyID:
-    return getValueType(Ty->getPointerElementType(), TypeName);
-  case Type::FixedVectorTyID:
-    return getValueType(cast<VectorType>(Ty)->getElementType(), TypeName);
-  default:
-    return ValueType::Struct;
-  }
-}
-
 std::string MetadataStreamerV2::getTypeName(Type *Ty, bool Signed) const {
   switch (Ty->getTypeID()) {
   case Type::IntegerTyID: {
     if (!Signed)
       return (Twine('u') + getTypeName(Ty, true)).str();
 
     auto BitWidth = Ty->getIntegerBitWidth();
     switch (BitWidth) {
void MetadataStreamerV2::emitKernelArg(const DataLayout &DL, Type *Ty,
   HSAMetadata.mKernels.back().mArgs.push_back(Kernel::Arg::Metadata());
   auto &Arg = HSAMetadata.mKernels.back().mArgs.back();
 
   Arg.mName = std::string(Name);
   Arg.mTypeName = std::string(TypeName);
   Arg.mSize = DL.getTypeAllocSize(Ty);
   Arg.mAlign = DL.getABITypeAlign(Ty).value();
   Arg.mValueKind = ValueKind;
-  Arg.mValueType = getValueType(Ty, BaseTypeName);
   Arg.mPointeeAlign = PointeeAlign ? PointeeAlign->value() : 0;
 
   if (auto PtrTy = dyn_cast<PointerType>(Ty))
     Arg.mAddrSpaceQual = getAddressSpaceQualifier(PtrTy->getAddressSpace());
 
   Arg.mAccQual = getAccessQualifier(AccQual);
 
   // TODO: Emit Arg.mActualAccQual.
return StringSwitch<StringRef>(BaseTypeName)
       .Case("queue_t", "queue")
       .Default(isa<PointerType>(Ty)
                    ? (Ty->getPointerAddressSpace() == AMDGPUAS::LOCAL_ADDRESS
                           ? "dynamic_shared_pointer"
                           : "global_buffer")
                    : "by_value");
 }
 
-StringRef MetadataStreamerV3::getValueType(Type *Ty, StringRef TypeName) const {
-  switch (Ty->getTypeID()) {
-  case Type::IntegerTyID: {
-    auto Signed = !TypeName.startswith("u");
-    switch (Ty->getIntegerBitWidth()) {
-    case 8:
-      return Signed ? "i8" : "u8";
-    case 16:
-      return Signed ? "i16" : "u16";
-    case 32:
-      return Signed ? "i32" : "u32";
-    case 64:
-      return Signed ? "i64" : "u64";
-    default:
-      return "struct";
-    }
-  }
-  case Type::HalfTyID:
-    return "f16";
-  case Type::FloatTyID:
-    return "f32";
-  case Type::DoubleTyID:
-    return "f64";
-  case Type::PointerTyID:
-    return getValueType(Ty->getPointerElementType(), TypeName);
-  case Type::FixedVectorTyID:
-    return getValueType(cast<VectorType>(Ty)->getElementType(), TypeName);
-  default:
-    return "struct";
-  }
-}
-
 std::string MetadataStreamerV3::getTypeName(Type *Ty, bool Signed) const {
   switch (Ty->getTypeID()) {
   case Type::IntegerTyID: {
     if (!Signed)
       return (Twine('u') + getTypeName(Ty, true)).str();
 
     auto BitWidth = Ty->getIntegerBitWidth();
     switch (BitWidth) {
if (!TypeName.empty())
     Arg[".type_name"] = Arg.getDocument()->getNode(TypeName, /*Copy=*/true);
   auto Size = DL.getTypeAllocSize(Ty);
   Align Alignment = DL.getABITypeAlign(Ty);
   Arg[".size"] = Arg.getDocument()->getNode(Size);
   Offset = alignTo(Offset, Alignment);
   Arg[".offset"] = Arg.getDocument()->getNode(Offset);
   Offset += Size;
   Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
-  Arg[".value_type"] =
-      Arg.getDocument()->getNode(getValueType(Ty, BaseTypeName), /*Copy=*/true);
   if (PointeeAlign)
     Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
 
   if (auto PtrTy = dyn_cast<PointerType>(Ty))
     if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
       Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true);
 
   if (auto AQ = getAccessQualifier(AccQual))

llvm/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-readelf --notes | FileCheck %s
 
-; CHECK:        - .args:           
+; 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", !""}

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

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck %s
 
 ; CHECK:      - Name:            test_ro_arg
 ; CHECK-NEXT:   SymbolName:      'test_ro_arg@kd'
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT: - Name:            in
 ; CHECK-NEXT:   TypeName:        'float*'
 ; CHECK-NEXT:   Size:            8
 ; CHECK-NEXT:   Align:           8
 ; CHECK-NEXT:   ValueKind:       GlobalBuffer
-; CHECK-NEXT:   ValueType:       F32
 ; CHECK-NEXT:   AddrSpaceQual:   Global
 ; CHECK-NEXT:   AccQual:         ReadOnly
 ; CHECK-NEXT:   IsConst:         true
 ; CHECK-NEXT:   IsRestrict:      true
 ; CHECK-NEXT: - Name:            out
 ; CHECK-NEXT:   TypeName:        'float*'
 ; CHECK-NEXT:   Size:            8
 ; CHECK-NEXT:   Align:           8
 ; CHECK-NEXT:   ValueKind:       GlobalBuffer
-; CHECK-NEXT:   ValueType:       F32
 ; CHECK-NEXT:   AddrSpaceQual:   Global
 ; CHECK-NEXT:   AccQual:         Default
 
 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", !""}

llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel-v3.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -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:        - .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) #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:        - .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) #1
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3

llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel.ll

 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
 ; CHECK-NOT:        ValueKind:     HiddenCompletionAction
 define amdgpu_kernel void @test_non_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:      - Name:            test_enqueue_kernel_caller
 ; CHECK-NEXT:   SymbolName:      'test_enqueue_kernel_caller@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenDefaultQueue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenCompletionAction
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #1
     !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
 }
 
 attributes #0 = { "amdgpu-implicitarg-num-bytes"="48" }

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

 ; 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-NOT:          .value_kind:     hidden_hostcall_buffer
 ; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK:              .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         80
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         96
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         104
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         112
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         160
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         168
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         176
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .args:
 ; CHECK-NEXT:         .name:           a
 ; CHECK-NEXT:         .offset:         0
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .type_name:      struct A
 ; CHECK-NEXT:         .value_kind:     by_value
-; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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 %a) #0
     !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:        - .args:
 ; CHECK-NEXT:         .name:           a
 ; CHECK-NEXT:         .offset:         0
 ; CHECK-NEXT:         .size:           32
 ; CHECK-NEXT:         .type_name:      struct A
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         80
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK:          .language:       OpenCL C
 ; CHECK-NEXT:     .language_version:
 ; CHECK-NEXT:       - 2
 ; CHECK-NEXT:       - 0
 ; CHECK:          .name:           test_array
 ; CHECK:          .symbol:         test_array.kd
 define amdgpu_kernel void @test_array([32 x i8] %a) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !26 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !27 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !28 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !29 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !30 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !31 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !32 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #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 !vec_type_hint !5
     !reqd_work_group_size !6 {
   ret void
 }
 
 ; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .args:
 ; CHECK-NEXT:         .name:           a
 ; CHECK-NEXT:         .offset:         0
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .type_name:      struct B
 ; CHECK-NEXT:         .value_kind:     by_value
-; 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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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 %a) #0
     !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:        - .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-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_multigrid_sync_arg
-; 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) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; 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_multigrid_sync_arg
-; 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) #0
     !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:        - .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:       - .address_space:  local
 ; CHECK-NEXT:         .name:           h
 ; CHECK-NEXT:         .offset:         32
 ; CHECK-NEXT:         .pointee_align:  1
 ; CHECK-NEXT:         .size:           4
 ; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:         .value_type:     struct
 ; CHECK-NEXT:       - .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:         .value_type:     i64
 ; CHECK-NEXT:       - .offset:         48
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:         .value_type:     i64
 ; CHECK-NEXT:       - .offset:         56
 ; 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:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         80
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         88
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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,

 
 ; 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:  8
 ; 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:  32
 ; 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:  64
 ; 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:  256
 ; 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:  128
 ; 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:  1024
 ; 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:       - .address_space:  local
 ; CHECK-NEXT:         .name:           h
 ; CHECK-NEXT:         .offset:         32
 ; CHECK-NEXT:         .pointee_align:  16
 ; CHECK-NEXT:         .size:           4
 ; CHECK-NEXT:         .value_kind:     dynamic_shared_pointer
-; CHECK-NEXT:         .value_type:     struct
 ; CHECK-NEXT:       - .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:         .value_type:     i64
 ; CHECK-NEXT:       - .offset:         48
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:         .value_type:     i64
 ; CHECK-NEXT:       - .offset:         56
 ; 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:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_printf_buffer
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         80
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         88
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK:          .language:       OpenCL C
 ; CHECK-NEXT:     .language_version:
 ; CHECK-NEXT:       - 2
 ; CHECK-NEXT:       - 0
 ; CHECK:          .name:           test_pointee_align_attribute
 ; CHECK:          .symbol:         test_pointee_align_attribute.kd
 define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a,
                                                         i8 addrspace(3)* align 8 %b,
                                                         <2 x i8> addrspace(3)* align 32 %c,
                                                         <3 x i8> addrspace(3)* align 64 %d,
                                                         <4 x i8> addrspace(3)* align 256 %e,
                                                         <8 x i8> addrspace(3)* align 128 %f,
                                                         <16 x i8> addrspace(3)* align 1024 %g,
                                                         {} addrspace(3)* align 16 %h) #0
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         80
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #1
     !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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         56
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; 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) #2
     !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:        - .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.printf:
 ; CHECK-NEXT: - '1:1:4:%d\n'

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

 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NOT:        ValueKind:     HiddenHostcallBuffer
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
 ; CHECK-NOT:        ValueKind:     HiddenCompletionAction
 ; CHECK:            ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_char(i8 %a) #0
     !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:      - Name:            test_ushort2
 ; CHECK-NEXT:   SymbolName:      'test_ushort2@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      ushort2
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     U16
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0
     !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:      - Name:            test_int3
 ; CHECK-NEXT:   SymbolName:      'test_int3@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int3
 ; CHECK-NEXT:       Size:          16
 ; CHECK-NEXT:       Align:         16
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_int3(<3 x i32> %a) #0
     !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:      - Name:            test_ulong4
 ; CHECK-NEXT:   SymbolName:      'test_ulong4@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      ulong4
 ; CHECK-NEXT:       Size:          32
 ; CHECK-NEXT:       Align:         32
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     U64
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0
     !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:      - Name:            test_half8
 ; CHECK-NEXT:   SymbolName:      'test_half8@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      half8
 ; CHECK-NEXT:       Size:          16
 ; CHECK-NEXT:       Align:         16
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     F16
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_half8(<8 x half> %a) #0
     !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:      - Name:            test_float16
 ; CHECK-NEXT:   SymbolName:      'test_float16@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      float16
 ; CHECK-NEXT:       Size:          64
 ; CHECK-NEXT:       Align:         64
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     F32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_float16(<16 x float> %a) #0
     !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:      - Name:            test_double16
 ; CHECK-NEXT:   SymbolName:      'test_double16@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      double16
 ; CHECK-NEXT:       Size:          128
 ; CHECK-NEXT:       Align:         128
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     F64
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_double16(<16 x double> %a) #0
     !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:      - Name:            test_pointer
 ; CHECK-NEXT:   SymbolName:      'test_pointer@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a) #0
     !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:      - Name:            test_image
 ; CHECK-NEXT:   SymbolName:      'test_image@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      image2d_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Image
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0
     !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:      - Name:            test_sampler
 ; CHECK-NEXT:   SymbolName:      'test_sampler@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      sampler_t
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     Sampler
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_sampler(i32 %a) #0
     !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:      - Name:            test_queue
 ; CHECK-NEXT:   SymbolName:      'test_queue@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      queue_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Queue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0
     !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:      - Name:            test_struct
 ; CHECK-NEXT:   SymbolName:      'test_struct@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      struct A
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_struct(%struct.A %a) #0
     !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:      - Name:            test_array
 ; CHECK-NEXT:   SymbolName:      'test_array@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      struct A
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_array([8 x i8] %a) #0
     !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:      - Name:            test_i128
 ; CHECK-NEXT:   SymbolName:      'test_i128@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      i128
 ; CHECK-NEXT:       Size:          16
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_i128(i128 %a) #0
     !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:      - Name:            test_multi_arg
 ; CHECK-NEXT:   SymbolName:      'test_multi_arg@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          b
 ; CHECK-NEXT:       TypeName:      short2
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I16
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          c
 ; CHECK-NEXT:       TypeName:      char3
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0
     !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:      - Name:            test_addr_space
 ; CHECK-NEXT:   SymbolName:      'test_addr_space@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          g
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          c
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Constant
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          l
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       PointeeAlign:  4
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
                                            i32 addrspace(4)* %c,
                                            i32 addrspace(3)* %l) #0
     !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:      - Name:            test_type_qual
 ; CHECK-NEXT:   SymbolName:      'test_type_qual@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:       IsVolatile:    true
 ; CHECK-NEXT:     - Name:          b
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:       IsConst:       true
 ; CHECK-NEXT:       IsRestrict:    true
 ; CHECK-NEXT:     - Name:          c
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Pipe
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:       IsPipe:        true
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a,
                                           i32 addrspace(1)* %b,
                                           %opencl.pipe_t addrspace(1)* %c) #0
     !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:      - Name:            test_access_qual
 ; CHECK-NEXT:   SymbolName:      'test_access_qual@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          ro
 ; CHECK-NEXT:       TypeName:      image1d_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Image
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       ReadOnly
 ; CHECK-NEXT:     - Name:          wo
 ; CHECK-NEXT:       TypeName:      image2d_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Image
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       WriteOnly
 ; CHECK-NEXT:     - Name:          rw
 ; CHECK-NEXT:       TypeName:      image3d_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     Image
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       ReadWrite
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 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) #0
     !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:      - Name:            test_vec_type_hint_half
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_half@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   half
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_half(i32 %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 !vec_type_hint !26 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_float
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_float@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   float
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_float(i32 %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 !vec_type_hint !27 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_double
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_double@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   double
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_double(i32 %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 !vec_type_hint !28 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_char
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_char@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   char
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_char(i32 %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 !vec_type_hint !29 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_short
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_short@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   short
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_short(i32 %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 !vec_type_hint !30 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_long
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_long@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   long
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_long(i32 %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 !vec_type_hint !31 {
   ret void
 }
 
 ; CHECK:      - Name:            test_vec_type_hint_unknown
 ; CHECK-NEXT:   SymbolName:      'test_vec_type_hint_unknown@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       VecTypeHint:   unknown
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      int
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %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 !vec_type_hint !32 {
   ret void
 }
 
 ; CHECK:      - Name:            test_reqd_wgs_vec_type_hint
 ; CHECK-NEXT:   SymbolName:      'test_reqd_wgs_vec_type_hint@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       ReqdWorkGroupSize: [ 1, 2, 4 ]
 ; CHECK-NEXT:       VecTypeHint:       int
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:              a
 ; CHECK-NEXT:       TypeName:          int
 ; CHECK-NEXT:       Size:              4
 ; CHECK-NEXT:       Align:             4
 ; CHECK-NEXT:       ValueKind:         ByValue
-; CHECK-NEXT:       ValueType:         I32
 ; CHECK-NEXT:       AccQual:           Default
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenNone
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenNone
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %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 !vec_type_hint !5
     !reqd_work_group_size !6 {
   ret void
 }
 
 ; CHECK:      - Name:            test_wgs_hint_vec_type_hint
 ; CHECK-NEXT:   SymbolName:      'test_wgs_hint_vec_type_hint@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       WorkGroupSizeHint: [ 8, 16, 32 ]
 ; CHECK-NEXT:       VecTypeHint:       uint4
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:              a
 ; CHECK-NEXT:       TypeName:          int
 ; CHECK-NEXT:       Size:              4
 ; CHECK-NEXT:       Align:             4
 ; CHECK-NEXT:       ValueKind:         ByValue
-; CHECK-NEXT:       ValueType:         I32
 ; CHECK-NEXT:       AccQual:           Default
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:         I64
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenNone
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenNone
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 ; CHECK-NEXT:     - Size:              8
 ; CHECK-NEXT:       Align:             8
 ; CHECK-NEXT:       ValueKind:         HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:         I8
 ; CHECK-NEXT:       AddrSpaceQual:     Global
 define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %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 !vec_type_hint !7
     !work_group_size_hint !8 {
   ret void
 }
 
 ; CHECK:      - Name:            test_arg_ptr_to_ptr
 ; CHECK-NEXT:   SymbolName:      'test_arg_ptr_to_ptr@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'int  addrspace(5)* addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a) #0
     !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:      - Name:            test_arg_struct_contains_ptr
 ; CHECK-NEXT:   SymbolName:      'test_arg_struct_contains_ptr@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      struct B
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B %a) #0
     !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:      - Name:            test_arg_vector_of_ptr
 ; CHECK-NEXT:   SymbolName:      'test_arg_vector_of_ptr@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'global int addrspace(5)* __attribute__((ext_vector_type(2)))'
 ; CHECK-NEXT:       Size:          16
 ; CHECK-NEXT:       Align:         16
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I32
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a) #0
     !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:      - Name:            test_arg_unknown_builtin_type
 ; CHECK-NEXT:   SymbolName:      'test_arg_unknown_builtin_type@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      clk_event_t
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_arg_unknown_builtin_type(
     %opencl.clk_event_t addrspace(1)* %a) #0
     !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:      - Name:            test_pointee_align
 ; CHECK-NEXT:   SymbolName:      'test_pointee_align@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'long  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          b
 ; CHECK-NEXT:       TypeName:      'char  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  1
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          c
 ; CHECK-NEXT:       TypeName:      'char2  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  2
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          d
 ; CHECK-NEXT:       TypeName:      'char3  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  4
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          e
 ; CHECK-NEXT:       TypeName:      'char4  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  4
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          f
 ; CHECK-NEXT:       TypeName:      'char8  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  8
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          g
 ; CHECK-NEXT:       TypeName:      'char16  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  16
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          h
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       PointeeAlign:  1
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 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,
                                               {} addrspace(3)* %h) #0
     !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:      - Name:            test_pointee_align_attribute
 ; CHECK-NEXT:   SymbolName:      'test_pointee_align_attribute@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      'long  addrspace(5)*'
 ; CHECK-NEXT:       Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     GlobalBuffer
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          b
 ; CHECK-NEXT:       TypeName:      'char  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  8
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          c
 ; CHECK-NEXT:       TypeName:      'char2  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  32
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          d
 ; CHECK-NEXT:       TypeName:      'char3  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  64
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          e
 ; CHECK-NEXT:       TypeName:      'char4  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  256
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          f
 ; CHECK-NEXT:       TypeName:      'char8  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  128
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Name:          g
 ; CHECK-NEXT:       TypeName:      'char16  addrspace(5)*'
 ; CHECK-NEXT:       Size:          4
 ; CHECK-NEXT:       Align:         4
 ; CHECK-NEXT:       ValueKind:     DynamicSharedPointer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       PointeeAlign:  1024
 ; CHECK-NEXT:       AddrSpaceQual: Local
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:      - Name:            h
 ; CHECK-NEXT:        Size:            4
 ; CHECK-NEXT:        Align:           4
 ; CHECK-NEXT:        ValueKind:       DynamicSharedPointer
-; CHECK-NEXT:        ValueType:       Struct
 ; CHECK-NEXT:        PointeeAlign:    16
 ; CHECK-NEXT:        AddrSpaceQual:   Local
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a,
                                                         i8 addrspace(3)* align 8 %b,
                                                         <2 x i8> addrspace(3)* align 32 %c,
                                                         <3 x i8> addrspace(3)* align 64 %d,
                                                         <4 x i8> addrspace(3)* align 256 %e,
                                                         <8 x i8> addrspace(3)* align 128 %f,
                                                         <16 x i8> addrspace(3)* align 1024 %g,

 ; CHECK-NEXT:   Attrs:
 ; CHECK-NEXT:       RuntimeHandle: __test_block_invoke_kernel_runtime_handle
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          arg
 ; CHECK-NEXT:       TypeName:      __block_literal
 ; CHECK-NEXT:       Size:          25
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     Struct
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenNone
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @__test_block_invoke_kernel(
     <{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #1
     !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:      - Name:            test_enqueue_kernel_caller
 ; CHECK-NEXT:   SymbolName:      'test_enqueue_kernel_caller@kd'
 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenPrintfBuffer
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenDefaultQueue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenCompletionAction
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AddrSpaceQual: Global
 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2
     !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: - Name:            unknown_addrspace_kernarg
 ; CHECK: Args:
 ; CHECK-NEXT: - Name:            ptr
 ; CHECK-NEXT: Size:            8
 ; CHECK-NEXT: Align:           8
 ; CHECK-NEXT: ValueKind:       GlobalBuffer
-; CHECK-NEXT: ValueType:       I32
 define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
   ret void
 }
 
 attributes #0 = { "amdgpu-implicitarg-num-bytes"="56" }
 attributes #1 = { "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
 attributes #2 = { "amdgpu-implicitarg-num-bytes"="56" "calls-enqueue-kernel" }
 

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

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -mattr=+code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX700 --check-prefix=NOTES %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -mattr=+code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX803 --check-prefix=NOTES %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=+code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 
 ; 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:          .name:           test0
 ; CHECK:          .symbol:         test0.kd
 define amdgpu_kernel void @test0(
     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:        - .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:          .name:           test8
 ; CHECK:          .symbol:         test8.kd
 define amdgpu_kernel void @test8(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #0 {
 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:        - .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:          .name:           test16
 ; CHECK:          .symbol:         test16.kd
 define amdgpu_kernel void @test16(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #1 {
 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:        - .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:          .name:           test24
 ; CHECK:          .symbol:         test24.kd
 define amdgpu_kernel void @test24(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #2 {
 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:        - .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:          .name:           test32
 ; CHECK:          .symbol:         test32.kd
 define amdgpu_kernel void @test32(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #3 {
 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:        - .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:          .name:           test48
 ; CHECK:          .symbol:         test48.kd
 define amdgpu_kernel void @test48(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #4 {
 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:        - .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-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         72
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_multigrid_sync_arg
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK:          .name:           test56
 ; CHECK:          .symbol:         test56.kd
 define amdgpu_kernel void @test56(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #5 {
 entry:
   %a.val = load half, half addrspace(1)* %a

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

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -mattr=-code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX700 --check-prefix=NOTES %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 -mattr=-code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX803 --check-prefix=NOTES %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-code-object-v3 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck --check-prefix=CHECK --check-prefix=GFX900 --check-prefix=NOTES %s
 
 ; CHECK: ---
 ; CHECK:  Version: [ 1, 0 ]
 ; CHECK:  Kernels:
 
 ; CHECK:      - Name:       test0
 ; CHECK:        SymbolName: 'test0@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test0(
     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:      - Name:       test8
 ; CHECK:        SymbolName: 'test8@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test8(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #0 {
 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:      - Name:       test16
 ; CHECK:        SymbolName: 'test16@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test16(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #1 {
 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:      - Name:       test24
 ; CHECK:        SymbolName: 'test24@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test24(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #2 {
 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:      - Name:       test32
 ; CHECK:        SymbolName: 'test32@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test32(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #3 {
 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:      - Name:       test48
 ; CHECK:        SymbolName: 'test48@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test48(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #4 {
 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:      - Name:       test56
 ; CHECK:        SymbolName: 'test56@kd'
 ; CHECK:        Args:
 ; CHECK-NEXT:     - Name:            r
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            a
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Name:            b
 ; CHECK-NEXT:       Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       GlobalBuffer
-; CHECK-NEXT:       ValueType:       F16
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:       I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenMultiGridSyncArg
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test56(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #5 {
 entry:
   %a.val = load half, half addrspace(1)* %a

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-absent-v3.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -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:        - .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_hostcall_buffer
 
 ; CHECK:          .language:       OpenCL C
 ; CHECK-NEXT:     .language_version:
 ; CHECK-NEXT:       - 2
 ; CHECK-NEXT:       - 0
 ; CHECK:          .name:           test_kernel

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-absent.ll

 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NOT:        ValueKind:     HiddenHostcallBuffer
 ; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
 ; CHECK-NOT:        ValueKind:     HiddenCompletionAction
 
 define amdgpu_kernel void @test_kernel(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

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present-v3.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes | FileCheck %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mattr=+code-object-v3 -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:        - .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_hostcall_buffer
-; CHECK-NEXT:         .value_type:     i8
 ; CHECK:          .language:       OpenCL C
 ; CHECK-NEXT:     .language_version:
 ; CHECK-NEXT:       - 2
 ; CHECK-NEXT:       - 0
 ; CHECK:          .name:           test_kernel
 ; CHECK:          .symbol:         test_kernel.kd
 
 declare <2 x i64> @__ockl_hostcall_internal(i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64)

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present.ll

 ; CHECK-NEXT:   Language:        OpenCL C
 ; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
 ; CHECK-NEXT:   Args:
 ; CHECK-NEXT:     - Name:          a
 ; CHECK-NEXT:       TypeName:      char
 ; CHECK-NEXT:       Size:          1
 ; CHECK-NEXT:       Align:         1
 ; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       ValueType:     I8
 ; CHECK-NEXT:       AccQual:       Default
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:       ValueType:     I64
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenHostcallBuffer
-; CHECK-NEXT:       ValueType:       I8
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
 ; CHECK-NOT:        ValueKind:     HiddenCompletionAction
 
 declare <2 x i64> @__ockl_hostcall_internal(i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64)
 
 define amdgpu_kernel void @test_kernel(i8 %a) #0
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3

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

 %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:   - .args:           
+; 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,

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

 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx700 -mattr=-code-object-v3 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX700 %s
 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx800 -mattr=-code-object-v3 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX800 %s
 // RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-code-object-v3 -show-encoding %s | FileCheck --check-prefix=CHECK --check-prefix=GFX900 %s
 
+// The legacy ValueType field should be parsed without error, but not
+// re-emitted.
+
 // CHECK:  .amd_amdgpu_hsa_metadata
 // CHECK:    Version: [ 1, 0 ]
 // CHECK:    Printf:
 // CHECK:      - '1:1:4:%d\n'
 // CHECK:      - '2:1:8:%g\n'
 // CHECK:    Kernels:
 // CHECK:      - Name:            test_kernel
 // CHECK:        SymbolName:      'test_kernel@kd'
 // CHECK:        Language:        OpenCL C
 // CHECK:        LanguageVersion: [ 2, 0 ]
 // CHECK:        Args:
 // CHECK:          - TypeName:      char
 // CHECK:            Size:          1
 // CHECK:            Align:         1
 // CHECK:            ValueKind:     ByValue
-// CHECK:            ValueType:     I8
 // CHECK:            AccQual:       Default
 // CHECK:          - Size:          8
 // CHECK:            Align:         8
 // CHECK:            ValueKind:     HiddenGlobalOffsetX
-// CHECK:            ValueType:     I64
 // CHECK:          - Size:          8
 // CHECK:            Align:         8
 // CHECK:            ValueKind:     HiddenGlobalOffsetY
-// CHECK:            ValueType:     I64
 // CHECK:          - Size:          8
 // CHECK:            Align:         8
 // CHECK:            ValueKind:     HiddenGlobalOffsetZ
-// CHECK:            ValueType:     I64
 // CHECK:          - Size:          8
 // CHECK:            Align:         8
 // CHECK:            ValueKind:     HiddenPrintfBuffer
-// CHECK:            ValueType:     I8
 // CHECK:            AddrSpaceQual: Global
 // CHECK:  .end_amd_amdgpu_hsa_metadata
 .amd_amdgpu_hsa_metadata
   Version: [ 1, 0 ]
   Printf: [ '1:1:4:%d\n', '2:1:8:%g\n' ]
   Kernels:
     - Name:            test_kernel
       SymbolName:      test_kernel@kd