AMDGPU/llvm-readobj: Add missing tests for note parsing/displaying

[kzhuravl]

Summary

This is a follow up review/change for https://reviews.llvm.org/D95638

Add valid note tests for code object v2 notes:

• NT_AMD_HSA_CODE_OBJECT_VERSION (required yaml2obj update)
• NT_AMD_HSA_HSAIL (required yaml2obj update)
• NT_AMD_HSA_ISA_VERSION (required yaml2obj update)
• NT_AMD_HSA_METADATA
• NT_AMD_HSA_ISA_NAME
• NT_AMD_PAL_METADATA

Add valid note tests for code object v3 notes:

• NT_AMDGPU_METADATA

Add invalid note tests for code object v2 notes:

• NT_AMD_HSA_CODE_OBJECT_VERSION (required yaml2obj update)
• NT_AMD_HSA_HSAIL (required yaml2obj update)
• NT_AMD_HSA_ISA_VERSION (required yaml2obj update)

Add invalid note tests for code object v3 notes:

• NT_AMDGPU_METADATA

[kzhuravl]

I believe this testing covers all cases (both valid and invalid) of note parsing/displaying for llvm-readobj

[jhenderson]

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v2.s

1

Please add a brief top-level comment to this test explaining the test's purpose.

It would also be helpful to have some inline comments next to each test case (possibly by the YAML) explaining why each case is invalid for future reference.

Same comments apply to each of the new tests.

Also, these should probably have a test file name ending in .test, since they aren't written in assembly.

121

It might be a good idea to have one that's too big as well as too small, to show that the size has to match exactly.

Same goes for the other cases where the desc size is incorrect.

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v3.s

30

Is this the minimum the data can be in order to trigger the code path you are testing here?

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.s

2–3

Here and elsewhere, you might want something like --match-full-lines for the FileCheck invocation. That will ensure no garbage is printed after the expected string on any given line of output, which is especially important for where strings are being printed (in case something is wrong regarding null termination in the code).

165

If I'm reading this correctly, this is just a string, in which case it can be any string, e.g. "foo" rather than some very large binary blob. I recommend simplifying drastically - it doesn't need to be a real-world case for testing purposes. Also, I believe you need a case for this note type where Desc.size() == 0. That may be tested elsewhere, but it would make sense to explicitly check it in llvm-readobj too.

Same goes for the NT_AMD_HSA_ISA_NAME case.

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s

225

Such a large (and opaque) blob might mean it's easier to write this test in assembly and use llvm-mc to generate the output. The same goes for the invalid version of this case.

llvm/tools/llvm-readobj/ELFDumper.cpp

5058

Nit: run clang-format.

5114

I missed this before, but what happens if Desc isn't a multiple of sizeof(PALMetadata)? It looks to me like you'd read off the end of the note data, which should trigger some sort of error (but currently isn't checked). Please fix and add a test case.

[jhenderson]

Thanks for the tests and fixes. Hopefully my comments make sense to you!

[kzhuravl]

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v3.s

30

I have trimmed it to the minimum.

[kzhuravl]

Address review feedback.

[jhenderson]

Just one nit, otherwise looks good, assuming all the tests are passing etc.

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v2.test

1–2 ↗

(On Diff #340893)

Most new llvm-readobj tests use ## for comments, to help them stand out from lit and FileCheck directives. Same goes with your other comments in this and other tests.

[kzhuravl]

Address review feedback.

[jhenderson]

LGTM, once the clang-format issue has been addressed. Note: If you know anybody who knows AMDGPU more than me (i.e. at all), it might not hurt to get them to give the tests to give a once over before landing this, but don't worry if you don't.

llvm/tools/llvm-readobj/ELFDumper.cpp

5058

Ping: clang-format appears to be complaining still here.

[kzhuravl]

clang-format.

[kzhuravl]

+ @t-tye to maybe look at note tests...

[t-tye]

LGTM

[kzhuravl]

@jhenderson , do you have any additional comments?

[jhenderson]

In D101304#2729044, @kzhuravl wrote:

@jhenderson , do you have any additional comments?

Nope!

[stella.stamenova]

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s

5

This test assumes that the triple is available, but that's not always the case which results in failures such as:

FAIL: LLVM :: tools/llvm-readobj/ELF/note-amd-valid-v3.s (66018 of 76097)
******************** TEST 'LLVM :: tools/llvm-readobj/ELF/note-amd-valid-v3.s' FAILED ********************
Script:
--
: 'RUN: at line 4';   /mnt/vss/_work/1/b/llvm/bin/llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj < /mnt/vss/_work/1/s/llvm-project/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s | /mnt/vss/_work/1/b/llvm/bin/llvm-readobj --notes - | /mnt/vss/_work/1/b/llvm/bin/FileCheck /mnt/vss/_work/1/s/llvm-project/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s --match-full-lines --check-prefix=LLVM
: 'RUN: at line 5';   /mnt/vss/_work/1/b/llvm/bin/llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj < /mnt/vss/_work/1/s/llvm-project/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s | /mnt/vss/_work/1/b/llvm/bin/llvm-readelf --notes - | /mnt/vss/_work/1/b/llvm/bin/FileCheck /mnt/vss/_work/1/s/llvm-project/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s --match-full-lines --check-prefix=GNU
--
Exit Code: 2

Command Output (stderr):
--
/mnt/vss/_work/1/b/llvm/bin/llvm-mc: error: : error: unable to get target for 'amdgcn-amd-amdhsa', see --version and --triple.
/mnt/vss/_work/1/b/llvm/bin/llvm-readobj: error: '<stdin>': The file was not recognized as a valid object file
FileCheck error: '<stdin>' is empty.
FileCheck command line:  /mnt/vss/_work/1/b/llvm/bin/FileCheck /mnt/vss/_work/1/s/llvm-project/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s --match-full-lines --check-prefix=LLVM

You either need to use REQUIRES: to specify what is needed or you need to move it to one of the directories that already do (such as AMDGPU).

[nickdesaulniers]

Hi, llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s is failing for me.
/android0/llvm-project/llvm/build/bin/llvm-mc: error: : error: unable to get target for 'amdgcn-amd-amdhsa', see --version and --triple.
/android0/llvm-project/llvm/build/bin/llvm-readobj: error: '<stdin>': The file was not recognized as a valid object file
FileCheck error: '<stdin>' is empty.

[kzhuravl]

Thanks, I am working on a fix.

[kzhuravl]

I have reverted the change while I am testing the fix.

[Conanap]

Hello,

The reland version of this patch for which I cannot find the differential of has broken LLVM::note-amd-valid-v2.test on PowerPC:

/home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test:17:14: error: LLVM-NEXT: expected string not found in input
# LLVM-NEXT: AMD HSA Code Object Version: [Major: 2, Minor: 1]
             ^
<stdin>:15:68: note: scanning from here
 Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
                                                                   ^
<stdin>:16:2: note: possible intended match here
 AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216]
 ^
Input file: <stdin>
Check file: /home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
-dump-input=help explains the following input dump.
Input was:
<<<<<<
           .
           .
           .
          10:  Offset: 0x40 
          11:  Size: 0x18 
          12:  Note { 
          13:  Owner: AMD 
          14:  Data size: 0x8 
          15:  Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version) 
next:17'0                                                                        X error: no match found
          16:  AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216] 
