public inbox for binutils@sourceware.org
 help / color / mirror / Atom feed
* [PATCH 0/7] Add AMDGCN support to readelf
@ 2022-03-15 19:42 Simon Marchi
  2022-03-15 19:42 ` [PATCH 1/7] bfd: add AMDGCN architecture Simon Marchi
                   ` (7 more replies)
  0 siblings, 8 replies; 10+ messages in thread
From: Simon Marchi @ 2022-03-15 19:42 UTC (permalink / raw)
  To: binutils; +Cc: 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


^ permalink raw reply	[flat|nested] 10+ messages in thread

end of thread, other threads:[~2022-03-16 13:02 UTC | newest]

Thread overview: 10+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2022-03-15 19:42 [PATCH 0/7] Add AMDGCN support to readelf Simon Marchi
2022-03-15 19:42 ` [PATCH 1/7] bfd: add AMDGCN architecture Simon Marchi
2022-03-15 19:42 ` [PATCH 2/7] opcodes: handle bfd_amdgcn_arch in configure script Simon Marchi
2022-03-15 19:42 ` [PATCH 3/7] binutils/readelf: handle AMDGPU OS ABIs Simon Marchi
2022-03-15 19:43 ` [PATCH 4/7] binutils/readelf: decode AMDGPU-specific e_flags Simon Marchi
2022-03-15 19:43 ` [PATCH 5/7] binutils/readelf: handle NT_AMDGPU_METADATA note name Simon Marchi
2022-03-15 19:43 ` [PATCH 6/7] binutils/readelf: build against msgpack, dump NT_AMDGPU_METADATA note contents Simon Marchi
2022-03-15 19:43 ` [PATCH 7/7] binutils/readelf: handle AMDGPU relocation types Simon Marchi
2022-03-16  0:08 ` [PATCH 0/7] Add AMDGCN support to readelf Alan Modra
2022-03-16 13:02   ` Simon Marchi

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).