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.

From-SVN: r280149

9 files changed, 89 insertions, 32 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 9e9f822..a195863 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 00d264f..fe7236d 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 9a36192..eb3fe96 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 6576aee..cccc285 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 658b7a6..ac24446 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 2ac1c0d..558e800 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 62130f4..765ea2a 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 e895a4f..fe2e62b 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 9ca9ab5..9342220 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)