next:17'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:17'1      ?                                                                possible intended match

Seems like a small fix but if you could have a push for it soon that'd be great. Thanks! FYI @rksharma

[kzhuravl]

In D101304#2733289, @Conanap wrote:

Hello,

The reland version of this patch for which I cannot find the differential of has broken LLVM::note-amd-valid-v2.test on PowerPC:

/home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test:17:14: error: LLVM-NEXT: expected string not found in input
# LLVM-NEXT: AMD HSA Code Object Version: [Major: 2, Minor: 1]
             ^
<stdin>:15:68: note: scanning from here
 Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
                                                                   ^
<stdin>:16:2: note: possible intended match here
 AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216]
 ^
Input file: <stdin>
Check file: /home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
-dump-input=help explains the following input dump.
Input was:
<<<<<<
           .
           .
           .
          10:  Offset: 0x40 
          11:  Size: 0x18 
          12:  Note { 
          13:  Owner: AMD 
          14:  Data size: 0x8 
          15:  Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version) 
next:17'0                                                                        X error: no match found
          16:  AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216] 
next:17'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:17'1      ?                                                                possible intended match

Seems like a small fix but if you could have a push for it soon that'd be great. Thanks! FYI @rksharma

Hi,

Reland patch fixed an issue if AMDGPU backend was not built, which was a one liner (I reverted original patch to not create the noise).

Seems like it is an endianness issue.

Unfortunately I don't have PowerPC, will adding these lines in note-amd-valid-v2.test mark it as XFAIL for you:

diff --git a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
index 688bbd50994c..f9cac212ce9f 100644
--- a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
+++ b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
@@ -1,6 +1,9 @@
 ## This test is checking the handling of valid note entries for AMDGPU code
 ## object v2.

+## Big endian not supported.
+# XFAIL: host-byteorder-big-endian
+
 # RUN: yaml2obj %s -o %t.o
 # RUN: llvm-readobj --notes %t.o | FileCheck %s --match-full-lines --check-prefix=LLVM
 # RUN: llvm-readelf --notes %t.o | FileCheck %s --match-full-lines --check-prefix=GNU

?

[Conanap]

In D101304#2733314, @kzhuravl wrote:

In D101304#2733289, @Conanap wrote:

Hello,

The reland version of this patch for which I cannot find the differential of has broken LLVM::note-amd-valid-v2.test on PowerPC:

/home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test:17:14: error: LLVM-NEXT: expected string not found in input
# LLVM-NEXT: AMD HSA Code Object Version: [Major: 2, Minor: 1]
             ^
<stdin>:15:68: note: scanning from here
 Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
                                                                   ^
<stdin>:16:2: note: possible intended match here
 AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216]
 ^
Input file: <stdin>
Check file: /home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
-dump-input=help explains the following input dump.
Input was:
<<<<<<
           .
           .
           .
          10:  Offset: 0x40 
          11:  Size: 0x18 
          12:  Note { 
          13:  Owner: AMD 
          14:  Data size: 0x8 
          15:  Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version) 
next:17'0                                                                        X error: no match found
          16:  AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216] 
next:17'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:17'1      ?                                                                possible intended match

Seems like a small fix but if you could have a push for it soon that'd be great. Thanks! FYI @rksharma

Hi,

Reland patch fixed an issue if AMDGPU backend was not built, which was a one liner (I reverted original patch to not create the noise).

Seems like it is an endianness issue.

Unfortunately I don't have PowerPC, will adding these lines in note-amd-valid-v2.test mark it as XFAIL for you:

diff --git a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
index 688bbd50994c..f9cac212ce9f 100644
--- a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
+++ b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
@@ -1,6 +1,9 @@
 ## This test is checking the handling of valid note entries for AMDGPU code
 ## object v2.

