diff options
author | vries <vries@138bc75d-0d04-0410-961f-82ee72b054a4> | 2018-12-14 13:48:56 +0000 |
---|---|---|
committer | vries <vries@138bc75d-0d04-0410-961f-82ee72b054a4> | 2018-12-14 13:48:56 +0000 |
commit | 07c14f81e79fe4e4dfbd1176bcf10937f1fdf38f (patch) | |
tree | 2d768b42e3a31f9210ad81e2e41b2466bd284d01 /libgomp/testsuite | |
parent | 76c21b271247ccbd681bdb4530426d2fe35dbfa5 (diff) |
[offloading] Error on missing symbols
When compiling an OpenMP or OpenACC program containing a reference in the
offloaded code to a symbol that has not been included in the offloaded code,
the offloading compiler may ICE in lto1.
Fix this by erroring out instead, mentioning the problematic symbol:
...
error: variable 'var' has been referenced in offloaded code but hasn't
been marked to be included in the offloaded code
lto1: fatal error: errors during merging of translation units
compilation terminated.
...
Build x86_64 with nvptx accelerator and reg-tested libgomp.
Build x86_64 and reg-tested libgomp.
2018-12-14 Tom de Vries <tdevries@suse.de>
* lto-cgraph.c (verify_node_partition): New function.
(input_overwrite_node, input_varpool_node): Use verify_node_partition.
* testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test.
* testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test.
* testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test.
* testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test.
* testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@267134 138bc75d-0d04-0410-961f-82ee72b054a4
Diffstat (limited to 'libgomp/testsuite')
5 files changed, 82 insertions, 0 deletions
diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c new file mode 100644 index 00000000000..b8aa3da48a1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c @@ -0,0 +1,12 @@ +/* { dg-skip-if "" { *-*-* } } */ + +#pragma omp declare target +extern int var; +#pragma omp end declare target + +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} + diff --git a/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c new file mode 100644 index 00000000000..9e59ef8864e --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/function-not-offloaded.c @@ -0,0 +1,16 @@ +/* { dg-do link } */ +/* { dg-excess-errors "unresolved symbol foo, lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */ +/* { dg-additional-sources "function-not-offloaded-aux.c" } */ + +#pragma omp declare target +int var; +#pragma omp end declare target + +extern void foo (); + +int +main () +{ +#pragma omp target + foo (); +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c new file mode 100644 index 00000000000..bc4b916e9a4 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/variable-not-offloaded.c @@ -0,0 +1,19 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target offload_device_nonshared_as } } */ + +int var; /* { dg-error "variable 'var' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target offload_device_nonshared_as } } */ + +#pragma omp declare target +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} +#pragma omp end declare target + +int +main () +{ +#pragma omp target + foo (); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c new file mode 100644 index 00000000000..c94f268462f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c @@ -0,0 +1,18 @@ +/* { dg-do link } */ +/* { dg-excess-errors "lto1, mkoffload and lto-wrapper fatal errors" { target openacc_nvidia_accel_configured } } */ + +int var; +#pragma acc declare create (var) + +void __attribute__((noinline, noclone)) +foo () /* { dg-error "function 'foo' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code" "" { target openacc_nvidia_accel_configured } } */ +{ + var++; +} + +int +main () +{ +#pragma acc parallel + foo (); +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c new file mode 100644 index 00000000000..8e10271b97f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c @@ -0,0 +1,17 @@ +/* { dg-do link } */ + +int var; /* { dg-error "'var' requires a 'declare' directive for use in a 'routine' function" } */ + +#pragma acc routine +void __attribute__((noinline, noclone)) +foo (void) +{ + var++; +} + +int +main () +{ +#pragma acc parallel + foo (); +} |