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

3 files changed, 42 insertions, 17 deletions

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