[Patch][gcn, nvptx, offloading] mkoffload – handle -fpic/-fPIC

[Thomas Schwinge]

Hi Kwok!

On 2020-07-08T18:35:43+0100, Kwok Cheung Yeung <kcy@codesourcery.com> wrote:
>  > I tried out the patch with one test-case and -pie -fPIC/-fpic already
>  > seems to works, so perhaps we could have at least one test-case
>  > exercising this in libgomp?  That sounds easier to do than the
>  > shared-lib test-case.
>
> I've created a simple testcase which tries to generate a shared library with
> offloaded code.

Thanks.

> Without the mkoffload patch, it fails during linking with:
>
> ld: /tmp/ccNaT7fO.target.o: relocation R_X86_64_32 against `.rodata' can not be
> used when making a shared object; recompile with -fPIC
>
> I have tested this on a x64 host with offloading to nvptx and gcn.

Confirmed, and also that it actually works; using the two files attached:

    $ install/bin/gcc -Wall -Wextra -fopenacc -o libf.so shared-lib.c -shared -fPIC
    $ install/bin/gcc -Wall -Wextra -fopenacc shared-lib-main.c -Wl,-rpath,$PWD -L$PWD -lf

..., and execute 'a.out' on a GCN and a nvptx GPU machine.

(As discussed before, it's non-trivial to get such a two-step compilation
process into the libgomp testsuite, so not doing that now.)

> On AMD GCN,
> it also produces a couple of extra linker warnings that I have added dg-warning
> entries for.

Isn't that unexpected to see such warnings?  (Would you declare this
issue "done" if users then see such warnings?)  That said, I'm not seeing
these, thus get:

    FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/shared-lib.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa  -O0   (test for warnings, line )
    FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/shared-lib.c -DACC_DEVICE_TYPE_radeon=1 -DACC_MEM_SHARED=0 -foffload=amdgcn-amdhsa  -O0   (test for warnings, line )

> Okay for trunk/OG10 together with the previous mkoffload patch?

In principle yes, with the 'dg-warning' directives removed.  Please
verify why you're seeing these warnings with GCN 'mkoffload'.

> commit 43238117c261285a6b95d881bcc2f9efd9f752ad
> Author: Kwok Cheung Yeung <kcy@codesourcery.com>
> Date:   Wed Jul 8 03:28:08 2020 -0700
>
>     Add test case

Try to make 'git log --oneline' meaningful, so a bit more verbose,
please.  Maybe: "Add test case 'libgomp.oacc-c-c++-common/shared-lib.c'"?

> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/shared-lib.c
> @@ -0,0 +1,16 @@

Maybe add a header line comment, like: "Verify that we can compile
offloading code into a shared library".

> +/* { dg-do link } */
> +/* { dg-additional-options "-shared -fPIC" { target fpic } } */

Let's hope this does the expected thing for all offloading/libgomp
testing configurations.  Not sure if (additionally?) we should
'dg-require-effective-target shared'?  See 'g++.dg/tls/pr85400.C', for
example.  (Jakub?)

> +
> +#define N 512
> +
> +void f(int a[])
> +{
> +  int i;
> +
> +  #pragma acc parallel
> +  for (i = 0; i < N; i++)
> +    a[i]++;
> +}
> +
> +/* { dg-warning "relocation against `.*' in read-only section `\.rodata'" "" { target openacc_radeon_accel_selected } 0 } */
> +/* { dg-warning "creating DT_TEXTREL in a shared object" "" { target openacc_radeon_accel_selected } 0 } */


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
-------------- next part --------------
A non-text attachment was scrubbed...
Name: shared-lib.c
Type: text/x-csrc
Size: 211 bytes
Desc: not available
URL: <https://gcc.gnu.org/pipermail/gcc-patches/attachments/20200709/e04efbe7/attachment.bin>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: shared-lib-main.c
Type: text/x-csrc
Size: 174 bytes
Desc: not available
URL: <https://gcc.gnu.org/pipermail/gcc-patches/attachments/20200709/e04efbe7/attachment-0001.bin>