diff options
author | Jakub Jelinek <jakub@redhat.com> | 2018-07-26 18:12:02 +0200 |
---|---|---|
committer | Jakub Jelinek <jakub@gcc.gnu.org> | 2018-07-26 18:12:02 +0200 |
commit | 5883c5ccc9339a2cbce7754179bed7279d98d4aa (patch) | |
tree | aeea1b4aee66b98c517921562febedce95453201 | |
parent | c83b4b824214039fea696083e6a888aa7c9063ce (diff) | |
download | gcc-5883c5ccc9339a2cbce7754179bed7279d98d4aa.zip gcc-5883c5ccc9339a2cbce7754179bed7279d98d4aa.tar.gz gcc-5883c5ccc9339a2cbce7754179bed7279d98d4aa.tar.bz2 |
re PR middle-end/86660 (libgomp.c++/for-15.C ICEs with nvptx offloading)
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: r263010
-rw-r--r-- | gcc/ChangeLog | 7 | ||||
-rw-r--r-- | gcc/omp-low.c | 5 | ||||
-rw-r--r-- | libgomp/ChangeLog | 5 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.c/pr86660.c | 28 |
4 files changed, 44 insertions, 1 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d9e8e10..ceaa404 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,10 @@ +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-07-26 Martin Liska <mliska@suse.cz> PR lto/86548 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 714490d..843c66f 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -1183,13 +1183,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 27ff705..0a2a051 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2018-07-26 Jakub Jelinek <jakub@redhat.com> + + PR middle-end/86660 + * testsuite/libgomp.c/pr86660.c: New test. + 2018-07-26 Cesar Philippidis <cesar@codesourcery.com> Tom de Vries <tdevries@suse.de> diff --git a/libgomp/testsuite/libgomp.c/pr86660.c b/libgomp/testsuite/libgomp.c/pr86660.c new file mode 100644 index 0000000..bea6b15 --- /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; +} |