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

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index d9e8e1083312e964355e36b76284ed0c22d32376..ceaa404c9ba9e1238541b91dc27f4454a039aa5a 100644 (file)
--- 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 714490d6921dd10ce36c2ad8a49b5954e12db889..843c66fd22199b3b77737030c5f489d7c2f1c482 100644 (file)
--- 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 27ff705e90997df6c2eb0993cf56c30403edaaa8..0a2a051c9e8aaa2e506f098640910c518571bda4 100644 (file)
--- 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 (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;
+}