Index: llvm/lib/ObjectYAML/ELFYAML.cpp =================================================================== --- llvm/lib/ObjectYAML/ELFYAML.cpp +++ llvm/lib/ObjectYAML/ELFYAML.cpp @@ -155,6 +155,9 @@ 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); Index: llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v2.s =================================================================== --- /dev/null +++ llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v2.s @@ -0,0 +1,151 @@ +# 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' Index: llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v3.s =================================================================== --- /dev/null +++ llvm/test/tools/llvm-readobj/ELF/note-amd-invalid-v3.s @@ -0,0 +1,30 @@ +# 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 Index: llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.s =================================================================== --- /dev/null +++ llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v2.s @@ -0,0 +1,177 @@ +# 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' Index: llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s =================================================================== --- /dev/null +++ llvm/test/tools/llvm-readobj/ELF/note-amd-valid-v3.s @@ -0,0 +1,225 @@ +# 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' Index: llvm/tools/llvm-readobj/ELFDumper.cpp =================================================================== --- llvm/tools/llvm-readobj/ELFDumper.cpp +++ llvm/tools/llvm-readobj/ELFDumper.cpp @@ -5055,9 +5055,12 @@ 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: { @@ -5108,7 +5111,7 @@ auto Isa = reinterpret_cast(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};