[Proposal]: Per-Architecture DWP Packaging for Split-DWARF

[qxy11]

Background

Most production binaries are built with split-DWARF to minimize the size of binaries. During development on a local machine, the debugger can locate these .dwo files relative to the build tree. However, this original build tree is gone in any production or remote debugging scenario and coredump analysis on SEVs. Without a .dwp file, the debugger has no debug info to retrieve variable names and source locations.

DWP packaging makes split-DWARF debug info usable outside the original build environment. Since we currently have no standard for DWP packaging of GPU DWOs and can't include device-side debug info without greatly bloating production binary sizes, GPU debugging in production is non-functional.

Architecture Mismatch

When we compile a HIP source file with split-dwarf for multiple GPU architectures, we get DWO files for distinct architectures. For example, if we run hipcc -gsplit-dwarf -g --offload-arch=gfx942 --offload-arch=gfx950 -c foo.hip:

• Host compilation (x86_64): produces foo.dwo
• Device compilation per GPU arch: produces foo_gfx942.dwo, foo_gfx950.dwo, etc.

The resulting .dwo files target different CPU/GPU architectures, and must all be correctly packaged into DWPs for both host and device debugging to work in production scenarios.
For example:

$ hipcc -gsplit-dwarf -g --offload-arch=gfx942 --offload-arch=gfx950 -c saxpy.hip

$ llvm-dwarfdump --debug-info saxpy.dwo | head -5
saxpy.dwo:	file format elf64-x86-64
.debug_info.dwo contents:
0x00000000: Compile Unit: ... unit_type = DW_UT_split_compile, ...
  DWO_id = 0x2ba31cef503e1158

$ llvm-dwarfdump --debug-info saxpy_gfx942.dwo | head -5
saxpy_gfx942.dwo:	file format elf64-amdgpu
.debug_info.dwo contents:
0x00000000: Compile Unit: ... unit_type = DW_UT_split_compile, ...
  DWO_id = 0xad21fe8a2dfeb5bb

$ llvm-dwarfdump --debug-info saxpy_gfx950.dwo | head -5
saxpy_gfx950.dwo:	file format elf64-amdgpu
.debug_info.dwo contents:
0x00000000: Compile Unit: ... unit_type = DW_UT_split_compile, ...
  DWO_id = 0x10ccaff1dd367ead

We need to have some standardized way of packaging these CPU and GPU DWOs.

Options

Option 1: Treat Architectures as Compatible

When we tested on ROCm 6.2.1, 6.4.0, 6.4.2, all produced identical DWO IDs across GPU architectures, causing llvm-dwp to fail with error: duplicate DWO ID (FA993CFF029EEF25). If we use ROCm 7.0, there are no issues with combining all of these DWOs into one file. However, the order of inputs may matter, as llvm-dwp seems to determine its output ELF container architecture from the first input file, even if the DWP contains DWARF CUs from both x86-64 and AMDGCN targets.

# Host DWO first → x86-64 container
$ llvm-dwp -o host_first.dwp saxpy.dwo saxpy_gfx942.dwo saxpy_gfx950.dwo
$ readelf -h host_first.dwp | grep "Machine:"
  Machine: Advanced Micro Devices X86-64

# Device DWO first → AMDGPU container
$ llvm-dwp -o device_first.dwp saxpy_gfx942.dwo saxpy.dwo saxpy_gfx950.dwo
$ readelf -h device_first.dwp | grep "Machine:"
  Machine: AMD GPU

ROCgdb seems to tolerate this architecture mismatch in the DWP load, but still needs to learn to search this DWP for GPU code objects. Right now, it loads the GPU code object as file:///...saxpy#offset=8192&size=9800, and tries to load file:///...saxpy#offset=8192&size=9800.dwp.

Option 2: Separate DWP Per Architecture

We could produce separate DWP files per architecture. This way, the host DWP is elf64-86-64, and each device DWP is elf64-amdgpu. There are no architecture mismatches, but it produces 1 + N DWPs for each GPU architecture. For a binary foo compiled for both gfx942 and gfx950:

