This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[committed] Don't ignore OpenMP map clauses for declare target to vars if there is always modifier (PR middle-end/86660)


Hi!

We can't ignore map clauses for variables that are declare target to
if the map clause has always modifier, because we need to copy the data to
and/or from the device as user requested.  This has slightly undesirable
effect that the vars inside of the construct are remapped, perhaps with more
work we could keep using the device var directly.

Bootstrapped/regtested on x86_64-linux and i686-linux, committed to trunk
and 8.3.

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.

--- gcc/omp-low.c.jj	2018-07-17 12:54:13.543991017 +0200
+++ gcc/omp-low.c	2018-07-26 13:43:08.453714154 +0200
@@ -1183,13 +1183,16 @@ scan_sharing_clauses (tree clauses, omp_
 	  /* 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/testsuite/libgomp.c/pr86660.c.jj	2018-07-26 14:37:49.269202230 +0200
+++ libgomp/testsuite/libgomp.c/pr86660.c	2018-07-26 14:35:31.938979831 +0200
@@ -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;
+}

	Jakub


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]