OpenACC – support "if" + "if_present" clauses with "host_data"

2020-01-10  Gergö Barany  <gergo@codesourcery.com>
	    Thomas Schwinge <thomas@codesourcery.com>
	    Julian Brown  <julian@codesourcery.com>
	    Tobias Burnus  <tobias@codesourcery.com>

        gcc/c/
        * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

        gcc/cp/
        * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

        gcc/fortran/
        * openmp.c (OACC_HOST_DATA_CLAUSES): Add PRAGMA_OACC_CLAUSE_IF
        and PRAGMA_OACC_CLAUSE_IF_PRESENT.

	gcc/
	* omp-low.c (lower_omp_target): Use GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
	if PRAGMA_OACC_CLAUSE_IF_PRESENT exist.

	gcc/testsuite/
	* c-c++-common/goacc/host_data-1.c: Added tests of if and if_present
	clauses on host_data.
	* gfortran.dg/goacc/host_data-tree.f95: Likewise.

	include/
	* gomp-constants.h (enum gomp_map_kind): New enumeration constant
	GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT.
        
	libgomp/
	* oacc-parallel.c (GOACC_data_start): Handle
	GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT.
	* target.c (gomp_map_vars_async): Likewise.
	* testsuite/libgomp.oacc-c-c++-common/host_data-7.c: New.
	* testsuite/libgomp.oacc-fortran/host_data-5.F90: New.

From-SVN: r280115

6 files changed, 51 insertions, 5 deletions

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ea06069..bf9e3e0 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -16701,7 +16701,9 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter)
 */
 
 #define OACC_HOST_DATA_CLAUSE_MASK					\
-	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
+	( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE)          \
+	 | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)                  \
+	 | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) )
 
 static tree
 c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p)
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 595c185..2ddbe13 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40487,7 +40487,9 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
   structured-block  */
 
 #define OACC_HOST_DATA_CLAUSE_MASK					\
-  ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
+  ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE)                \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF)                        \
+   | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) )
 
 static tree
 cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index c105657..1062212 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2031,7 +2031,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
   (OACC_LOOP_CLAUSES | OACC_KERNELS_CLAUSES)
 #define OACC_SERIAL_LOOP_CLAUSES \
   (OACC_LOOP_CLAUSES | OACC_SERIAL_CLAUSES)
-#define OACC_HOST_DATA_CLAUSES omp_mask (OMP_CLAUSE_USE_DEVICE)
+#define OACC_HOST_DATA_CLAUSES \
+  (omp_mask (OMP_CLAUSE_USE_DEVICE)					      \
+   | OMP_CLAUSE_IF							      \
+   | OMP_CLAUSE_IF_PRESENT)
 #define OACC_DECLARE_CLAUSES \
   (omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT	      \
    | OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT    \
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index e692a53..9a36192 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -12006,6 +12006,9 @@ 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;
 	    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/c-c++-common/goacc/host_data-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
index 0c7a857..658b7a6 100644
--- a/gcc/testsuite/c-c++-common/goacc/host_data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c
@@ -7,6 +7,9 @@ f (void)
 {
 #pragma acc host_data use_device(v1)
   ;
+
+#pragma acc host_data use_device(v1) if_present
+  ;
 }
 
 
@@ -16,9 +19,32 @@ void
 foo (float *x, float *y)
 {
   int n = 1 << 10;
-#pragma acc data create(x[0:n]) copyout(y[0:n])
+#pragma acc data create(x[0:n])
   {
+    bar (x, y);
+
+    /* This should fail at run time because y is not mapped.  */
 #pragma acc host_data use_device(x,y)
     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
+    bar (x, y);
+
+#pragma acc data copyout(y[0:n])
+    {
+#pragma acc host_data use_device(x,y)
+      bar (x, y);
+
+#pragma acc host_data use_device(x,y) if_present
+      bar (x, y);
+
+#pragma acc host_data use_device(x,y) if(x != y)
+      bar (x, y);
+
+#pragma acc host_data use_device(x,y) if_present if(x != y)
+      bar (x, y);
+    }
   }
 }
diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
index d44ca58..2ac1c0d 100644
--- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
@@ -7,5 +7,15 @@ program test
 
   !$acc host_data use_device(p)
   !$acc end host_data
+
+  !$acc host_data use_device(p) if (p == 42)
+  !$acc end host_data
+
+  !$acc host_data use_device(p) if_present if (p == 43)
+  !$acc end host_data
 end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 1 "original" } }
+! { 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" } }