• foo.dwp: host debug info only (x86-64)
• foo.gfx942.dwp: gfx942 device debug info only (AMDGPU)
• foo.gfx950.dwp: gfx950 device debug info only (AMDGPU)

This option would require us to have some industry standard naming convention, and we would need to patch GPU debuggers to know to look for <binary>.<arch>.dwp.

Proposal

We'd like to propose adopting Option 2 as standard, and creating a patch so that when debugging a GPU code object for architecture gfxNNNN embedded in a host binary, the debugger searches for <host_binary>.gfxNNNN.dwp alongside the host binary.

[clayborg]

Since DWP files are ELF files, and these files have architectures in the ELF header, I would prefer to see the "Separate DWP Per Architecture" approach.

[simark]

It's not clear to me if there is anything really architecture-specific in the DWP format itself (the headers, the CU/TU index sections, etc). So I don't see any technical reason why it wouldn't work to put everything in one DWP. But I do agree that it would make things a bit awkward, and that doing one DWP per architecture seems safer.

Some questions I have:

• is it possible to have more than one embedded ELF for a given arch? In this case, would you still have one DWP containing the debug info for all the ELFs of that arch? Because if not, there would be a clash if the naming scheme was just foo.gfx942.dwp.
• Would it be possible to re-use the ELF-in-ELF offload bundle format for DWPs? The top-level DWP would be the for top-level (host) ELF, and it could contain the DWPs corresponding to the offload ELF files. The debugger would need a way to know which embedded DWP matches which embedded ELF.

See this proposal [DWARF Package file (.dwp) .debug_dwp ID] that has been accepted for DWARF 6:

https://dwarfstd.org/issues/260116.1.html

I think that could be useful for us, to match ELF files vs DWP files.

[qxy11]

is it possible to have more than one embedded ELF for a given arch? In this case, would you still have one DWP containing the debug info for all the ELFs of that arch?

It's possible to have more than one embedded ELF for a given arch. When you link multiple .hip translation units, each TU that contains device code (i.e. __global__ or __device__ functions) produces its own offload bundle in the .hip_fatbin section, and each bundle contains a separate code object per GPU arch. So linking 2 TUs with kernels targeting gfx942 gives you 2 embedded gfx942 ELFs in the final binary.

This doesn't seem to cause a naming clash with the foo.gfx942.dwp scheme and you can still have just one DWP per arch. It would contain the CUs from all TUs for that arch. From my simple tests, it seems like DWO IDs are unique across TUs (at least on ROCm 7.0+), so the debugger can match skeleton CU to DWP CU entry by DWO ID regardless of how many code objects exist for that arch.

Would it be possible to re-use the ELF-in-ELF offload bundle format for DWPs?

Using .debug_dwp seems like the right approach. It seems like we would need some approach for llvm-dwp to generate an ID and embed it in a .debug_dwp section in the output DWP

[clayborg]

It's not clear to me if there is anything really architecture-specific in the DWP format itself (the headers, the CU/TU index sections, etc). So I don't see any technical reason why it wouldn't work to put everything in one DWP. But I do agree that it would make things a bit awkward, and that doing one DWP per architecture seems safer.

Some questions I have:

• is it possible to have more than one embedded ELF for a given arch? In this case, would you still have one DWP containing the debug info for all the ELFs of that arch? Because if not, there would be a clash if the naming scheme was just foo.gfx942.dwp.

It is ok for all of the .dwo files for any contained ELF files be in one .dwp file as they are looked up by a DWO signature. The main issue is given the executable path, how do you find the right .dwp file?

