From 5883c5ccc9339a2cbce7754179bed7279d98d4aa Mon Sep 17 00:00:00 2001 From: Jakub Jelinek Date: Thu, 26 Jul 2018 18:12:02 +0200 Subject: [PATCH] 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 --- gcc/ChangeLog | 7 +++++++ gcc/omp-low.c | 5 ++++- libgomp/ChangeLog | 5 +++++ libgomp/testsuite/libgomp.c/pr86660.c | 28 +++++++++++++++++++++++++++ 4 files changed, 44 insertions(+), 1 deletion(-) create mode 100644 libgomp/testsuite/libgomp.c/pr86660.c diff --git a/gcc/ChangeLog b/gcc/ChangeLog index d9e8e108331..ceaa404c9ba 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,10 @@ +2018-07-26 Jakub Jelinek + + 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 PR lto/86548 diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 714490d6921..843c66fd221 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 27ff705e909..0a2a051c9e8 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,8 @@ +2018-07-26 Jakub Jelinek + + PR middle-end/86660 + * testsuite/libgomp.c/pr86660.c: New test. + 2018-07-26 Cesar Philippidis Tom de Vries 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; +} -- 2.30.2