[llvm-readobj][AMDGPU] Bypass MD verification for PAL

[Pierre-vh]

Summary

Small split change from D146023.

Migrate elf-notes to v4 and fix llvm-readobj to work with PAL metadata.

[arsenm]

llvm/tools/llvm-readobj/ELFDumper.cpp

5435

Why, I don't understand?

5438

Why isn't this on only under an assert?

5438

Nevermind, this isn't codegen

[jhenderson]

One nit from me. Somebody who understands AMDGPU will need to do a proper review though.

llvm/tools/llvm-readobj/ELFDumper.cpp

5436

[Pierre-vh]

Comment

llvm/tools/llvm-readobj/ELFDumper.cpp

5435

It's only designed for AMDHSA, it looks for keys that have amdhsa (e.g. amdhsa.version) and fails when it doesn't find them

[Pierre-vh]

Rebase (ping?)

[Pierre-vh]

Ping

[foad]

llvm/tools/llvm-readobj/ELFDumper.cpp

5437

Since this is AMDHSA stuff, could you check for the presence of "amdhsa." instead of the absence of "amdpal."?

[Pierre-vh]

Change check

[Pierre-vh]

Ping

[kzhuravl]

looks good

llvm/test/CodeGen/AMDGPU/elf-notes.ll

 ; OSABI-UNK-ELF: AMD HSA ISA Name:
 ; OSABI-UNK-ELF: amdgcn-amd-unknown--gfx802
 ; OSABI-UNK-ELF-NOT: Unknown note type
 ; OSABI-UNK-ELF-NOT: NT_AMD_HSA_METADATA (AMD HSA Metadata)
 ; OSABI-UNK-ELF-NOT: Unknown note type
 ; 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@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
 ; OSABI-PAL-ELF-NOT: NT_AMD_HSA_METADATA (AMD HSA Metadata)
 ; OSABI-PAL-ELF: NT_AMDGPU_METADATA (AMDGPU Metadata)
 ; OSABI-PAL-ELF: AMDGPU Metadata:
 ; OSABI-PAL-ELF: amdpal.pipelines:
 ; OSABI-PAL-ELF:   - .hardware_stages:
 ; OSABI-PAL-ELF:       .cs:
 ; OSABI-PAL-ELF:         .entry_point:    elf_notes
 ; OSABI-PAL-ELF:         .scratch_memory_size: 0
 ; OSABI-PAL-ELF:         .sgpr_count:     96
 ; OSABI-PAL-ELF:         .vgpr_count:     1
 ; 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
 ; R600-NOT: .amd_amdgpu_isa
 ; R600-NOT: .amd_amdgpu_hsa_metadata
 ; R600-NOT: .amd_amdgpu_pal_metadata
 
 define amdgpu_kernel void @elf_notes() {
   ret void
 }
 
 !llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_code_object_version", i32 200}
+!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}

llvm/tools/llvm-readobj/ELFDumper.cpp

default:
     return {"", ""};
   case ELF::NT_AMDGPU_METADATA: {
     StringRef MsgPackString =
         StringRef(reinterpret_cast<const char *>(Desc.data()), Desc.size());
     msgpack::Document MsgPackDoc;
     if (!MsgPackDoc.readFromBlob(MsgPackString, /*Multi=*/false))
       return {"", ""};
 
-    AMDGPU::HSAMD::V3::MetadataVerifier Verifier(true);
     std::string MetadataString;
+
+    // FIXME: Metadata Verifier doesn't work with AMDPAL MD.
+    //  This is an ugly workaround to avoid the verifier.
+    if (MsgPackString.find("amdpal.") == 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()) {
       // TODO: passing a scalar root to toYAML() asserts:
       // (PolymorphicTraits<T>::getKind(Val) != NodeKind::Scalar &&
       //    "plain scalar documents are not supported")
       // To avoid this crash we print the raw data instead.
       return {"", ""};