Lets say we have an executable /tmp/a.out which is a native executable, but it also contains 3 ELF files for gfx942, each of which have 2 .dwo files. The current DWARF code will look in various areas for the .dwp file. It starts by appending .dwp to the executable name (/tmp/a.out.dwp). If it finds it, then it will use it. So if we wanted to keep things in the same .dwp file for both native and GPU, this is the easiest solution. The main problem with the single file solution is the DWARF information has information that specifies registers (the DWARF location expressions) and the unwind info uses register numbers (.debug_frame), so it isn't a perfect solution in that regard.

When debugging the GPU side, the ELF file's path will also be /tmp/a.out, but it will be at a file offset and size within the main executable. But if we just take the path /tmp/a.out and append .dwp, we will find a file, but will need to check the architecture to verify it matches. If it doesn't, we can append an architecture to the name and try looking for that. The issue we have right now is gfx942 isn't an architecture, but an environment: "amdgcn-amd--gfx942". This is how we identify an exact AMD GPU type right now. So we could try appending the architecture name amdgcn to get /tmp/a.out.amdgcn.dwp, and if that fails, try appending the arch name and the environment /tmp/a.out.amdgcn.gfx942.dwp. Either way, we will need a way to locate this second file.

Can a binary contian multiple GPU embedded ELF files? Or will there only ever be one? Like, can an executable contain 2 ELF files for gfx942, and 2 ELF files for gfx801?

• Would it be possible to re-use the ELF-in-ELF offload bundle format for DWPs? The top-level DWP would be the for top-level (host) ELF, and it could contain the DWPs corresponding to the offload ELF files. The debugger would need a way to know which embedded DWP matches which embedded ELF.

See this proposal [DWARF Package file (.dwp) .debug_dwp ID] that has been accepted for DWARF 6:

https://dwarfstd.org/issues/260116.1.html

I think that could be useful for us, to match ELF files vs DWP files.

Yes, this seems to be a good way to match things up. We should use this. Then the contained ELF files can have a .debug_dwp section that points to the right .dwp file for the right GPU and the name doesn't matter as it will be specified.

[lancesix]

Another thing to note is that with our runtime, code objects can be reported as being loaded from the inferior memory:

(gdb) info sharedlibrary mem
From                To                  Syms Read   Shared Object Library
0x00007ffff5f26e00  0x00007ffff5f29f00  Yes (*)     memory://803731#offset=0x555555de3090&size=45104

In such case, we don't really have a file we can manipulate to add some architecture name suffix. The .debug_dwp section would look like the most flexible approach, but might require a bit of tooling work to make sure the GPU code objects embeded in the fat binary have the proper section appended.

[simark]

It's not clear to me if there is anything really architecture-specific in the DWP format itself (the headers, the CU/TU index sections, etc). So I don't see any technical reason why it wouldn't work to put everything in one DWP. But I do agree that it would make things a bit awkward, and that doing one DWP per architecture seems safer.
Some questions I have:

• is it possible to have more than one embedded ELF for a given arch? In this case, would you still have one DWP containing the debug info for all the ELFs of that arch? Because if not, there would be a clash if the naming scheme was just foo.gfx942.dwp.

It is ok for all of the .dwo files for any contained ELF files be in one .dwp file as they are looked up by a DWO signature. The main issue is given the executable path, how do you find the right .dwp file?

Lets say we have an executable /tmp/a.out which is a native executable, but it also contains 3 ELF files for gfx942, each of which have 2 .dwo files. The current DWARF code will look in various areas for the .dwp file. It starts by appending .dwp to the executable name (/tmp/a.out.dwp). If it finds it, then it will use it. So if we wanted to keep things in the same .dwp file for both native and GPU, this is the easiest solution. The main problem with the single file solution is the DWARF information has information that specifies registers (the DWARF location expressions) and the unwind info uses register numbers (.debug_frame), so it isn't a perfect solution in that regard.

