aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorJakub Jelinek <jakub@redhat.com>2018-07-26 18:12:02 +0200
committerJakub Jelinek <jakub@gcc.gnu.org>2018-07-26 18:12:02 +0200
commit5883c5ccc9339a2cbce7754179bed7279d98d4aa (patch)
treeaeea1b4aee66b98c517921562febedce95453201
parentc83b4b824214039fea696083e6a888aa7c9063ce (diff)
downloadgcc-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/ChangeLog7
-rw-r--r--gcc/omp-low.c5
-rw-r--r--libgomp/ChangeLog5
-rw-r--r--libgomp/testsuite/libgomp.c/pr86660.c28
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;
+}