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 911c49b31eb80ec26b8855650f535cc192de8282..ebb7f2c19e10ea5648c7396db688278e618eaf7e 100644 (file)
--- 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 e92f53a59e72205b2bfb8cd91a4a6a4695b94a7b..04cbf7416670b32e82d8e940e5df8d547a9c1944 100644 (file)
--- 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 a1e2a3cb179657004364baad29462b2dc9807357..a9bbc5aaa56f5205f98b0fabf92a281deaf4a5ac 100644 (file)
--- 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 (file)
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;
+}