gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT for VLA helper variables on target data even if...

	* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
	GOVD_EXPLICIT flags.
	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
	call install_var_field with mask 11 instead of 3.
	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
gcc/c/
	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
gcc/cp/
	* semantics.c (finish_omp_clauses): For C_ORT_OMP
	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
	instead of generic_head to track duplicates.
libgomp/
	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
	perform the lookup in the first loop only if !not_found_cnt, otherwise
	perform lookups for it in the second loop guarded with
	if (not_found_cnt || has_firstprivate).
	* testsuite/libgomp.c/target-37.c: New test.
	* testsuite/libgomp.c++/target-22.C: New test.

From-SVN: r274206

11 files changed, 287 insertions, 31 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 45050ed..00d8030 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,16 @@
+2019-08-08  Jakub Jelinek  <jakub@redhat.com>
+
+	* gimplify.c (omp_add_variable): Use GOVD_PRIVATE | GOVD_EXPLICIT
+	for VLA helper variables on target data even if not GOVD_FIRSTPRIVATE.
+	(gimplify_scan_omp_clauses): For OMP_CLAUSE_USE_DEVICE_* use just
+	GOVD_EXPLICIT flags.
+	(gimplify_omp_workshare): For OMP_TARGET_DATA move all
+	OMP_CLAUSE_USE_DEVICE_* clauses to the end of clauses chain.
+	* omp-low.c (scan_sharing_clauses): For OMP_CLAUSE_USE_DEVICE_*
+	call install_var_field with mask 11 instead of 3.
+	(lower_omp_target): For OMP_CLAUSE_USE_DEVICE_* use pass
+	(splay_tree_key) &DECL_UID (var) to build_sender_ref instead of var.
+
 2019-08-07  Richard Sandiford  <richard.sandiford@arm.com>
 
 	* config/aarch64/constraints.md (Z): Handle floating-point zeros too.
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index 18c1ce0..4944696 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,9 @@
+2019-08-08  Jakub Jelinek  <jakub@redhat.com>
+
+	* c-typeck.c (c_finish_omp_clauses): For C_ORT_OMP
+	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
+	instead of generic_head to track duplicates.
+
 2019-08-07  Jakub Jelinek  <jakub@redhat.com>
 
 	* c-parser.c (c_parser_omp_clause_name): Parse use_device_addr clause.
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index bda90fd..1d865aa 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13680,7 +13680,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead.  */
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
-  /* If ort == C_ORT_OMP used as nontemporal_head instead.  */
+  /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
+     instead.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
   if (ort & C_ORT_ACC)
@@ -14072,13 +14073,19 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (ort == C_ORT_ACC
-		   && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	  else if ((ort == C_ORT_ACC
+		    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+		   || (ort == C_ORT_OMP
+		       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+			   || (OMP_CLAUSE_CODE (c)
+			       == OMP_CLAUSE_USE_DEVICE_ADDR))))
 	    {
 	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
-			    "%qD appears more than once in reduction clauses",
+			    ort == C_ORT_ACC
+			    ? "%qD appears more than once in reduction clauses"
+			    : "%qD appears more than once in data clauses",
 			    t);
 		  remove = true;
 		}
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index bae5147..f937bf5 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2019-08-08  Jakub Jelinek  <jakub@redhat.com>
+
+	* semantics.c (finish_omp_clauses): For C_ORT_OMP
+	OMP_CLAUSE_USE_DEVICE_* clauses use oacc_reduction_head bitmap
+	instead of generic_head to track duplicates.
+
 2019-08-07  Marek Polacek  <polacek@redhat.com>
 
 	PR c++/81429 - wrong parsing of constructor with C++11 attribute.
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 8fe632f..9fd759c 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6148,7 +6148,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead.  */
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
-  /* If ort == C_ORT_OMP used as nontemporal_head instead.  */
+  /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
+     instead.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
 
   if (ort & C_ORT_ACC)
@@ -6406,13 +6407,19 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			  omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
-	  else if (ort == C_ORT_ACC
-		   && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+	  else if ((ort == C_ORT_ACC
+		    && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
+		   || (ort == C_ORT_OMP
+		       && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+			   || (OMP_CLAUSE_CODE (c)
+			       == OMP_CLAUSE_USE_DEVICE_ADDR))))
 	    {
 	      if (bitmap_bit_p (&oacc_reduction_head, DECL_UID (t)))
 		{
 		  error_at (OMP_CLAUSE_LOCATION (c),
-			    "%qD appears more than once in reduction clauses",
+			    ort == C_ORT_ACC
+			    ? "%qD appears more than once in reduction clauses"
+			    : "%qD appears more than once in data clauses",
 			    t);
 		  remove = true;
 		}
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 978df86..daa0b71 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -6932,8 +6932,10 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
 	    nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
 	  else if (flags & GOVD_PRIVATE)
 	    nflags = GOVD_PRIVATE;
-	  else if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
-		   && (flags & GOVD_FIRSTPRIVATE))
+	  else if (((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0
+		    && (flags & GOVD_FIRSTPRIVATE))
+		   || (ctx->region_type == ORT_TARGET_DATA
+		       && (flags & GOVD_DATA_SHARE_CLASS) == 0))
 	    nflags = GOVD_PRIVATE | GOVD_EXPLICIT;
 	  else
 	    nflags = GOVD_FIRSTPRIVATE;
@@ -9016,6 +9018,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
+	  flags = GOVD_EXPLICIT;
+	  goto do_add;
+
 	case OMP_CLAUSE_IS_DEVICE_PTR:
 	  flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
 	  goto do_add;
@@ -12404,8 +12409,27 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
 				      OMP_CLAUSES (expr));
       break;
     case OMP_TARGET_DATA:
-      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
-				      OMP_CLAUSES (expr));
+      /* Put use_device_{ptr,addr} clauses last, as map clauses are supposed
+	 to be evaluated before the use_device_{ptr,addr} clauses if they
+	 refer to the same variables.  */
+      {
+	tree use_device_clauses;
+	tree *pc, *uc = &use_device_clauses;
+	for (pc = &OMP_CLAUSES (expr); *pc; )
+	  if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_PTR
+	      || OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_USE_DEVICE_ADDR)
+	    {
+	      *uc = *pc;
+	      *pc = OMP_CLAUSE_CHAIN (*pc);
+	      uc = &OMP_CLAUSE_CHAIN (*uc);
+	    }
+	  else
+	    pc = &OMP_CLAUSE_CHAIN (*pc);
+	*uc = NULL_TREE;
+	*pc = use_device_clauses;
+	stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_DATA,
+					OMP_CLAUSES (expr));
+      }
       break;
     case OMP_TEAMS:
       stmt = gimple_build_omp_teams (body, OMP_CLAUSES (expr));
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ca8edde..444610b 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1243,9 +1243,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	  if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
 	       && !omp_is_reference (decl))
 	      || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
-	    install_var_field (decl, true, 3, ctx);
+	    install_var_field (decl, true, 11, ctx);
 	  else
-	    install_var_field (decl, false, 3, ctx);
+	    install_var_field (decl, false, 11, ctx);
 	  if (DECL_SIZE (decl)
 	      && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
 	    {
@@ -11857,11 +11857,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    ovar = OMP_CLAUSE_DECL (c);
 	    var = lookup_decl_in_outer_ctx (ovar, ctx);
-	    x = build_sender_ref (ovar, ctx);
 	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
-	      tkind = GOMP_MAP_USE_DEVICE_PTR;
+	      {
+		tkind = GOMP_MAP_USE_DEVICE_PTR;
+		x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx);
+	      }
 	    else
-	      tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+	      {
+		tkind = GOMP_MAP_FIRSTPRIVATE_INT;
+		x = build_sender_ref (ovar, ctx);
+	      }
 	    type = TREE_TYPE (ovar);
 	    if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_ADDR
 		 && !omp_is_reference (ovar))
@@ -12032,7 +12037,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
 	    var = OMP_CLAUSE_DECL (c);
 	    if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR)
-	      x = build_sender_ref (var, ctx);
+	      x = build_sender_ref ((splay_tree_key) &DECL_UID (var), ctx);
 	    else
 	      x = build_receiver_ref (var, false, ctx);
 	    if (is_variable_sized (var))
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 01440f7..5356d00 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,12 @@
+2019-08-08  Jakub Jelinek  <jakub@redhat.com>
+
+	* target.c (gomp_map_vars_internal): For GOMP_MAP_USE_DEVICE_PTR
+	perform the lookup in the first loop only if !not_found_cnt, otherwise
+	perform lookups for it in the second loop guarded with
+	if (not_found_cnt || has_firstprivate).
+	* testsuite/libgomp.c/target-37.c: New test.
+	* testsuite/libgomp.c++/target-22.C: New test.
+
 2019-08-07  Jakub Jelinek  <jakub@redhat.com>
 
 	* testsuite/libgomp.c/target-18.c (struct S): New type.
diff --git a/libgomp/target.c b/libgomp/target.c
index 9416401..1c9ca68 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -580,20 +580,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	}
       else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
 	{
-	  cur_node.host_start = (uintptr_t) hostaddrs[i];
-	  cur_node.host_end = cur_node.host_start;
-	  splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
-	  if (n == NULL)
+	  tgt->list[i].key = NULL;
+	  if (!not_found_cnt)
 	    {
-	      gomp_mutex_unlock (&devicep->lock);
-	      gomp_fatal ("use_device_ptr pointer wasn't mapped");
 	    }
-	  cur_node.host_start -= n->host_start;
-	  hostaddrs[i]
-	    = (void *) (n->tgt->tgt_start + n->tgt_offset
-			+ cur_node.host_start);
-	  tgt->list[i].key = NULL;
-	  tgt->list[i].offset = ~(uintptr_t) 0;
+	  else
+	    tgt->list[i].offset = 0;
 	  continue;
 	}
       else if ((kind & typemask) == GOMP_MAP_STRUCT)
@@ -791,9 +783,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 		tgt_size += len;
 		continue;
 	      case GOMP_MAP_FIRSTPRIVATE_INT:
