diff options
| author | Jakub Jelinek <jakub@redhat.com> | 2018-10-12 19:28:51 +0200 |
|---|---|---|
| committer | Jakub Jelinek <jakub@gcc.gnu.org> | 2018-10-12 19:28:51 +0200 |
| commit | 37bc00e7f05602ffbee267693a522b1fda7f5212 (patch) | |
| tree | 93dc641dcb93b279219ec94f11ed96255d8469f0 | |
| parent | 30de465eef6d3ac3bd08638305ebc508a3243203 (diff) | |
backport: re PR middle-end/86660 (libgomp.c++/for-15.C ICEs with nvptx offloading)
Backported from mainline
2018-07-26 Jakub Jelinek <jakub@redhat.com>
PR middle-end/86660
* omp-low.c (scan_sharing_clauses): Don't ignore map clauses for
declare target to variables if they have always,{to,from,tofrom} map
kinds.
* testsuite/libgomp.c/pr86660.c: New test.
From-SVN: r265116
| -rw-r--r-- | gcc/ChangeLog | 10 | ||||
| -rw-r--r-- | gcc/omp-low.c | 5 | ||||
| -rw-r--r-- | libgomp/ChangeLog | 8 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.c/pr86660.c | 28 |
4 files changed, 50 insertions, 1 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 911c49b31eb..ebb7f2c19e1 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,13 @@ +2018-10-12 Jakub Jelinek <jakub@redhat.com> + + Backported from mainline + 2018-07-26 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/86660 + * omp-low.c (scan_sharing_clauses): Don't ignore map clauses for + declare target to variables if they have always,{to,from,tofrom} map + kinds. + 2018-10-12 Richard Biener <rguenther@suse.de> PR c++/54278 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index e92f53a59e7..04cbf741667 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2063,13 +2063,16 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, /* Global variables with "omp declare target" attribute don't need to be copied, the receiver side will use them directly. However, global variables with "omp declare target link" - attribute need to be copied. */ + attribute need to be copied. Or when ALWAYS modifier is used. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && DECL_P (decl) && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE)) || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx)) && varpool_node::get_create (decl)->offloadable && !lookup_attribute ("omp declare target link", diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index a1e2a3cb179..a9bbc5aaa56 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,11 @@ +2018-10-12 Jakub Jelinek <jakub@redhat.com> + + Backported from mainline + 2018-07-26 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/86660 + * testsuite/libgomp.c/pr86660.c: New test. + 2018-06-26 Jakub Jelinek <jakub@redhat.com> PR c++/86291 diff --git a/libgomp/testsuite/libgomp.c/pr86660.c b/libgomp/testsuite/libgomp.c/pr86660.c new file mode 100644 index 00000000000..bea6b15270b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr86660.c @@ -0,0 +1,28 @@ +/* PR middle-end/86660 */ + +#pragma omp declare target +int v[20]; + +void +foo (void) +{ + if (v[7] != 2) + __builtin_abort (); + v[7] = 1; +} +#pragma omp end declare target + +int +main () +{ + v[5] = 8; + v[7] = 2; + #pragma omp target map (always, tofrom: v) + { + foo (); + v[5] = 3; + } + if (v[7] != 1 || v[5] != 3) + __builtin_abort (); + return 0; +} |