+## Big endian not supported.
+# XFAIL: host-byteorder-big-endian
+
 # RUN: yaml2obj %s -o %t.o
 # RUN: llvm-readobj --notes %t.o | FileCheck %s --match-full-lines --check-prefix=LLVM
 # RUN: llvm-readelf --notes %t.o | FileCheck %s --match-full-lines --check-prefix=GNU

?

That would fix it yes, thank you!
Another possible solution would be to add a BE specific runline, but I'm not really familiar with the AMD test files so it's just a suggestion.

[kzhuravl]

In D101304#2733333, @Conanap wrote:

In D101304#2733314, @kzhuravl wrote:

In D101304#2733289, @Conanap wrote:

Hello,

The reland version of this patch for which I cannot find the differential of has broken LLVM::note-amd-valid-v2.test on PowerPC:

/home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test:17:14: error: LLVM-NEXT: expected string not found in input
# LLVM-NEXT: AMD HSA Code Object Version: [Major: 2, Minor: 1]
             ^
<stdin>:15:68: note: scanning from here
 Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
                                                                   ^
<stdin>:16:2: note: possible intended match here
 AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216]
 ^
Input file: <stdin>
Check file: /home/buildbots/ppc64be-clang-test/clang-ppc64be/llvm/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
-dump-input=help explains the following input dump.
Input was:
<<<<<<
           .
           .
           .
          10:  Offset: 0x40 
          11:  Size: 0x18 
          12:  Note { 
          13:  Owner: AMD 
          14:  Data size: 0x8 
          15:  Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version) 
next:17'0                                                                        X error: no match found
          16:  AMD HSA Code Object Version: [Major: 33554432, Minor: 16777216] 
next:17'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:17'1      ?                                                                possible intended match

Seems like a small fix but if you could have a push for it soon that'd be great. Thanks! FYI @rksharma

Hi,

Reland patch fixed an issue if AMDGPU backend was not built, which was a one liner (I reverted original patch to not create the noise).

Seems like it is an endianness issue.

Unfortunately I don't have PowerPC, will adding these lines in note-amd-valid-v2.test mark it as XFAIL for you:

diff --git a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
index 688bbd50994c..f9cac212ce9f 100644
--- a/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
+++ b/llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.test
@@ -1,6 +1,9 @@
 ## This test is checking the handling of valid note entries for AMDGPU code
 ## object v2.

+## Big endian not supported.
+# XFAIL: host-byteorder-big-endian
+
 # RUN: yaml2obj %s -o %t.o
 # RUN: llvm-readobj --notes %t.o | FileCheck %s --match-full-lines --check-prefix=LLVM
 # RUN: llvm-readelf --notes %t.o | FileCheck %s --match-full-lines --check-prefix=GNU

?

That would fix it yes, thank you!
Another possible solution would be to add a BE specific runline, but I'm not really familiar with the AMD test files so it's just a suggestion.

Thanks, I have pushed that code snippet for now to make build bots happy. Let me look into adding big endian specific run lines today...

[kzhuravl]

In D101304#2729931, @nickdesaulniers wrote:

Hi, llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s is failing for me.
/android0/llvm-project/llvm/build/bin/llvm-mc: error: : error: unable to get target for 'amdgcn-amd-amdhsa', see --version and --triple.
/android0/llvm-project/llvm/build/bin/llvm-readobj: error: '<stdin>': The file was not recognized as a valid object file
FileCheck error: '<stdin>' is empty.

Thanks, I have re-landed original patch + requires amdgpu-registered-target

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s

5

Thanks, I have re-landed original patch + requires amdgpu-registered-target

llvm/lib/ObjectYAML/ELFYAML.cpp

#define ECase(X) IO.enumCase(Value, #X, ELF::X)
   ECase(NT_FREEBSD_PROCSTAT_VMMAP);
   ECase(NT_FREEBSD_PROCSTAT_GROUPS);
   ECase(NT_FREEBSD_PROCSTAT_UMASK);
   ECase(NT_FREEBSD_PROCSTAT_RLIMIT);
   ECase(NT_FREEBSD_PROCSTAT_OSREL);
   ECase(NT_FREEBSD_PROCSTAT_PSSTRINGS);
   ECase(NT_FREEBSD_PROCSTAT_AUXV);
   // AMD specific notes. (Code Object V2)
+  ECase(NT_AMD_HSA_CODE_OBJECT_VERSION);
+  ECase(NT_AMD_HSA_HSAIL);
+  ECase(NT_AMD_HSA_ISA_VERSION);
   ECase(NT_AMD_HSA_METADATA);
   ECase(NT_AMD_HSA_ISA_NAME);
   ECase(NT_AMD_PAL_METADATA);
   // AMDGPU specific notes. (Code Object V3)
   ECase(NT_AMDGPU_METADATA);
 #undef ECase
   IO.enumFallback<Hex32>(Value);
 }

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v2.s