When debugging the GPU side, the ELF file's path will also be /tmp/a.out, but it will be at a file offset and size within the main executable. But if we just take the path /tmp/a.out and append .dwp, we will find a file, but will need to check the architecture to verify it matches. If it doesn't, we can append an architecture to the name and try looking for that. The issue we have right now is gfx942 isn't an architecture, but an environment: "amdgcn-amd--gfx942". This is how we identify an exact AMD GPU type right now. So we could try appending the architecture name amdgcn to get /tmp/a.out.amdgcn.dwp, and if that fails, try appending the arch name and the environment /tmp/a.out.amdgcn.gfx942.dwp. Either way, we will need a way to locate this second file.

To be clear, my understanding is that a single .dwp file could very well hold .dwo files targetting different amdgcn processors (gfxNNNN), or different flags, assuming that there is no DWO ID clash. But it might just make the workflow easier, if the binaries (and their associated .dwp files) are produced independently, to have the possibility to have separate .dwp files for each (gfxNNNN, flags) tuple. Otherwise, we would need the ability to merge multiple .dwp files into one, I don't think we have a tool for that today.

Can a binary contian multiple GPU embedded ELF files? Or will there only ever be one? Like, can an executable contain 2 ELF files for gfx942, and 2 ELF files for gfx801?

Yes it can. My test case has two, it is generated like this:

/opt/rocm/bin/hipcc -O0 -g3 -gsplit-dwarf --offload-arch=gfx1010 --offload-arch=gfx90a -c main.hip -o main.o
/opt/rocm/bin/hipcc -O0 -g3 -gsplit-dwarf --offload-arch=gfx1010 --offload-arch=gfx90a -c kernel_add.hip -o kernel_add.o
/opt/rocm/bin/hipcc -O0 -g3 -gsplit-dwarf --offload-arch=gfx1010 --offload-arch=gfx90a -c kernel_scale.hip -o kernel_scale.o
/opt/rocm/bin/hipcc -g3 -gsplit-dwarf main.o kernel_add.o kernel_scale.o -o rocm-dwp-test

When I extract the resulting binary:

$ llvm-objdump --offloading rocm-dwp-test                        

rocm-dwp-test:  file format elf64-x86-64
Extracting offload bundle: rocm-dwp-test.0.host-x86_64-unknown-linux-gnu-
Extracting offload bundle: rocm-dwp-test.0.hipv4-amdgcn-amd-amdhsa--gfx1010
Extracting offload bundle: rocm-dwp-test.0.hipv4-amdgcn-amd-amdhsa--gfx90a
Extracting offload bundle: rocm-dwp-test.1.host-x86_64-unknown-linux-gnu-
Extracting offload bundle: rocm-dwp-test.1.hipv4-amdgcn-amd-amdhsa--gfx1010
Extracting offload bundle: rocm-dwp-test.1.hipv4-amdgcn-amd-amdhsa--gfx90a

The .hip_fatbin section consists of two independent bundles concatenated, one for kernel_add and one for kernel_scale. Each bundle contains code objects for the 3 arches (when we include the host).

If we used the naming scheme described above, this would imply that the .dwo files for both gfx1010 programs would have to be in the same .dwp file.

See this proposal [DWARF Package file (.dwp) .debug_dwp ID] that has been accepted for DWARF 6:
https://dwarfstd.org/issues/260116.1.html
I think that could be useful for us, to match ELF files vs DWP files.

Yes, this seems to be a good way to match things up. We should use this. Then the contained ELF files can have a .debug_dwp section that points to the right .dwp file for the right GPU and the name doesn't matter as it will be specified.

Yeah, it seems to me like it would be good to implement .debug_dwp as the preferred method of finding the .dwp files. We could still implement a name-based approach as a fallback if we find that we need to.

I think that llvm-dwp would be a good candidate for handling all of this. I think it could handle:

• Recursing into bundled code objects to create .dwp files for the top-level binary as well as for all bundled code objects. It could have options to create one .dwp file per architecture (per bundle entry ID), or one .dwp per ELF file, whatever we want.
• Inserting a .debug_dwp section in all of those files.

It would be even seem better if packaging the .dwp file and inserting the .debug_dwp section could be done by the toolchain when you link, and not as a post processing step, but that would be a bigger change.