gcc-14 offload support for amd gfx103x and gfx110x gpus is broken

Bug #2060619 reported by oscarbg
8
This bug affects 1 person
Affects Status Importance Assigned to Milestone
gcc-14 (Ubuntu)
Fix Released
Undecided
Unassigned

Bug Description

Hi,

for context bug was posted here first:
https://github.com/gcc-mirror/gcc/commit/1bf18629c54adf4893c8db5227a36e1952ee69a3#commitcomment-140648051

gcc Andrew Stubbs author mentions:
"Yes, it looks like the libgomp device-side library is missing. Probably not listed in the packaging script?"

using gcc-14:
gcc-14 --version
gcc-14 (Ubuntu 14-20240330-1ubuntu2) 14.0.1 20240330 (experimental) [master r14-9728-g6fc84f680d0]

using the code posted below as sum.c

testing building with offload support on gfx1036:

gcc-14 -fopenmp -foffload=amdgcn-amdhsa -foffload-options=-march=gfx1036 -fcf-protection=none sum.c -o sumamd1030.exe

fails with:
ld: error: incompatible mach: /usr/lib/gcc/x86_64-linux-gnu/14/accel/amdgcn-amdhsa/libgomp.a(parallel.o)
collect2: error: ld returned 1 exit status
gcn mkoffload: fatal error: x86_64-linux-gnu-accel-amdgcn-amdhsa-gcc-14 returned 1 exit status
compilation terminated.
lto-wrapper: fatal error: /usr/libexec/gcc/x86_64-linux-gnu/14//accel/amdgcn-amdhsa/mkoffload returned 1 exit status
compilation terminated.
/usr/bin/ld: error: lto-wrapper failed
collect2: error: ld returned 1 exit status

instead selecting gfx900 works:
gcc-14 -fopenmp -foffload=amdgcn-amdhsa -foffload-options=-march=gfx900 -fcf-protection=none sum.c -o sumamd900.exe

note gfx103x and gfx11xx targets are new to GCC-14..

inspecting files seems the error is due to package
gcc-14-offload-amdgcn_14-20240330-1ubuntu2_amd64.deb
missing libgomp.so files in
\usr\lib\gcc\x86_64-linux-gnu\14\accel\amdgcn-amdhsa\gfx1036 folder
but having this file in gfx9xx folders for ex:
\usr\lib\gcc\x86_64-linux-gnu\14\accel\amdgcn-amdhsa\gfx906 folder

code used:
sum.c:
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <omp.h>

int main(int argc, char *argv[]) {
 long long int j, largeN;
 double sum = 0.0;
 clock_t tstart, tend;
 int num_dev;

 num_dev = omp_get_num_devices();

 printf("%d\n", num_dev);

 largeN = (long long int) atoll(argv[1]);

 printf("Large Number = %lld\n", largeN);

 tstart = clock();
#pragma omp parallel for private(j) reduction(+:sum)
 for(j=1; j<largeN; j++) {
  sum += (double) (1.0/j);
 }
 tend = clock();
 printf("Sum [1, ... , %lld] = %.6f\n", largeN, sum);
 printf("\nCompleted in %.2f seconds%45s\n\n", ( (double) (tend - tstart) ) / CLOCKS_PER_SEC, "");

 sum = 0.0;
 tstart = clock();
#pragma omp target map(to:largeN) map(tofrom:sum)
#pragma omp parallel for simd private(j) reduction(+:sum)
 for (j=1; j<largeN; j++) {
  sum += (double) (1.0/j);
 }
 tend = clock();
 printf("Sum [1, ... , %lld] = %.6f\n", largeN, sum);
 printf("\nCompleted in %.2f seconds%45s\n\n", ( (double) (tend - tstart) ) / CLOCKS_PER_SEC, "");

 return(0);
}

Revision history for this message
oscarbg (rtfss1) wrote :

so seems fedora 40 has packaged it right

https://packages.fedoraproject.org/pkgs/gcc/gcc-offload-nvptx/fedora-40.html#provides

x86_64-redhat-linux/14/accel/amdgcn-amdhsa/gfx1036/

    crt0.o
    libc.a
    libcaf_single.a
    libg.a
    libgcc.a
    libgcov.a
    libgfortran.a
    libgfortran.spec
    libgomp.a
    libgomp.spec
    libm.a
    libssp.a
    libssp_nonshared.a

Revision history for this message
Thomas Schwinge (tschwinge) wrote :

(<https://packages.fedoraproject.org/pkgs/gcc/gcc-offload-amdgcn/fedora-40.html#provides> you mean.)

---

I see in <https://git.launchpad.net/ubuntu/+source/gcc-14/tree/debian/rules2>, in the '$(install_gcn_stamp)' target:

    for d in . gfx906 gfx908 gfx90a; do \

I suspect that hard-coded list of GCN multilib directories is the reason why the 'gfx1036' one is missing, and the compiler thus defaults to '.' (a proper error would be appropriate for this incompatible-ISA setting), and thus the link fails.

Revision history for this message
Launchpad Janitor (janitor) wrote :

This bug was fixed in the package gcc-14 - 14-20240412-0ubuntu1

---------------
gcc-14 (14-20240412-0ubuntu1) noble; urgency=medium

  * Snapshot, taken from the trunk 20240412.

  * Apply proposed patch for PR libquadmath/114533 (Simon Chopin).
    Closes: #1064426. LP: #2052929.
  * Don't build the libstdc++-doc package when the nodoc profile is enabled.
  * Clean up some version restrictions on build-dependencies.
  * Don't set the _FILE_OFFSET_BITS and _TIME_BITS macros, if the
    _DISTRO_EVADE_TIME_BITS macro is defined.
    See https://sourceware.org/bugzilla/show_bug.cgi?id=31624.
  * Build-depend on amdgcn-tools-18, required for new gcn offload targets.
  * Package the new gfx1030, gfx1036, gfx1100 and gfx1103 offload targets.
    LP: #2060619.

 -- Matthias Klose <email address hidden> Fri, 12 Apr 2024 14:30:14 +0200

Changed in gcc-14 (Ubuntu):
status: New → Fix Released
To post a comment you must log in.
This report contains Public information  
Everyone can see this information.

Other bug subscribers

Remote bug watches

Bug watches keep track of this bug in other bug trackers.