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

+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

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

libgomp/ChangeLog

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

libgomp/testsuite/libgomp.c/pr86660.c

+/* 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;
+}