[llvm] 415956f - [llvm-readobj][AMDGPU] Bypass MD verification for PAL
via llvm-commits
llvm-commits at lists.llvm.org
Tue May 2 23:45:30 PDT 2023
Author: pvanhout
Date: 2023-05-03T08:45:24+02:00
New Revision: 415956fe7ef276c4dd1d3e4a689c47d24aa034dc
URL: https://github.com/llvm/llvm-project/commit/415956fe7ef276c4dd1d3e4a689c47d24aa034dc
DIFF: https://github.com/llvm/llvm-project/commit/415956fe7ef276c4dd1d3e4a689c47d24aa034dc.diff
LOG: [llvm-readobj][AMDGPU] Bypass MD verification for PAL
Small split change from D146023.
Migrate elf-notes to v4 and fix llvm-readobj to work with PAL metadata.
Reviewed By: kzhuravl
Differential Revision: https://reviews.llvm.org/D146119
Added:
Modified:
llvm/test/CodeGen/AMDGPU/elf-notes.ll
llvm/tools/llvm-readobj/ELFDumper.cpp
Removed:
################################################################################
diff --git a/llvm/test/CodeGen/AMDGPU/elf-notes.ll b/llvm/test/CodeGen/AMDGPU/elf-notes.ll
index 7c78a92486494..37299366e13c3 100644
--- a/llvm/test/CodeGen/AMDGPU/elf-notes.ll
+++ b/llvm/test/CodeGen/AMDGPU/elf-notes.ll
@@ -25,40 +25,38 @@
; OSABI-UNK-ELF-NOT: NT_AMD_PAL_METADATA (AMD PAL Metadata)
; OSABI-UNK-ELF-NOT: Unknown note type
-; OSABI-HSA: .hsa_code_object_version
-; OSABI-HSA: .hsa_code_object_isa
-; OSABI-HSA: .amd_amdgpu_isa "amdgcn-amd-amdhsa--gfx802"
-; OSABI-HSA: .amd_amdgpu_hsa_metadata
+; OSABI-HSA: amdhsa.target: amdgcn-amd-amdhsa--gfx802
+; OSABI-HSA: amdhsa.version:
+; OSABI-HSA: .end_amdgpu_metadata
; OSABI-HSA-NOT: .amd_amdgpu_pal_metadata
-; OSABI-HSA-ELF: NT_AMD_HSA_CODE_OBJECT_VERSION (AMD HSA Code Object Version)
-; OSABI-HSA-ELF: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
-; OSABI-HSA-ELF: NT_AMD_HSA_ISA_NAME (AMD HSA ISA Name)
-; OSABI-HSA-ELF: AMD HSA ISA Name:
-; OSABI-HSA-ELF: amdgcn-amd-amdhsa--gfx802
-; OSABI-HSA-ELF: NT_AMD_HSA_METADATA (AMD HSA Metadata)
-; OSABI-HSA-ELF: HSA Metadata:
+; OSABI-HSA-ELF: NT_AMDGPU_METADATA (AMDGPU Metadata)
; OSABI-HSA-ELF: ---
-; OSABI-HSA-ELF: Version: [ 1, 0 ]
-; OSABI-HSA-ELF: Kernels:
-; OSABI-HSA-ELF: - Name: elf_notes
-; OSABI-HSA-ELF: SymbolName: 'elf_notes at kd'
-; OSABI-HSA-ELF: CodeProps:
-; OSABI-HSA-ELF: KernargSegmentSize: 0
-; OSABI-HSA-ELF: GroupSegmentFixedSize: 0
-; OSABI-HSA-ELF: PrivateSegmentFixedSize: 0
-; OSABI-HSA-ELF: KernargSegmentAlign: 4
-; OSABI-HSA-ELF: WavefrontSize: 64
-; OSABI-HSA-ELF: NumSGPRs: 96
+; OSABI-HSA-ELF: amdhsa.kernels:
+; OSABI-HSA-ELF: - .args: []
+; OSABI-HSA-ELF: .group_segment_fixed_size: 0
+; OSABI-HSA-ELF: .kernarg_segment_align: 4
+; OSABI-HSA-ELF: .kernarg_segment_size: 0
+; OSABI-HSA-ELF: .max_flat_workgroup_size: 1024
+; OSABI-HSA-ELF: .name: elf_notes
+; OSABI-HSA-ELF: .private_segment_fixed_size: 0
+; OSABI-HSA-ELF: .sgpr_count: 96
+; OSABI-HSA-ELF: .sgpr_spill_count: 0
+; OSABI-HSA-ELF: .symbol: elf_notes.kd
+; OSABI-HSA-ELF: .vgpr_count: 0
+; OSABI-HSA-ELF: .vgpr_spill_count: 0
+; OSABI-HSA-ELF: .wavefront_size: 64
+; OSABI-HSA-ELF: amdhsa.target: amdgcn-amd-amdhsa--gfx802
+; OSABI-HSA-ELF: amdhsa.version:
+; OSABI-HSA-ELF: - 1
+; OSABI-HSA-ELF: - 1
; OSABI-HSA-ELF: ...
; OSABI-HSA-ELF-NOT: NT_AMD_PAL_METADATA (AMD PAL Metadata)
-; OSABI-PAL-NOT: .hsa_code_object_version
-; OSABI-PAL: .hsa_code_object_isa
-; OSABI-PAL: .amd_amdgpu_isa "amdgcn-amd-amdpal--gfx802"
+; OSABI-PAL: .amd_amdgpu_isa "amdgcn-amd-amdpal--gfx802"
+; OSABI-PAL: .amdgpu_pal_metadata
; OSABI-PAL-NOT: .amd_amdgpu_hsa_metadata
-; OSABI-PAL-ELF: NT_AMD_HSA_ISA_VERSION (AMD HSA ISA Version)
; OSABI-PAL-ELF: NT_AMD_HSA_ISA_NAME (AMD HSA ISA Name)
; OSABI-PAL-ELF: AMD HSA ISA Name:
; OSABI-PAL-ELF: amdgcn-amd-amdpal--gfx802
@@ -75,16 +73,6 @@
; OSABI-PAL-ELF: .registers:
; OSABI-PAL-ELF: 11794: 11469504
; OSABI-PAL-ELF: 11795: 128
-; OSABI-PAL: amdpal.pipelines:
-; OSABI-PAL: - .hardware_stages:
-; OSABI-PAL: .cs:
-; OSABI-PAL: .entry_point: elf_notes
-; OSABI-PAL: .scratch_memory_size: 0
-; OSABI-PAL: .sgpr_count: 0x60
-; OSABI-PAL: .vgpr_count: 0x1
-; OSABI-PAL: .registers:
-; OSABI-PAL: 0x2e12 (COMPUTE_PGM_RSRC1): 0xaf02c0
-; OSABI-PAL: 0x2e13 (COMPUTE_PGM_RSRC2): 0x80
; R600-NOT: .hsa_code_object_version
; R600-NOT: .hsa_code_object_isa
@@ -97,4 +85,4 @@ define amdgpu_kernel void @elf_notes() {
}
!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 200}
+!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
diff --git a/llvm/tools/llvm-readobj/ELFDumper.cpp b/llvm/tools/llvm-readobj/ELFDumper.cpp
index 43d9a0f576aca..0a32085dfa649 100644
--- a/llvm/tools/llvm-readobj/ELFDumper.cpp
+++ b/llvm/tools/llvm-readobj/ELFDumper.cpp
@@ -5514,10 +5514,16 @@ static AMDGPUNote getAMDGPUNote(uint32_t NoteType, ArrayRef<uint8_t> Desc) {
if (!MsgPackDoc.readFromBlob(MsgPackString, /*Multi=*/false))
return {"", ""};
- AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
std::string MetadataString;
- if (!Verifier.verify(MsgPackDoc.getRoot()))
- MetadataString = "Invalid AMDGPU Metadata\n";
+
+ // FIXME: Metadata Verifier only works with AMDHSA.
+ // This is an ugly workaround to avoid the verifier for other MD
+ // formats (e.g. amdpal)
+ if (MsgPackString.find("amdhsa.") != StringRef::npos) {
+ AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
+ if (!Verifier.verify(MsgPackDoc.getRoot()))
+ MetadataString = "Invalid AMDGPU Metadata\n";
+ }
raw_string_ostream StrOS(MetadataString);
if (MsgPackDoc.getRoot().isScalar()) {
More information about the llvm-commits
mailing list