-	      case GOMP_MAP_USE_DEVICE_PTR:
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
 		continue;
+	      case GOMP_MAP_USE_DEVICE_PTR:
+		if (tgt->list[i].offset == 0)
+		  {
+		    cur_node.host_start = (uintptr_t) hostaddrs[i];
+		    cur_node.host_end = cur_node.host_start;
+		    n = gomp_map_lookup (mem_map, &cur_node);
+		    if (n == NULL)
+		      {
+			gomp_mutex_unlock (&devicep->lock);
+			gomp_fatal ("use_device_ptr pointer wasn't mapped");
+		      }
+		    cur_node.host_start -= n->host_start;
+		    hostaddrs[i]
+		      = (void *) (n->tgt->tgt_start + n->tgt_offset
+				  + cur_node.host_start);
+		    tgt->list[i].offset = ~(uintptr_t) 0;
+		  }
+		continue;
 	      case GOMP_MAP_STRUCT:
 		first = i + 1;
 		last = i + sizes[i];
diff --git a/libgomp/testsuite/libgomp.c++/target-22.C b/libgomp/testsuite/libgomp.c++/target-22.C
new file mode 100644
index 0000000..9d9dea0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-22.C
@@ -0,0 +1,99 @@
+extern "C" void abort (void);
+struct S { int e, f; };
+
+void
+foo (int *&p, int (&s)[5], int &t, S &u, int n)
+{
+  int a[4] = { 7, 8, 9, 10 }, b[n], c[3] = { 20, 21, 22 };
+  int *r = a + 1, *q = p - 1, i, err;
+  int v = 27;
+  S w = { 28, 29 };
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data map(to:a) use_device_ptr(r) map(from:err)
+  #pragma omp target is_device_ptr(r) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (r[i - 1] != 7 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_ptr(p) map(from:err) map(to:q[:4])
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:b) use_device_addr(b) map(from:err)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < n; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr(c) map(to:c) map(from:err)
+  #pragma omp target is_device_ptr(c) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 3; i++)
+      if (c[i] != 20 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:s[:5]) use_device_addr(s) map(from:err)
+  #pragma omp target is_device_ptr(s) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 5; i++)
+      if (s[i] != 17 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr (v) map(to: v) map(to:u) use_device_addr (u) map(from:err)
+  {
+    int *z = &v;
+    S *x = &u;
+    #pragma omp target is_device_ptr (z, x) map(from:err)
+    {
+      err = 0;
+      if (*z != 27 || x->e != 25 || x->f != 26)
+	err = 1;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to: t) use_device_addr (t, w) map (to: w) map(from:err)
+  {
+    int *z = &t;
+    S *x = &w;
+    #pragma omp target is_device_ptr (z) is_device_ptr (x) map(from:err)
+    {
+      err = 0;
+      if (*z != 24 || x->e != 28 || x->f != 29)
+	err = 1;
+    }
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  int a[4] = { 0, 1, 2, 3 }, b[5] = { 17, 18, 19, 20, 21 };
+  int *p = a + 1;
+  int t = 24;
+  S u = { 25, 26 };
+  foo (p, b, t, u, 9);
+}
diff --git a/libgomp/testsuite/libgomp.c/target-37.c b/libgomp/testsuite/libgomp.c/target-37.c
new file mode 100644
index 0000000..b3cc6a2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c/target-37.c
@@ -0,0 +1,71 @@
+extern void abort (void);
+struct S { int e, f; };
+
+void
+foo (int n)
+{
+  int a[4] = { 0, 1, 2, 3 }, b[n], c = 4;
+  struct S d = { 5, 6 };
+  int *p = a + 1, i, err;
+  for (i = 0; i < n; i++)
+    b[i] = 9 + i;
+  #pragma omp target data use_device_ptr(p) map(from:err) map(to:a)
+  #pragma omp target is_device_ptr(p) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (p[i - 1] != i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  for (i = 0; i < 4; i++)
+    a[i] = 23 + i;
+  #pragma omp target data map(to:a) use_device_addr(a) map(from:err)
+  #pragma omp target is_device_ptr(a) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (a[i] != 23 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr(b) map(from:err) map(to:b)
+  #pragma omp target is_device_ptr(b) private(i) map(from:err)
+  {
+    err = 0;
+    for (i = 0; i < 4; i++)
+      if (b[i] != 9 + i)
+	err = 1;
+  }
+  if (err)
+    abort ();
+  #pragma omp target data map(to:c) use_device_addr(c) map(from:err)
+  {
+    int *q = &c;
+    #pragma omp target is_device_ptr(q) map(from:err)
+    {
+      err = *q != 4;
+    }
+  }
+  if (err)
+    abort ();
+  #pragma omp target data use_device_addr(d) map(to:d) map(from:err)
+  {
+    struct S *r = &d;
+    #pragma omp target is_device_ptr(r) map(from:err)
+    {
+      err = r->e != 5 || r->f != 6;
+    }
+  }
+  if (err)
+    abort ();
+}
+
+int
+main ()
+{
+  foo (9);
+  return 0;
+}