[PATCH 0/7] Add AMDGCN support to readelf

[Simon Marchi]

This goal of this series is to add to readelf the capability to inspect
ELF files of the AMDGCN architecture.  The AMDGCN architecture is an AMD
GPU architecture that can be used to execute compute kernelsi.  ELF
files produced for that architecture (using the HIP compiler [3])
typically don't exist by themselves, but are bundled inside the ELF file
of the host program that will submit the kernels for execution on the
GPU.  But other than that, they are regular ELF files.

The documentation for the ELF file format for the architecture can be
found here [1].  I refer to it a few times throughout the patches.

Our end goal is to upstream the GDB port for AMDGCN.  But the readelf
bits are relatively simple, independent of GDB and useful on their own,
which is why I submit them first.

I uploaded this series to a users branch for convenience [4].

Since it may not be trivial for those trying for the first time, here
are instructions on how to install the ROCm stack and generate a binary
that can be used to test what's included in this series.  Is it based on
the instructions found here [2].  I suggest using Ubuntu 20.04, since
that's what I tested with.  This can all be done in a Docker container
to avoid polluting the system (it ends up installed a lot of packages).

 - Install the amdgpu-install tool

    $ wget https://repo.radeon.com/amdgpu-install/21.50/ubuntu/focal/amdgpu-install_21.50.50000-1_all.deb
    $ sudo apt-get install ./amdgpu-install_21.50.50000-1_all.deb

 - Install the ROCm packages

    $ sudo amdgpu-install --usecase=rocm

 - Checkout the users/simark/amdgpu-readelf binutils-gdb branch

    $ git checkout users/simark/amdgpu-readelf

 - Configure and build

    $ ./configure --target=amdgcn-amd-amdhsa --disable-gas
    $ make all-binutils

 - Write a test source file with

    #include "hip/hip_runtime.h"
    __global__ void the_kernel() { __builtin_amdgcn_s_sleep(1); }
    int main() { hipLaunchKernelGGL(the_kernel, dim3(1), dim3(1), 0, 0); }

 - Compile the test source file

    $ hipcc test.cpp

 - Extract the bundled AMDGCN ELF file from the host ELF file (adjust
   the URI accordingly, based on the output of roc-obj-ls)

    $ /opt/rocm/bin/roc-obj-ls a.out
    1       host-x86_64-unknown-linux                                           file://a.out#offset=8192&size=0
    1       hipv4-amdgcn-amd-amdhsa--gfx803                                     file://a.out#offset=8192&size=9536
    $ /opt/rocm/bin/roc-obj-extract "file://a.out#offset=8192&size=9536"

 - Run readelf on the extracted binary

    $ ./binutils/readelf --header a.out-offset8192-size9536.co

[1] https://llvm.org/docs/AMDGPUUsage.html#elf-code-object
[2] https://github.com/RadeonOpenCompute/ROCm/blob/fac29ca466c7362a659ca4cd53fc2abc798a3c90/ROCm_Installation_Guide%20v5.0.pdf
[3] https://github.com/ROCm-Developer-Tools/HIP
[4] https://sourceware.org/git/?p=binutils-gdb.git;a=shortlog;h=refs/heads/users/simark/amdgpu-readelf

Simon Marchi (7):
  bfd: add AMDGCN architecture
  opcodes: handle bfd_amdgcn_arch in configure script
  binutils/readelf: handle AMDGPU OS ABIs
  binutils/readelf: decode AMDGPU-specific e_flags
  binutils/readelf: handle NT_AMDGPU_METADATA note name
  binutils/readelf: build against msgpack, dump NT_AMDGPU_METADATA note
    contents
  binutils/readelf: handle AMDGPU relocation types

 bfd/Makefile.am       |   4 +
 bfd/Makefile.in       |   6 +
 bfd/archures.c        |  15 ++
 bfd/bfd-in2.h         |  13 ++
 bfd/config.bfd        |   5 +
 bfd/configure         |   1 +
 bfd/configure.ac      |   1 +
 bfd/cpu-amdgcn.c      |  59 +++++++
 bfd/elf-bfd.h         |   1 +
 bfd/elf64-amdgcn.c    |  80 +++++++++
 bfd/targets.c         |   5 +
 binutils/Makefile.am  |  14 +-
 binutils/Makefile.in  |  11 +-
 binutils/config.in    |   3 +
 binutils/configure    | 137 ++++++++++++++-
 binutils/configure.ac |  17 ++
 binutils/readelf.c    | 379 ++++++++++++++++++++++++++++++++++++++++--
 include/elf/amdgpu.h  | 115 +++++++++++++
 include/elf/common.h  |   3 +
 opcodes/configure     |   1 +
 opcodes/configure.ac  |   1 +
 21 files changed, 852 insertions(+), 19 deletions(-)
 create mode 100644 bfd/cpu-amdgcn.c
 create mode 100644 bfd/elf64-amdgcn.c
 create mode 100644 include/elf/amdgpu.h


base-commit: af481f01297a61af30f3f3dbaa11e6636c587dcb
-- 
2.35.1