From 9bc3b95dfefd37d860c5dc0004f8a53f6290fbb1 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Sun, 9 Feb 2020 08:17:10 +0100 Subject: openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regions DECL_IN_CONSTANT_POOL are shared and thus don't really get emitted in the BLOCK where they are used, so for OpenMP target regions that have initializers gimplified into copying from them we actually map them at runtime from host to offload devices. This patch instead marks them as "omp declare target", so that they are on the target device from the beginning and don't need to be copied there. 2020-02-09 Jakub Jelinek * gimplify.c (gimplify_adjust_omp_clauses_1): Promote DECL_IN_CONSTANT_POOL variables into "omp declare target" to avoid copying them around between host and target. * testsuite/libgomp.c/target-38.c: New test. --- libgomp/testsuite/libgomp.c/target-38.c | 28 ++++++++++++++++++++++++++++ 1 file changed, 28 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/target-38.c (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/target-38.c b/libgomp/testsuite/libgomp.c/target-38.c new file mode 100644 index 0000000..8169972 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-38.c @@ -0,0 +1,28 @@ +#define A(n) n##0, n##1, n##2, n##3, n##4, n##5, n##6, n##7, n##8, n##9 +#define B(n) A(n##0), A(n##1), A(n##2), A(n##3), A(n##4), A(n##5), A(n##6), A(n##7), A(n##8), A(n##9) + +int +foo (int x) +{ + int b[] = { B(4), B(5), B(6) }; + return b[x]; +} + +int v[] = { 1, 2, 3, 4, 5, 6 }; +#pragma omp declare target to (foo, v) + +int +main () +{ + int i = 5; + asm ("" : "+g" (i)); + #pragma omp target map(tofrom:i) + { + int a[] = { B(1), B(2), B(3) }; + asm ("" : : "m" (a) : "memory"); + i = a[i] + foo (i) + v[i & 63]; + } + if (i != 105 + 405 + 6) + __builtin_abort (); + return 0; +} -- cgit v1.1 From fd789c816b06235b04698636db69e302b24c83ba Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Mon, 10 Feb 2020 08:08:00 +0100 Subject: Add xfails to libgomp tests target-{33,34}.c, target-link-1.c Add xfails for nvptx offloading because "no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c" (https://gcc.gnu.org/PR81688) and because "omp target link not implemented for nvptx" (https://gcc.gnu.org/PR81689). libgomp/ * testsuite/libgomp.c/target-33.c: Add xfail for execution on offload_target_nvptx, cf. https://gcc.gnu.org/PR81688. * testsuite/libgomp.c/target-34.c: Likewise. * testsuite/libgomp.c/target-link-1.c: Add xfail for offload_target_nvptx, cf. https://gcc.gnu.org/PR81689. --- libgomp/testsuite/libgomp.c/target-33.c | 3 +++ libgomp/testsuite/libgomp.c/target-34.c | 3 +++ libgomp/testsuite/libgomp.c/target-link-1.c | 3 +++ 3 files changed, 9 insertions(+) (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c index 1bed4b6..15d2d7e 100644 --- a/libgomp/testsuite/libgomp.c/target-33.c +++ b/libgomp/testsuite/libgomp.c/target-33.c @@ -1,3 +1,6 @@ +/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81688. */ + extern void abort (void); int diff --git a/libgomp/testsuite/libgomp.c/target-34.c b/libgomp/testsuite/libgomp.c/target-34.c index 66d9f54..5a35964 100644 --- a/libgomp/testsuite/libgomp.c/target-34.c +++ b/libgomp/testsuite/libgomp.c/target-34.c @@ -1,3 +1,6 @@ +/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81688. */ + extern void abort (void); int diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c index 681677c..99ce33b 100644 --- a/libgomp/testsuite/libgomp.c/target-link-1.c +++ b/libgomp/testsuite/libgomp.c/target-link-1.c @@ -1,3 +1,6 @@ +/* { dg-xfail-if "#pragma omp target link not implemented" { offload_target_nvptx } } + Cf. https://gcc.gnu.org/PR81689. */ + struct S { int s, t; }; int a = 1, b = 1; -- cgit v1.1 From 001ab12e620c6f117b2e93c77d188bd62fe7ba03 Mon Sep 17 00:00:00 2001 From: Frederik Harwath Date: Thu, 13 Feb 2020 07:30:16 +0100 Subject: openmp: ignore nowait if async execution is unsupported [PR93481] An OpenMP "nowait" clause on a target construct currently leads to a call to GOMP_OFFLOAD_async_run in the plugin that is used for offloading at execution time. The nvptx plugin contains only a stub of this function that always produces a fatal error if called. This commit changes the "nowait" implementation to ignore the clause if the executing device's plugin does not implement GOMP_OFFLOAD_async_run. The stub in the nvptx plugin is removed which effectively means that programs containing "nowait" can now be executed with nvptx offloading as if the clause had not been used. This behavior is consistent with the OpenMP specification which says that "[...] execution of the target task *may* be deferred" (emphasis added), cf. OpenMP 5.0, page 172. libgomp/ * plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub. * target.c (gomp_load_plugin_for_device): Make "async_run" loading optional. (gomp_target_task_fn): Assert "devicep->async_run_func". (clear_unsupported_flags): New function to remove unsupported flags (right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored. (GOMP_target_ext): Apply clear_unsupported_flags to flags. * testsuite/libgomp.c/target-33.c: Remove xfail for offload_target_nvptx. * testsuite/libgomp.c/target-34.c: Likewise. --- libgomp/testsuite/libgomp.c/target-33.c | 3 --- libgomp/testsuite/libgomp.c/target-34.c | 3 --- 2 files changed, 6 deletions(-) (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/target-33.c b/libgomp/testsuite/libgomp.c/target-33.c index 15d2d7e..1bed4b6 100644 --- a/libgomp/testsuite/libgomp.c/target-33.c +++ b/libgomp/testsuite/libgomp.c/target-33.c @@ -1,6 +1,3 @@ -/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } - Cf. https://gcc.gnu.org/PR81688. */ - extern void abort (void); int diff --git a/libgomp/testsuite/libgomp.c/target-34.c b/libgomp/testsuite/libgomp.c/target-34.c index 5a35964..66d9f54 100644 --- a/libgomp/testsuite/libgomp.c/target-34.c +++ b/libgomp/testsuite/libgomp.c/target-34.c @@ -1,6 +1,3 @@ -/* { dg-xfail-run-if "GOMP_OFFLOAD_async_run not implemented" { offload_target_nvptx } } - Cf. https://gcc.gnu.org/PR81688. */ - extern void abort (void); int -- cgit v1.1 From 9c3cdb43c2bdaf8a8d2e62db010b04f6086d76b7 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Sun, 15 Mar 2020 01:27:40 +0100 Subject: tree-nested: Fix handling of *reduction clauses with C array sections [PR93566] tree-nested.c didn't handle C array sections in {,task_,in_}reduction clauses. 2020-03-14 Jakub Jelinek PR middle-end/93566 * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Handle {,in_,task_}reduction clauses with C/C++ array sections. * testsuite/libgomp.c/pr93566.c: New test. --- libgomp/testsuite/libgomp.c/pr93566.c | 113 ++++++++++++++++++++++++++++++++++ 1 file changed, 113 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/pr93566.c (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/pr93566.c b/libgomp/testsuite/libgomp.c/pr93566.c new file mode 100644 index 0000000..3334bd57 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr93566.c @@ -0,0 +1,113 @@ +/* PR middle-end/93566 */ +/* { dg-additional-options "-std=c99" } */ + +extern void abort (void); + +void +foo (int *x) +{ + void nest (void) { + #pragma omp parallel for reduction(+:x[:10]) + for (int i = 0; i < 1024; i++) + for (int j = 0; j < 10; j++) + x[j] += j * i; + } + nest (); + for (int i = 0; i < 10; i++) + if (x[i] != 1023 * 1024 / 2 * i) + abort (); +} + +void +bar (void) +{ + int x[10] = {}; + void nest (void) { + #pragma omp parallel for reduction(+:x[:10]) + for (int i = 0; i < 1024; i++) + for (int j = 0; j < 10; j++) + x[j] += j * i; + } + nest (); + for (int i = 0; i < 10; i++) + if (x[i] != 1023 * 1024 / 2 * i) + abort (); +} + +void +baz (void) +{ + int x[10] = {}; + void nest (void) { + #pragma omp parallel for reduction(+:x[2:5]) + for (int i = 0; i < 1024; i++) + for (int j = 2; j < 7; j++) + x[j] += j * i; + } + nest (); + for (int i = 2; i < 7; i++) + if (x[i] != 1023 * 1024 / 2 * i) + abort (); +} + +void +qux (int *x) +{ + void nest (void) { x++; } + nest (); + #pragma omp parallel for reduction(+:x[:9]) + for (int i = 0; i < 1024; i++) + for (int j = 0; j < 9; j++) + x[j] += j * i; + nest (); + for (int i = 0; i < 9; i++) + if (x[i - 1] != 1023 * 1024 / 2 * i) + abort (); +} + +void +quux (void) +{ + int x[10]; + void nest (void) { for (int i = 0; i < 10; i++) x[i] = 0; } + int nest2 (int i) { return x[i]; } + nest (); + #pragma omp parallel for reduction(+:x[:7]) + for (int i = 0; i < 1024; i++) + for (int j = 0; j < 7; j++) + x[j] += j * i; + for (int i = 0; i < 7; i++) + if (nest2 (i) != 1023 * 1024 / 2 * i) + abort (); +} + +void +corge (void) +{ + int x[10]; + void nest (void) { for (int i = 0; i < 10; i++) x[i] = 0; } + int nest2 (int i) { return x[i]; } + nest (); + #pragma omp parallel for reduction(+:x[2:4]) + for (int i = 0; i < 1024; i++) + for (int j = 2; j < 6; j++) + x[j] += j * i; + for (int i = 2; i < 6; i++) + if (nest2 (i) != 1023 * 1024 / 2 * i) + abort (); +} + +int +main () +{ + int a[10] = {}; + foo (a); + bar (); + baz (); + for (int i = 0; i < 10; i++) + a[i] = 0; + qux (a); + quux (); + corge (); + return 0; +} -- cgit v1.1 From c2211a60ff05b7a0289d3e287e72c181bb4d5d8b Mon Sep 17 00:00:00 2001 From: Tobias Burnus Date: Tue, 24 Mar 2020 15:13:56 +0100 Subject: Fix OpenMP offload handling for target-link variables for nvptx (PR81689) PR libgomp/81689 * lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state. PR libgomp/81689 * omp-offload.c (omp_finish_file): Fix target-link handling if targetm_common.have_named_sections is false. PR libgomp/81689 * testsuite/libgomp.c/target-link-1.c: Remove xfail. --- libgomp/testsuite/libgomp.c/target-link-1.c | 3 --- 1 file changed, 3 deletions(-) (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/target-link-1.c b/libgomp/testsuite/libgomp.c/target-link-1.c index 99ce33b..681677c 100644 --- a/libgomp/testsuite/libgomp.c/target-link-1.c +++ b/libgomp/testsuite/libgomp.c/target-link-1.c @@ -1,6 +1,3 @@ -/* { dg-xfail-if "#pragma omp target link not implemented" { offload_target_nvptx } } - Cf. https://gcc.gnu.org/PR81689. */ - struct S { int s, t; }; int a = 1, b = 1; -- cgit v1.1 From dc703151d4f4560e647649506d5b4ceb0ee11e90 Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Tue, 12 May 2020 09:17:09 +0200 Subject: openmp: Implement discovery of implicit declare target to clauses This attempts to implement what the OpenMP 5.0 spec in declare target section says as ammended by the 5.1 changes so far (related to device_type(host)), except that it doesn't have the device(ancestor: ...) handling yet because we do not support it yet, and I've left so far out the except lambda note, because I need that clarified. 2020-05-12 Jakub Jelinek * omp-offload.h (omp_discover_implicit_declare_target): Declare. * omp-offload.c: Include context.h. (omp_declare_target_fn_p, omp_declare_target_var_p, omp_discover_declare_target_fn_r, omp_discover_declare_target_var_r, omp_discover_implicit_declare_target): New functions. * cgraphunit.c (analyze_functions): Call omp_discover_implicit_declare_target. * testsuite/libgomp.c/target-39.c: New test. --- libgomp/testsuite/libgomp.c/target-39.c | 47 +++++++++++++++++++++++++++++++++ 1 file changed, 47 insertions(+) create mode 100644 libgomp/testsuite/libgomp.c/target-39.c (limited to 'libgomp/testsuite/libgomp.c') diff --git a/libgomp/testsuite/libgomp.c/target-39.c b/libgomp/testsuite/libgomp.c/target-39.c new file mode 100644 index 0000000..4442f43 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-39.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-options "-O0" } */ + +extern void abort (void); +volatile int v; +#pragma omp declare target to (v) +typedef void (*fnp1) (void); +typedef fnp1 (*fnp2) (void); +void f1 (void) { v++; } +void f2 (void) { v += 4; } +void f3 (void) { v += 16; f1 (); } +fnp1 f4 (void) { v += 64; return f2; } +int a = 1; +int *b = &a; +int **c = &b; +fnp2 f5 (void) { f3 (); return f4; } +#pragma omp declare target to (c, f5) + +int +main () +{ + int err = 0; + #pragma omp target map(from:err) + { + volatile int xa; + int *volatile xb; + int **volatile xc; + fnp2 xd; + fnp1 xe; + err = 0; + xa = a; + err |= xa != 1; + xb = b; + err |= xb != &a; + xc = c; + err |= xc != &b; + xd = f5 (); + err |= v != 17; + xe = xd (); + err |= v != 81; + xe (); + err |= v != 85; + } + if (err) + abort (); + return 0; +} -- cgit v1.1