+# RUN: yaml2obj %s -o %t.o
+# RUN: llvm-readobj --notes %t.o | FileCheck %s --check-prefix=LLVM
+# RUN: llvm-readelf --notes %t.o | FileCheck %s --check-prefix=GNU
+
+# LLVM:      Notes [
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_code_object_version
+# LLVM-NEXT:     Offset: 0x40
+# LLVM-NEXT:     Size: 0x14
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x4
+# LLVM-NEXT:       Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
+# LLVM-NEXT:       AMD HSA Code Object Version: Invalid AMD HSA Code Object Version
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_hsail
+# LLVM-NEXT:     Offset: 0x54
+# LLVM-NEXT:     Size: 0x1C
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0xA
+# LLVM-NEXT:       Type: NT_AMD_HSA_HSAIL (AMD HSA HSAIL Properties)
+# LLVM-NEXT:       AMD HSA HSAIL Properties: Invalid AMD HSA HSAIL Properties
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_version_0
+# LLVM-NEXT:     Offset: 0x70
+# LLVM-NEXT:     Size: 0x18
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x8
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# LLVM-NEXT:       AMD HSA ISA Version: Invalid AMD HSA ISA Version
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_version_1
+# LLVM-NEXT:     Offset: 0x88
+# LLVM-NEXT:     Size: 0x28
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x17
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# LLVM-NEXT:       AMD HSA ISA Version: Invalid AMD HSA ISA Version
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_version_2
+# LLVM-NEXT:     Offset: 0xB0
+# LLVM-NEXT:     Size: 0x28
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x17
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# LLVM-NEXT:       AMD HSA ISA Version: Invalid AMD HSA ISA Version
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_version_3
+# LLVM-NEXT:     Offset: 0xD8
+# LLVM-NEXT:     Size: 0x28
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x17
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# LLVM-NEXT:       AMD HSA ISA Version: Invalid AMD HSA ISA Version
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT: ]
+
+# GNU:      Displaying notes found in: .note.nt_amd_hsa_code_object_version
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x00000004       NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
+# GNU-NEXT:     AMD HSA Code Object Version:
+# GNU-NEXT:         Invalid AMD HSA Code Object Version
+# GNU-EMPTY:
+# GNU-NEXT: Displaying notes found in: .note.nt_amd_hsa_hsail
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x0000000a       NT_AMD_HSA_HSAIL (AMD HSA HSAIL Properties)
+# GNU-NEXT:     AMD HSA HSAIL Properties:
+# GNU-NEXT:         Invalid AMD HSA HSAIL Properties
+# GNU-EMPTY:
+# GNU-NEXT: Displaying notes found in: .note.nt_amd_hsa_isa_version_0
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x00000008       NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# GNU-NEXT:     AMD HSA ISA Version:
+# GNU-NEXT:         Invalid AMD HSA ISA Version
+# GNU-EMPTY:
+# GNU-NEXT: Displaying notes found in: .note.nt_amd_hsa_isa_version_1
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x00000017       NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# GNU-NEXT:     AMD HSA ISA Version:
+# GNU-NEXT:         Invalid AMD HSA ISA Version
+# GNU-EMPTY:
+# GNU-NEXT: Displaying notes found in: .note.nt_amd_hsa_isa_version_2
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x00000017       NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# GNU-NEXT:     AMD HSA ISA Version:
+# GNU-NEXT:         Invalid AMD HSA ISA Version
+# GNU-EMPTY:
+# GNU-NEXT: Displaying notes found in: .note.nt_amd_hsa_isa_version_3
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMD                  0x00000017       NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# GNU-NEXT:     AMD HSA ISA Version:
+# GNU-NEXT:         Invalid AMD HSA ISA Version
+
+--- !ELF
+FileHeader:
+  Class: ELFCLASS64
+  Data:  ELFDATA2LSB
+  Type:  ET_REL
+Sections:
+  - Name:     .note.nt_amd_hsa_code_object_version
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_CODE_OBJECT_VERSION
+        Desc: '02000000'
+  - Name:     .note.nt_amd_hsa_hsail
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_HSAIL
+        Desc: '02000000010000000102'
+  - Name:     .note.nt_amd_hsa_isa_version_0
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_VERSION
+        Desc: '0400070008000000'
+  - Name:     .note.nt_amd_hsa_isa_version_1
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_VERSION
+        Desc: '04000700080000000000000002000000414d4400414d44'
+  - Name:     .note.nt_amd_hsa_isa_version_2
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_VERSION
+        Desc: '00000700080000000000000002000000414d4400414d44'
+  - Name:     .note.nt_amd_hsa_isa_version_3
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_VERSION
+        Desc: '04000000080000000000000002000000414d4400414d44'

llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v3.s

+# RUN: yaml2obj %s -o %t.o
+# RUN: llvm-readobj --notes %t.o | FileCheck %s --check-prefix=LLVM
+# RUN: llvm-readelf --notes %t.o | FileCheck %s --check-prefix=GNU
+
+# LLVM:      Notes [
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amdgpu_metadata
+# LLVM-NEXT:     Offset: 0x40
+# LLVM-NEXT:     Size: 0xE8
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMDGPU
+# LLVM-NEXT:       Data size: 0xD4
+# LLVM-NEXT:       Type: NT_AMDGPU_METADATA (AMDGPU Metadata)
+# LLVM-NEXT:       AMDGPU Metadata: Invalid AMDGPU Metadata
+
+# GNU:      Displaying notes found in: .note.nt_amdgpu_metadata
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMDGPU               0x000000d4       NT_AMDGPU_METADATA (AMDGPU Metadata)
+# GNU-NEXT:     AMDGPU Metadata:
+# GNU-NEXT:         Invalid AMDGPU Metadata
+
+--- !ELF
+FileHeader:
+  Class: ELFCLASS64
+  Data:  ELFDATA2LSB
+  Type:  ET_REL
+Sections:
+  - Name:     .note.nt_amdgpu_metadata
+    Type:     SHT_NOTE
+    Content:     07000000D400000020000000414D44475055000081ae616d646873612e6b65726e656c73918ab92e67726f75705f7365676d656e745f66697865645f73697a6502b62e6b65726e6172675f7365676d656e745f616c69676e04b52e6b65726e6172675f7365676d656e745f73697a6501b82e6d61785f666c61745f776f726b67726f75705f73697a6508a52e6e616d65a3666f6fbb2e707269766174655f7365676d656e745f66697865645f73697a6503ab2e736770725f636f756e7406a72e73796d626f6ca3666f6fab2e766770725f636f756e7407af2e7761766566726f6e745f73697a6505

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.s

