Hi Kwok!
On 2020-07-08T18:35:43+0100, Kwok Cheung Yeung <[email protected]> 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 <[email protected]>
> 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
/* { dg-do link } */
/* { dg-additional-options "-shared -fPIC" { target fpic } } */
#define N 512
void f(int a[])
{
int i;
#pragma acc parallel loop copy(a[0:N])
for (i = 0; i < N; i++)
a[i] = -i;
}
#include <assert.h>
extern void f(int a[]);
#define N 512
int main()
{
int a[N] = { 0 };
f(a);
for (int i = 0; i < N; i++)
assert (a[i] == -i);
return 0;
}
