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

1 files changed, 66 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c
new file mode 100644
index 00000000000..6830ef1e7ed
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* Test if, if_present clauses on host_data construct.  */
+/* C/C++ variant of 'libgomp.oacc-fortran/host_data-5.F90' */
+
+#include <assert.h>
+#include <stdint.h>
+
+void
+foo (float *p, intptr_t host_p, int cond)
+{
+  assert (p == (float *) host_p);
+
+#pragma acc data copyin(host_p)
+  {
+#pragma acc host_data use_device(p) if_present
+    /* p not mapped yet, so it will be equal to the host pointer.  */
+    assert (p == (float *) host_p);
+
+#pragma acc data copy(p[0:100])
+    {
+      /* Not inside a host_data construct, so p is still the host pointer.  */
+      assert (p == (float *) host_p);
+
+#pragma acc host_data use_device(p)
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* The device address is different from host address.  */
+	assert (p != (float *) host_p);
+#endif
+      }
+
+#pragma acc host_data use_device(p) if_present
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* p is present now, so this is the same as above.  */
+	assert (p != (float *) host_p);
+#endif
+      }
+
+#pragma acc host_data use_device(p) if(cond)
+      {
+#if ACC_MEM_SHARED
+	assert (p == (float *) host_p);
+#else
+	/* p is the device pointer iff cond is true.  */
+	assert ((p != (float *) host_p) == cond);
+#endif
+      }
+    }
+  }
+}
+
+int
+main (void)
+{
+  float arr[100];
+  foo (arr, (intptr_t) arr, 0);
+  foo (arr, (intptr_t) arr, 1);
+
+  return 0;
+}