Further changes for the OpenACC 'if_present' clause on the 'host_data' construct

	gcc/
	* tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition.
	* tree-core.h: Document it.
	* gimplify.c (gimplify_omp_workshare): Set it.
	* omp-low.c (lower_omp_target): Use it.
	* tree-pretty-print.c (dump_omp_clause): Print it.
	gcc/testsuite/
	* c-c++-common/goacc/host_data-1.c: Extend.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.

	gcc/
	* omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>:
	Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'.

	libgomp/
	* target.c (gomp_map_vars_internal)
	<GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code
	paths.

git-svn-id: https://gcc.gnu.org/svn/gcc/trunk@280149 138bc75d-0d04-0410-961f-82ee72b054a4

11 files changed, 133 insertions, 55 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9e9f8221af4f..a195863212ea 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,14 @@
+2020-01-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition.
+	* tree-core.h: Document it.
+	* gimplify.c (gimplify_omp_workshare): Set it.
+	* omp-low.c (lower_omp_target): Use it.
+	* tree-pretty-print.c (dump_omp_clause): Print it.
+
+	* omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>:
+	Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'.
+
 2020-01-10  David Malcolm  <dmalcolm@redhat.com>
 
 	* Makefile.in (OBJS): Add tree-diagnostic-path.o.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 00d264fc90f4..fe7236de4c3f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -12802,14 +12802,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
 				      OMP_CLAUSES (expr));
       break;
-    case OACC_KERNELS:
-      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
-				      OMP_CLAUSES (expr));
-      break;
     case OACC_HOST_DATA:
+      if (omp_find_clause (OMP_CLAUSES (expr), OMP_CLAUSE_IF_PRESENT))
+	{
+	  for (tree c = OMP_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+	    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
+	      OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c) = 1;
+	}
+
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
 				      OMP_CLAUSES (expr));
       break;
+    case OACC_KERNELS:
+      stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
+				      OMP_CLAUSES (expr));
+      break;
     case OACC_PARALLEL:
       stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
 				      OMP_CLAUSES (expr));
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 9a36192e8ef0..eb3fe9688fe4 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12006,9 +12006,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		tkind = GOMP_MAP_FIRSTPRIVATE_INT;
 		x = build_sender_ref (ovar, ctx);
 	      }
-	    if (tkind == GOMP_MAP_USE_DEVICE_PTR
-		&& omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT))
-	      tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
+
+	    if (is_gimple_omp_oacc (ctx->stmt))
+	      {
+		gcc_assert (tkind == GOMP_MAP_USE_DEVICE_PTR);
+
+		if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (c))
+		  tkind = GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT;
+	      }
+
 	    type = TREE_TYPE (ovar);
 	    if (lang_hooks.decls.omp_array_data (ovar, true))
 	      var = lang_hooks.decls.omp_array_data (ovar, false);
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index 6576aee81c81..cccc2853ed59 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,8 @@
+2020-01-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* c-c++-common/goacc/host_data-1.c: Extend.
+	* gfortran.dg/goacc/host_data-tree.f95: Likewise.
+
 2020-01-10  Jakub Jelinek  <jakub@redhat.com>
 
 	PR tree-optimization/93210
diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
index 658b7a677bc7..ac2444614672 100644
--- a/gcc/testsuite/c-c++-common/goacc/host_data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
@@ -1,14 +1,20 @@
 /* Test valid use of host_data directive.  */
 
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
 int v1[3][3];
 
 void
 f (void)
 {
 #pragma acc host_data use_device(v1)
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(v1\\)$" 1 "original" } }
+     { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(v1\\)$" 1 "gimple" } } */
   ;
 
 #pragma acc host_data use_device(v1) if_present
+  /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(v1\\)$" 1 "original" } }
+     { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:v1\\)$" 1 "gimple" } } */
   ;
 }
 
@@ -16,7 +22,7 @@ f (void)
 void bar (float *, float *);
 
 void
-foo (float *x, float *y)
+foo (float *x, float *y, float *yy)
 {
   int n = 1 << 10;
 #pragma acc data create(x[0:n])
@@ -25,26 +31,38 @@ foo (float *x, float *y)
 
     /* This should fail at run time because y is not mapped.  */
 #pragma acc host_data use_device(x,y)
+    /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
+       { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
     bar (x, y);
 
     /* y is still not mapped, but this should not fail at run time but
        continue execution with y remaining as the host address.  */
 #pragma acc host_data use_device(x,y) if_present
+    /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(y\\) use_device_ptr\\(x\\)$" 1 "original" } }
+       { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:y\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
     bar (x, y);
 
-#pragma acc data copyout(y[0:n])
+#pragma acc data copyout(yy[0:n])
     {
-#pragma acc host_data use_device(x,y)
-      bar (x, y);
+#pragma acc host_data use_device(x,yy)
+      /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+	 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
+      bar (x, yy);
 
-#pragma acc host_data use_device(x,y) if_present
-      bar (x, y);
+#pragma acc host_data use_device(x,yy) if_present
+      /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+	 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
+      bar (x, yy);
 
-#pragma acc host_data use_device(x,y) if(x != y)
-      bar (x, y);
+#pragma acc host_data use_device(x,yy) if(x != yy)
+      /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x \\!= yy\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+	 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "gimple" } } */
+      bar (x, yy);
 
-#pragma acc host_data use_device(x,y) if_present if(x != y)
-      bar (x, y);
+#pragma acc host_data use_device(x,yy) if_present if(x == yy)
+      /* { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data if\\(x == yy\\) if_present use_device_ptr\\(yy\\) use_device_ptr\\(x\\)$" 1 "original" } }
+	 { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data if\\(D\\.\[0-9\]+\\) if_present use_device_ptr\\(if_present:yy\\) use_device_ptr\\(if_present:x\\)$" 1 "gimple" } } */
+      bar (x, yy);
     }
   }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
index 2ac1c0d66d62..558e80014d75 100644
--- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
@@ -1,21 +1,23 @@
-! { dg-do compile } 
-! { dg-additional-options "-fdump-tree-original" } 
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
 
 program test
   implicit none
   integer, pointer :: p
 
   !$acc host_data use_device(p)
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\)$" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(p\\)$" 1 "gimple" } }
   !$acc end host_data
 
   !$acc host_data use_device(p) if (p == 42)
+  ! { dg-final { scan-tree-dump-times "(?n)D\\.\[0-9\]+ = \\*p == 42;$" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\)$" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\)$" 1 "gimple" } }
   !$acc end host_data
 
   !$acc host_data use_device(p) if_present if (p == 43)
+  ! { dg-final { scan-tree-dump-times "(?n)D\\.\[0-9\]+ = \\*p == 43;$" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma acc host_data use_device_ptr\\(p\\) if\\(D\\.\[0-9\]+\\) if_present$" 1 "original" } }
+  ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_host_data use_device_ptr\\(if_present:p\\) if\\(D\\.\[0-9\]+\\) if_present$" 1 "gimple" } }
   !$acc end host_data
 end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 3 "original" } }
-! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 42;" 1 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\)" 2 "original" } }
-! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 43;" 1 "original" } }
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\) if_present" 1 "original" } }
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 62130f488ccf..765ea2a95423 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -1175,6 +1175,9 @@ struct GTY(()) tree_base {
        OMP_CLAUSE_REDUCTION_OMP_ORIG_REF in
 	   OMP_CLAUSE_{,TASK_,IN_}REDUCTION
 
+       OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT in
+	   OMP_CLAUSE_USE_DEVICE_PTR
+
        TRANSACTION_EXPR_RELAXED in
 	   TRANSACTION_EXPR
 
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index e895a4f66099..fe2e62b31ba0 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -432,7 +432,7 @@ static void
 dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 {
   const char *name;
-
+  const char *modifier = NULL;
   switch (OMP_CLAUSE_CODE (clause))
     {
     case OMP_CLAUSE_PRIVATE:
@@ -446,13 +446,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       goto print_remap;
     case OMP_CLAUSE_LASTPRIVATE:
       name = "lastprivate";
-      if (!OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (clause))
-	goto print_remap;
-      pp_string (pp, "lastprivate(conditional:");
-      dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
-			 spc, flags, false);
-      pp_right_paren (pp);
-      break;
+      if (OMP_CLAUSE_LASTPRIVATE_CONDITIONAL (clause))
+	modifier = "conditional:";
+      goto print_remap;
     case OMP_CLAUSE_COPYIN:
       name = "copyin";
       goto print_remap;
@@ -464,6 +460,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
       goto print_remap;
     case OMP_CLAUSE_USE_DEVICE_PTR:
       name = "use_device_ptr";
+      if (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT (clause))
+	modifier = "if_present:";
       goto print_remap;
     case OMP_CLAUSE_USE_DEVICE_ADDR:
       name = "use_device_addr";
@@ -501,6 +499,8 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
   print_remap:
       pp_string (pp, name);
       pp_left_paren (pp);
+      if (modifier)
+	pp_string (pp, modifier);
       dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
 			 spc, flags, false);
       pp_right_paren (pp);
diff --git a/gcc/tree.h b/gcc/tree.h
index 9ca9ab58ec02..93422206b63d 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1627,6 +1627,11 @@ class auto_suppress_location_wrappers
 #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
   TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
 
+/* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
+   clause.  */
+#define OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_USE_DEVICE_PTR)->base.public_flag)
+
 #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \
   (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)
 
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 152e52c21c11..81d0c164a3a4 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2020-01-10  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* target.c (gomp_map_vars_internal)
+	<GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code
+	paths.
+
 2020-01-10  Jakub Jelinek  <jakub@redhat.com>
 
 	PR libgomp/93219
diff --git a/libgomp/target.c b/libgomp/target.c
index 522b69e6d5d1..38de1c0cf922 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -740,22 +740,24 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	      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)
+	      if (n != NULL)
+		{
+		  cur_node.host_start -= n->host_start;
+		  hostaddrs[i]
+		    = (void *) (n->tgt->tgt_start + n->tgt_offset
+				+ cur_node.host_start);
+		}
+	      else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
 		{
-		  if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
-		    {
-		      /* If not present, continue using the host address.  */
-		      tgt->list[i].offset = 0;
-		      continue;
-		    }
 		  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;
+	      else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
+		/* If not present, continue using the host address.  */
+		;
+	      else
+		__builtin_unreachable ();
+	      tgt->list[i].offset = OFFSET_INLINED;
 	    }
 	  else
 	    tgt->list[i].offset = 0;
@@ -980,27 +982,40 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
 	      case GOMP_MAP_FIRSTPRIVATE_INT:
 	      case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
 		continue;
-	      case GOMP_MAP_USE_DEVICE_PTR:
 	      case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
+		/* The OpenACC 'host_data' construct only allows 'use_device'
+		   "mapping" clauses, so in the first loop, 'not_found_cnt'
+		   must always have been zero, so all OpenACC 'use_device'
+		   clauses have already been handled.  (We can only easily test
+		   'use_device' with 'if_present' clause here.)  */
+		assert (tgt->list[i].offset == OFFSET_INLINED);
+		/* Nevertheless, FALLTHRU to the normal handling, to keep the
+		   code conceptually simple, similar to the first loop.  */
+	      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)
+		    if (n != NULL)
+		      {
+			cur_node.host_start -= n->host_start;
+			hostaddrs[i]
+			  = (void *) (n->tgt->tgt_start + n->tgt_offset
+				      + cur_node.host_start);
+		      }
+		    else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
 		      {
-			if ((kind & typemask)
-			    == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
-			  /* If not present, continue using the host address.  */
-			  continue;
 			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;
+		    else if ((kind & typemask)
+			     == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
+		      /* If not present, continue using the host address.  */
+		      ;
+		    else
+		      __builtin_unreachable ();
+		    tgt->list[i].offset = OFFSET_INLINED;
 		  }
 		continue;
 	      case GOMP_MAP_STRUCT: