MOLD removes GCC's offload symbols when linking

[tob2]

That's with a GCC compiler with enabled offload support.

For instance:

• For building manually, see https://gcc.gnu.org/wiki/Offloading
• On openSUSE/SUSE, install additionally cross-{nvptx,amdgcn}-gcc13 – newer GCC can be also found at https://build.opensuse.org/project/show/devel:gcc
• On Debian/Ubuntu, install additionally gcc-13-offload-{nvptx,amdgcn} – on Ubuntu, newer GCC can be also found at https://launchpad.net/~ubuntu-toolchain-r/+archive/ubuntu/test
• On Fedora/RHEL, install additionally {gcc,libgomp}-offload-{nvptx,amdgcn} (amdgcn support since Fedora 40)

When using:

• file1.c

#pragma omp declare target
int f() {
  return 42;
}
#pragma omp end declare target

• file2.c

extern int f();

int main() {
  int x = 0;
  #pragma omp target map(x)
    x = f();
  return x != 42;
}

And then running:

gcc -fopenmp file1.c file2.c  -fuse-ld=mold -o test.mold
gcc -fopenmp file1.c file2.c  -fuse-ld=bfd -o test.bfd

The MOLD version lacks the following symbol:

    3: 0000000000000000      0 FUNC    GLOBAL DEFAULT    UNDEF GOMP_offload_unregister_ver@GOMP_4.5 (3)
    5: 0000000000000000      0 FUNC    GLOBAL DEFAULT    UNDEF GOMP_offload_register_ver@GOMP_4.5 (3)

And, hence, it fails here (I do have an Nvidia card for offloading):

# Works:
$ OMP_TARGET_OFFLOAD=mandatory ./test.bfd

# Fails:
# Works:
$ OMP_TARGET_OFFLOAD=mandatory ./test.mold

libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot be used for offloading

[rui314]

Do I need an Nvidia card to reproduce the issue? I spin up Ubuntu 23.04 docker image, installed the compiler with apt-get install gcc-13-offload-{nvptx,amdgcn}, and built your file1.c and file2.c with mold as instructed. Then, OMP_TARGET_OFFLOAD=mandatory ./test.mold finished without any errors.

[tob2]

Do I need an Nvidia card to reproduce the issue?

Yes — if you want to reproduce the run-time result.

No — to see the issue as there inspecting the resulting binaries for the symbols — or looking at GCC's lto-plugin/lto-plugin.c either in the debugger or compiling it with some some printf debugging (separate compilation should be possible).

Namely, if you have the GOMP_offload_register_ver@GOMP_4.5 symbol in the binary, it is okay, if not, there os no offloading.

Offloading (in GCC, I think also in LLVM, but I mightbe wrong) piggybacks on the LTO, using special sections like ".gnu.offload_lto_.opts" to store the data.

That's primarily handled via the LDPT_REGISTER_CLAIM_FILE_HOOK or LDPT_REGISTER_CLAIM_FILE_HOOK_V2 hook. Namely, in the above linked linker plugin, claim_file_handler_v2 checks on one hand for normal LTO sections but also for files that contain offload sections. The former are claimed", the latter only remembered.

Once done, in all_symbols_read_handler, GCC's LTO wrapper is called if either files have been claimed (normal host side LTO) or offload data has been found or both.

My impression is that GCC's LTO plugin is never called and therefore cannot do the offloading code generation. Possibly, there is a check in MOLD that there is an LTO section, but for GCC offloading, also a ckeck for some offloading section is needed. The linker plugin itself only checks for ".gnu.offload_lto_.opts" (There are more but those are then handled by the LTO compiler itself.)

For the LTO handling in GCC and the section names, see also https://gcc.gnu.org/wiki/Offloading#Address_mapping_tables

I hope that helps.

Thanks!

[tob2]

(For LLVM, seemingly ".llvm.offloading" is used according to https://clang.llvm.org/docs/OffloadingDesign.html#linking-target-device-code — but I think they don't use linker plugins, but I have not digested LLVM's workflow.)

[rui314]

Does compiling and linking with -flto fix the issue? I.e. does the executables created with the following commands work?

gcc -fopenmp file1.c file2.c  -fuse-ld=mold -o test.mold -flto
gcc -fopenmp file1.c file2.c  -fuse-ld=bfd -o test.bfd -flto

[rui314]

Could you try again with the above commit?