+# RUN: yaml2obj %s -o %t.o
+# RUN: llvm-readobj --notes %t.o | FileCheck %s --check-prefix=LLVM
+# RUN: llvm-readelf --notes %t.o | FileCheck %s --check-prefix=GNU
+
+# LLVM:      Notes [
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_code_object_version
+# LLVM-NEXT:     Offset: 0x40
+# LLVM-NEXT:     Size: 0x18
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x8
+# LLVM-NEXT:       Type: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
+# LLVM-NEXT:       AMD HSA Code Object Version: [Major: 2, Minor: 1]
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_hsail
+# LLVM-NEXT:     Offset: 0x58
+# LLVM-NEXT:     Size: 0x1C
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0xC
+# LLVM-NEXT:       Type: NT_AMD_HSA_HSAIL (AMD HSA HSAIL Properties)
+# LLVM-NEXT:       AMD HSA HSAIL Properties: [HSAIL Major: 2, HSAIL Minor: 1, Profile: 1, Machine Model: 2, Default Float Round: 3]
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_version
+# LLVM-NEXT:     Offset: 0x74
+# LLVM-NEXT:     Size: 0x2C
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x1B
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# LLVM-NEXT:       AMD HSA ISA Version: [Vendor: AMD, Architecture: AMDGPU, Major: 8, Minor: 0, Stepping: 2]
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_metadata
+# LLVM-NEXT:     Offset: 0xA0
+# LLVM-NEXT:     Size: 0x15C
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x14B
+# LLVM-NEXT:       Type: NT_AMD_HSA_METADATA (AMD HSA Metadata)
+# LLVM-NEXT:       AMD HSA Metadata: ---
+# LLVM-NEXT: Version:         [ 1, 0 ]
+# LLVM-NEXT: Kernels:
+# LLVM-NEXT:   - Name:            elf_notes
+# LLVM-NEXT:     SymbolName:      'elf_notes@kd'
+# LLVM-NEXT:     CodeProps:
+# LLVM-NEXT:       KernargSegmentSize: 0
+# LLVM-NEXT:       GroupSegmentFixedSize: 0
+# LLVM-NEXT:       PrivateSegmentFixedSize: 0
+# LLVM-NEXT:       KernargSegmentAlign: 4
+# LLVM-NEXT:       WavefrontSize:   64
+# LLVM-NEXT:       NumSGPRs:        96
+# LLVM-NEXT:       MaxFlatWorkGroupSize: 1024
+# LLVM-NEXT: ...
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_hsa_isa_name
+# LLVM-NEXT:     Offset: 0x1FC
+# LLVM-NEXT:     Size: 0x2C
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x19
+# LLVM-NEXT:       Type: NT_AMD_HSA_ISA_NAME (AMD HSA ISA Name)
+# LLVM-NEXT:       AMD HSA ISA Name: amdgcn-amd-amdhsa--gfx802
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT:   NoteSection {
+# LLVM-NEXT:     Name: .note.nt_amd_pal_metadata
+# LLVM-NEXT:     Offset: 0x228
+# LLVM-NEXT:     Size: 0x28
+# LLVM-NEXT:     Note {
+# LLVM-NEXT:       Owner: AMD
+# LLVM-NEXT:       Data size: 0x18
+# LLVM-NEXT:       Type: NT_AMD_PAL_METADATA (AMD PAL Metadata)
+# LLVM-NEXT:       AMD PAL Metadata: [2: 1][4: 2][8: 4]
+# LLVM-NEXT:     }
+# LLVM-NEXT:   }
+# LLVM-NEXT: ]
+
+# GNU:       Displaying notes found in: .note.nt_amd_hsa_code_object_version
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x00000008       NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
+# GNU-NEXT:      AMD HSA Code Object Version:
+# GNU-NEXT:          [Major: 2, Minor: 1]
+# GNU-EMPTY:
+# GNU-NEXT:  Displaying notes found in: .note.nt_amd_hsa_hsail
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x0000000c       NT_AMD_HSA_HSAIL (AMD HSA HSAIL Properties)
+# GNU-NEXT:      AMD HSA HSAIL Properties:
+# GNU-NEXT:          [HSAIL Major: 2, HSAIL Minor: 1, Profile: 1, Machine Model: 2, Default Float Round: 3]
+# GNU-EMPTY:
+# GNU-NEXT:  Displaying notes found in: .note.nt_amd_hsa_isa_version
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x0000001b       NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
+# GNU-NEXT:      AMD HSA ISA Version:
+# GNU-NEXT:          [Vendor: AMD, Architecture: AMDGPU, Major: 8, Minor: 0, Stepping: 2]
+# GNU-EMPTY:
+# GNU-NEXT:  Displaying notes found in: .note.nt_amd_hsa_metadata
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x0000014b       NT_AMD_HSA_METADATA (AMD HSA Metadata)
+# GNU-NEXT:      AMD HSA Metadata:
+# GNU-NEXT:          ---
+# GNU-NEXT:  Version:         [ 1, 0 ]
+# GNU-NEXT:  Kernels:
+# GNU-NEXT:    - Name:            elf_notes
+# GNU-NEXT:      SymbolName:      'elf_notes@kd'
+# GNU-NEXT:      CodeProps:
+# GNU-NEXT:        KernargSegmentSize: 0
+# GNU-NEXT:        GroupSegmentFixedSize: 0
+# GNU-NEXT:        PrivateSegmentFixedSize: 0
+# GNU-NEXT:        KernargSegmentAlign: 4
+# GNU-NEXT:        WavefrontSize:   64
+# GNU-NEXT:        NumSGPRs:        96
+# GNU-NEXT:        MaxFlatWorkGroupSize: 1024
+# GNU-NEXT:  ...
+# GNU-EMPTY:
+# GNU-NEXT:  Displaying notes found in: .note.nt_amd_hsa_isa_name
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x00000019       NT_AMD_HSA_ISA_NAME (AMD HSA ISA Name)
+# GNU-NEXT:      AMD HSA ISA Name:
+# GNU-NEXT:          amdgcn-amd-amdhsa--gfx802
+# GNU-EMPTY:
+# GNU-NEXT:  Displaying notes found in: .note.nt_amd_pal_metadata
+# GNU-NEXT:    Owner                Data size        Description
+# GNU-NEXT:    AMD                  0x00000018       NT_AMD_PAL_METADATA (AMD PAL Metadata)
+# GNU-NEXT:      AMD PAL Metadata:
+# GNU-NEXT:          [2: 1][4: 2][8: 4]
+
+--- !ELF
+FileHeader:
+  Class: ELFCLASS64
+  Data:  ELFDATA2LSB
+  Type:  ET_REL
+Sections:
+  - Name:     .note.nt_amd_hsa_code_object_version
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_CODE_OBJECT_VERSION
+        Desc: '0200000001000000'
+  - Name:     .note.nt_amd_hsa_hsail
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_HSAIL
+        Desc: '020000000100000001020300'
+  - Name:     .note.nt_amd_hsa_isa_version
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_VERSION
+        Desc: '04000700080000000000000002000000414d4400414d4447505500'
+  - Name:     .note.nt_amd_hsa_metadata
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_METADATA
+        Desc: '2d2d2d0a56657273696f6e3a2020202020202020205b20312c2030205d0a4b65726e656c733a0a20202d204e616d653a202020202020202020202020656c665f6e6f7465730a2020202053796d626f6c4e616d653a20202020202027656c665f6e6f746573406b64270a20202020436f646550726f70733a0a2020202020204b65726e6172675365676d656e7453697a653a20300a20202020202047726f75705365676d656e74466978656453697a653a20300a202020202020507269766174655365676d656e74466978656453697a653a20300a2020202020204b65726e6172675365676d656e74416c69676e3a20340a2020202020205761766566726f6e7453697a653a20202036340a2020202020204e756d53475052733a202020202020202039360a2020202020204d6178466c6174576f726b47726f757053697a653a20313032340a2e2e2e0a'
+  - Name:     .note.nt_amd_hsa_isa_name 
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_HSA_ISA_NAME 
+        Desc: '616d6467636e2d616d642d616d646873612d2d676678383032'
+  - Name:     .note.nt_amd_pal_metadata 
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMD
+        Type: NT_AMD_PAL_METADATA 
+        Desc: '020000000100000004000000020000000800000004000000'

llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s

+# RUN: yaml2obj %s -o %t.o
+# RUN: llvm-readobj --notes %t.o | FileCheck %s --check-prefix=LLVM
+# RUN: llvm-readelf --notes %t.o | FileCheck %s --check-prefix=GNU
+
+# LLVM:       Notes [
+# LLVM-NEXT:    NoteSection {
+# LLVM-NEXT:      Name: .note.nt_amdgpu_metadata
+# LLVM-NEXT:      Offset: 0x40
+# LLVM-NEXT:      Size: 0x608
+# LLVM-NEXT:      Note {
+# LLVM-NEXT:        Owner: AMDGPU
+# LLVM-NEXT:        Data size: 0x5F1
+# LLVM-NEXT:        Type: NT_AMDGPU_METADATA (AMDGPU Metadata)
+# LLVM-NEXT:        AMDGPU Metadata: ---
+# LLVM-NEXT:  amdhsa.kernels:
+# LLVM-NEXT:    - .args:
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           r
+# LLVM-NEXT:          .offset:         0
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           a
+# LLVM-NEXT:          .offset:         8
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           b
+# LLVM-NEXT:          .offset:         16
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:      .group_segment_fixed_size: 0
+# LLVM-NEXT:      .kernarg_segment_align: 8
+# LLVM-NEXT:      .kernarg_segment_size: 24
+# LLVM-NEXT:      .max_flat_workgroup_size: 1024
+# LLVM-NEXT:      .name:           fadd
+# LLVM-NEXT:      .private_segment_fixed_size: 0
+# LLVM-NEXT:      .sgpr_count:     8
+# LLVM-NEXT:      .sgpr_spill_count: 0
+# LLVM-NEXT:      .symbol:         fadd.kd
+# LLVM-NEXT:      .vgpr_count:     3
+# LLVM-NEXT:      .vgpr_spill_count: 0
+# LLVM-NEXT:      .wavefront_size: 64
+# LLVM-NEXT:    - .args:
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           r
+# LLVM-NEXT:          .offset:         0
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           a
+# LLVM-NEXT:          .offset:         8
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           b
+# LLVM-NEXT:          .offset:         16
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:      .group_segment_fixed_size: 0
+# LLVM-NEXT:      .kernarg_segment_align: 8
+# LLVM-NEXT:      .kernarg_segment_size: 24
+# LLVM-NEXT:      .max_flat_workgroup_size: 1024
+# LLVM-NEXT:      .name:           fsub
+# LLVM-NEXT:      .private_segment_fixed_size: 0
+# LLVM-NEXT:      .sgpr_count:     8
+# LLVM-NEXT:      .sgpr_spill_count: 0
+# LLVM-NEXT:      .symbol:         fsub.kd
+# LLVM-NEXT:      .vgpr_count:     3
+# LLVM-NEXT:      .vgpr_spill_count: 0
+# LLVM-NEXT:      .wavefront_size: 64
+# LLVM-NEXT:    - .args:
+# LLVM-NEXT:        - .name:           i
+# LLVM-NEXT:          .offset:         0
+# LLVM-NEXT:          .size:           4
+# LLVM-NEXT:          .value_kind:     by_value
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           r
+# LLVM-NEXT:          .offset:         8
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           a
+# LLVM-NEXT:          .offset:         16
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:        - .address_space:  global
+# LLVM-NEXT:          .name:           b
+# LLVM-NEXT:          .offset:         24
+# LLVM-NEXT:          .size:           8
+# LLVM-NEXT:          .value_kind:     global_buffer
+# LLVM-NEXT:      .group_segment_fixed_size: 0
+# LLVM-NEXT:      .kernarg_segment_align: 8
+# LLVM-NEXT:      .kernarg_segment_size: 32
+# LLVM-NEXT:      .max_flat_workgroup_size: 1024
+# LLVM-NEXT:      .name:           empty
+# LLVM-NEXT:      .private_segment_fixed_size: 0
+# LLVM-NEXT:      .sgpr_count:     0
+# LLVM-NEXT:      .sgpr_spill_count: 0
+# LLVM-NEXT:      .symbol:         empty.kd
+# LLVM-NEXT:      .vgpr_count:     0
+# LLVM-NEXT:      .vgpr_spill_count: 0
+# LLVM-NEXT:      .wavefront_size: 64
+# LLVM-NEXT:  amdhsa.target:   amdgcn-amd-amdhsa--gfx803
+# LLVM-NEXT:  amdhsa.version:
+# LLVM-NEXT:    - 1
+# LLVM-NEXT:    - 1
+# LLVM-NEXT:  ...
+# LLVM-EMPTY:
+# LLVM-NEXT:      }
+# LLVM-NEXT:    }
+# LLVM-NEXT:  ]
+
+# GNU:      Displaying notes found in: .note.nt_amdgpu_metadata
+# GNU-NEXT:   Owner                Data size        Description
+# GNU-NEXT:   AMDGPU               0x000005f1       NT_AMDGPU_METADATA (AMDGPU Metadata)
+# GNU-NEXT:     AMDGPU Metadata:
+# GNU-NEXT:         ---
+# GNU-NEXT: amdhsa.kernels:
+# GNU-NEXT:   - .args:
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           r
+# GNU-NEXT:         .offset:         0
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           a
+# GNU-NEXT:         .offset:         8
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           b
+# GNU-NEXT:         .offset:         16
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:     .group_segment_fixed_size: 0
+# GNU-NEXT:     .kernarg_segment_align: 8
+# GNU-NEXT:     .kernarg_segment_size: 24
+# GNU-NEXT:     .max_flat_workgroup_size: 1024
+# GNU-NEXT:     .name:           fadd
+# GNU-NEXT:     .private_segment_fixed_size: 0
+# GNU-NEXT:     .sgpr_count:     8
+# GNU-NEXT:     .sgpr_spill_count: 0
+# GNU-NEXT:     .symbol:         fadd.kd
+# GNU-NEXT:     .vgpr_count:     3
+# GNU-NEXT:     .vgpr_spill_count: 0
+# GNU-NEXT:     .wavefront_size: 64
+# GNU-NEXT:   - .args:
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           r
+# GNU-NEXT:         .offset:         0
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           a
+# GNU-NEXT:         .offset:         8
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           b
+# GNU-NEXT:         .offset:         16
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:     .group_segment_fixed_size: 0
+# GNU-NEXT:     .kernarg_segment_align: 8
+# GNU-NEXT:     .kernarg_segment_size: 24
+# GNU-NEXT:     .max_flat_workgroup_size: 1024
+# GNU-NEXT:     .name:           fsub
+# GNU-NEXT:     .private_segment_fixed_size: 0
+# GNU-NEXT:     .sgpr_count:     8
+# GNU-NEXT:     .sgpr_spill_count: 0
+# GNU-NEXT:     .symbol:         fsub.kd
+# GNU-NEXT:     .vgpr_count:     3
+# GNU-NEXT:     .vgpr_spill_count: 0
+# GNU-NEXT:     .wavefront_size: 64
+# GNU-NEXT:   - .args:
+# GNU-NEXT:       - .name:           i
+# GNU-NEXT:         .offset:         0
+# GNU-NEXT:         .size:           4
+# GNU-NEXT:         .value_kind:     by_value
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           r
+# GNU-NEXT:         .offset:         8
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           a
+# GNU-NEXT:         .offset:         16
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:       - .address_space:  global
+# GNU-NEXT:         .name:           b
+# GNU-NEXT:         .offset:         24
+# GNU-NEXT:         .size:           8
+# GNU-NEXT:         .value_kind:     global_buffer
+# GNU-NEXT:     .group_segment_fixed_size: 0
+# GNU-NEXT:     .kernarg_segment_align: 8
+# GNU-NEXT:     .kernarg_segment_size: 32
+# GNU-NEXT:     .max_flat_workgroup_size: 1024
+# GNU-NEXT:     .name:           empty
+# GNU-NEXT:     .private_segment_fixed_size: 0
+# GNU-NEXT:     .sgpr_count:     0
+# GNU-NEXT:     .sgpr_spill_count: 0
+# GNU-NEXT:     .symbol:         empty.kd
+# GNU-NEXT:     .vgpr_count:     0
+# GNU-NEXT:     .vgpr_spill_count: 0
+# GNU-NEXT:     .wavefront_size: 64
+# GNU-NEXT: amdhsa.target:   amdgcn-amd-amdhsa--gfx803
+# GNU-NEXT: amdhsa.version:
+# GNU-NEXT:   - 1
+# GNU-NEXT:   - 1
+# GNU-NEXT: ...
+
+--- !ELF
+FileHeader:
+  Class: ELFCLASS64
+  Data:  ELFDATA2LSB
+  Type:  ET_REL
+Sections:
+  - Name:     .note.nt_amdgpu_metadata
+    Type:     SHT_NOTE
+    Notes:
+      - Name: AMDGPU
+        Type: NT_AMDGPU_METADATA
+        Desc: '83AE616D646873612E6B65726E656C73938DA52E617267739385AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A172A72E6F666673657400A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A161A72E6F666673657408A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A162A72E6F666673657410A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F627566666572B92E67726F75705F7365676D656E745F66697865645F73697A6500B62E6B65726E6172675F7365676D656E745F616C69676E08B52E6B65726E6172675F7365676D656E745F73697A6518B82E6D61785F666C61745F776F726B67726F75705F73697A65CD0400A52E6E616D65A466616464BB2E707269766174655F7365676D656E745F66697865645F73697A6500AB2E736770725F636F756E7408B12E736770725F7370696C6C5F636F756E7400A72E73796D626F6CA7666164642E6B64AB2E766770725F636F756E7403B12E766770725F7370696C6C5F636F756E7400AF2E7761766566726F6E745F73697A65408DA52E617267739385AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A172A72E6F666673657400A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A161A72E6F666673657408A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A162A72E6F666673657410A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F627566666572B92E67726F75705F7365676D656E745F66697865645F73697A6500B62E6B65726E6172675F7365676D656E745F616C69676E08B52E6B65726E6172675F7365676D656E745F73697A6518B82E6D61785F666C61745F776F726B67726F75705F73697A65CD0400A52E6E616D65A466737562BB2E707269766174655F7365676D656E745F66697865645F73697A6500AB2E736770725F636F756E7408B12E736770725F7370696C6C5F636F756E7400A72E73796D626F6CA7667375622E6B64AB2E766770725F636F756E7403B12E766770725F7370696C6C5F636F756E7400AF2E7761766566726F6E745F73697A65408DA52E617267739484A52E6E616D65A169A72E6F666673657400A52E73697A6504AB2E76616C75655F6B696E64A862795F76616C756585AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A172A72E6F666673657408A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A161A72E6F666673657410A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65A162A72E6F666673657418A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F627566666572B92E67726F75705F7365676D656E745F66697865645F73697A6500B62E6B65726E6172675F7365676D656E745F616C69676E08B52E6B65726E6172675F7365676D656E745F73697A6520B82E6D61785F666C61745F776F726B67726F75705F73697A65CD0400A52E6E616D65A5656D707479BB2E707269766174655F7365676D656E745F66697865645F73697A6500AB2E736770725F636F756E7400B12E736770725F7370696C6C5F636F756E7400A72E73796D626F6CA8656D7074792E6B64AB2E766770725F636F756E7400B12E766770725F7370696C6C5F636F756E7400AF2E7761766566726F6E745F73697A6540AD616D646873612E746172676574B9616D6467636E2D616D642D616D646873612D2D676678383033AE616D646873612E76657273696F6E920101'

llvm/tools/llvm-readobj/ELFDumper.cpp

case ELF::NT_AMD_HSA_HSAIL: {
     };
     if (Desc.size() != sizeof(HSAILProperties))
       return {"AMD HSA HSAIL Properties", "Invalid AMD HSA HSAIL Properties"};
     auto Properties = reinterpret_cast<const HSAILProperties *>(Desc.data());
     std::string HSAILPropetiesString;
     raw_string_ostream StrOS(HSAILPropetiesString);
     StrOS << "[HSAIL Major: " << Properties->HSAILMajorVersion
           << ", HSAIL Minor: " << Properties->HSAILMinorVersion
-          << ", Profile: " << Properties->Profile
-          << ", Machine Model: " << Properties->MachineModel
-          << ", Default Float Round: " << Properties->DefaultFloatRound << "]";
+          << ", Profile: "
+          << uint32_t(Properties->Profile)
+          << ", Machine Model: "
+          << uint32_t(Properties->MachineModel)
+          << ", Default Float Round: "
+          << uint32_t(Properties->DefaultFloatRound) << "]";
     return {"AMD HSA HSAIL Properties", HSAILPropetiesString};
   }
   case ELF::NT_AMD_HSA_ISA_VERSION: {
     struct IsaVersion {
       uint16_t VendorNameSize;
       uint16_t ArchitectureNameSize;
       uint32_t Major;
       uint32_t Minor;
static AMDNote getAMDNote(uint32_t NoteType, ArrayRef<uint8_t> Desc) {
   case ELF::NT_AMD_PAL_METADATA: {
     struct PALMetadata {
       uint32_t Key;
       uint32_t Value;
     };
     auto Isa = reinterpret_cast<const PALMetadata *>(Desc.data());
     std::string MetadataString;
     raw_string_ostream StrOS(MetadataString);
-    for (size_t I = 0, E = Desc.size() / sizeof(PALMetadata); I < E; ++E) {
+    for (size_t I = 0, E = Desc.size() / sizeof(PALMetadata); I < E; ++I) {
       StrOS << "[" << Isa[I].Key << ": " << Isa[I].Value << "]";
     }
     return {"AMD PAL Metadata", MetadataString};
   }
   }
 }
 
 struct AMDGPUNote {