From: Jakub Jelinek <jakub@redhat.com>
Date: Fri, 12 Oct 2018 17:28:51 +0000 (+0200)
Subject: backport: re PR middle-end/86660 (libgomp.c++/for-15.C ICEs with nvptx offloading)
X-Git-Tag: releases/gcc-6.5.0~31
X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=37bc00e7f05602ffbee267693a522b1fda7f5212;p=thirdparty%2Fgcc.git

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
---

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 911c49b31eb8..ebb7f2c19e10 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 e92f53a59e72..04cbf7416670 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 a1e2a3cb1796..a9bbc5aaa56f 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 000000000000..bea6b15270bd
--- /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;
+}