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

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ea060696c27e7b140170d1bacfbe95c933661b9a..bf9e3e05d3e83e9a6a2f40fa6fe4c3ec97dba8e0 100644 (file)
--- 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 595c185f2f8473bc003a8849341227db001d74ed..2ddbe138ae66ade924bf137c1a9348a16e957c43 100644 (file)
--- 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 c105657a91b41b4b022e98466ab9bfd166f8ed2e..10622129659b9db2df126ef65c05a2a1ab0a7438 100644 (file)
--- 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 e692a53a3de92ef875ff84baaf4a05523ea86edc..9a36192e8ef0489e2b5cf0e0479878bb38909532 100644 (file)
--- 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 0c7a857d0bf355cd988010357d9cedc17d2c1291..658b7a677bc70825a327099b21d7bd44ab0d8fb4 100644 (file)
--- 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 d44ca5870516fbb070391f543fb22f62af7a1654..2ac1c0d66d62a511178d9471edc77ca7f3d74a09 100644 (file)
--- 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" } }

diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9ca2178cbd4cb4bda9c6c0858ed5b7d97a0b70cf..5a7cc2c3f011adadf54eb9f5b33c4b5898072f95 100644 (file)
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -79,7 +79,7 @@ enum gomp_map_kind
     /* OpenACC link.  */
     GOMP_MAP_LINK =                    (GOMP_MAP_FLAG_SPECIAL_1 | 2),
     /* Use device data if present, fall back to host address otherwise.  */
-    GOMP_MAP_IF_PRESENT =                      (GOMP_MAP_FLAG_SPECIAL_1 | 3),
+    GOMP_MAP_IF_PRESENT =              (GOMP_MAP_FLAG_SPECIAL_1 | 3),
     /* Do not map, copy bits for firstprivate instead.  */
     GOMP_MAP_FIRSTPRIVATE =            (GOMP_MAP_FLAG_SPECIAL | 0),
     /* Similarly, but store the value in the pointer rather than
@@ -101,6 +101,10 @@ enum gomp_map_kind
     GOMP_MAP_FORCE_FROM =              (GOMP_MAP_FLAG_FORCE | GOMP_MAP_FROM),
     /* ..., and copy to and from device.  */
     GOMP_MAP_FORCE_TOFROM =            (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TOFROM),
+    /* Like GOMP_MAP_USE_DEVICE_PTR above, translate a host to a device
+       address.  If translation fails because the target is not mapped,
+       continue using the host address. */
+    GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_2 | 0),
     /* If not already present, allocate.  And unconditionally copy to
        device.  */
     GOMP_MAP_ALWAYS_TO =               (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_TO),

diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index edfc6067ee915f21e5d8aa5c95586bc1c3c2bd99..c7e46e35bd6b8d67a575329bb32af046b5894dd7 100644 (file)
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -415,7 +415,8 @@ GOACC_data_start (int flags_m, size_t mapnum,
        = _ACC_OTHER_EVENT_INFO_VALID_BYTES;
       enter_data_event_info.other_event.parent_construct = acc_construct_data;
       for (int i = 0; i < mapnum; ++i)
-       if ((kinds[i] & 0xff) == GOMP_MAP_USE_DEVICE_PTR)
+       if ((kinds[i] & 0xff) == GOMP_MAP_USE_DEVICE_PTR
+           || (kinds[i] & 0xff) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
          {
            /* If there is one such data mapping kind, then this is actually an
               OpenACC 'host_data' construct.  (GCC maps the OpenACC

diff --git a/libgomp/target.c b/libgomp/target.c
index 617baec8b401b19c7342e299c1aaad4702042336..522b69e6d5d17642041776a083058f614592394e 100644 (file)
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -720,7 +720,8 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
          tgt->list[i].offset = OFFSET_INLINED;
          continue;
        }
-      else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR)
+      else if ((kind & typemask) == GOMP_MAP_USE_DEVICE_PTR
+              || (kind & typemask) == GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT)
        {
          tgt->list[i].key = NULL;
          if (!not_found_cnt)
@@ -741,6 +742,12 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
              splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
              if (n == NULL)
                {
+                 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");
                }
@@ -974,6 +981,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
              case GOMP_MAP_ZERO_LEN_ARRAY_SECTION:
                continue;
              case GOMP_MAP_USE_DEVICE_PTR:
+             case GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT:
                if (tgt->list[i].offset == 0)
                  {
                    cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -981,6 +989,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
                    n = gomp_map_lookup (mem_map, &cur_node);
                    if (n == NULL)
                      {
+                       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");
                      }

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 (file)
index 0000000..6830ef1
--- /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;
+}

diff --git a/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
new file mode 100644 (file)
index 0000000..483ac3f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -0,0 +1,92 @@
+! { dg-do run }
+!
+! Test if, if_present clauses on host_data construct.
+!
+! Fortran variant of 'libgomp.oacc-c-c++-common/host_data-7.c'.
+!
+program main
+  use iso_c_binding
+  implicit none
+  real, target :: var, arr(100)
+  integer(c_intptr_t) :: host_p, host_parr
+  host_p = transfer(c_loc(var), host_p)
+  host_parr = transfer(c_loc(arr), host_parr)
+  call foo (var, arr, host_p, host_parr, .false.)
+  call foo (var, arr, host_p, host_parr, .true.)
+
+contains
+
+subroutine foo (p2, parr, host_p, host_parr, cond)
+  use openacc
+  implicit none
+  real, target, intent(in) :: parr(:), p2
+  integer(c_intptr_t), value, intent(in) :: host_p, host_parr
+  logical, value, intent(in) :: cond
+  real, pointer :: p
+  p => p2
+
+  if (host_p /= transfer(c_loc(p), host_p)) stop 1
+  if (host_parr /= transfer(c_loc(parr), host_parr)) stop 2
+#if !ACC_MEM_SHARED
+  if (acc_is_present(p, c_sizeof(p))) stop 3
+  if (acc_is_present(parr, 1)) stop 4
+#endif
+  
+  !$acc data copyin(host_p, host_parr)
+#if !ACC_MEM_SHARED
+    if (acc_is_present(p, c_sizeof(p))) stop 5
+    if (acc_is_present(parr, 1)) stop 6
+#endif
+    !$acc host_data use_device(p, parr) if_present
+      ! not mapped yet, so it will be equal to the host pointer.
+      if (transfer(c_loc(p), host_p) /= host_p) stop 7
+      if (transfer(c_loc(parr), host_parr) /= host_parr) stop 8
+    !$acc end host_data
+#if !ACC_MEM_SHARED
+    if (acc_is_present(p, c_sizeof(p))) stop 9
+    if (acc_is_present(parr, 1)) stop 10
+#endif
+
+    !$acc data copy(p, parr)
+      if (.not. acc_is_present(p, c_sizeof(p))) stop 11
+      if (.not. acc_is_present(parr, 1)) stop 12
+      ! Not inside a host_data construct, so still the host pointer.
+      if (transfer(c_loc(p), host_p) /= host_p) stop 13
+      if (transfer(c_loc(parr), host_parr) /= host_parr) stop 14
+      
+      !$acc host_data use_device(p, parr)
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 15
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 16
+#else
+        ! The device address is different from host address.
+        if (transfer(c_loc(p), host_p) == host_p) stop 17
+        if (transfer(c_loc(parr), host_parr) == host_parr) stop 18
+#endif
+      !$acc end host_data
+
+      !$acc host_data use_device(p, parr) if_present
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 19
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 20
+#else
+        ! is present now, so this is the same as above.
+        if (transfer(c_loc(p), host_p) == host_p) stop 21
+        if (transfer(c_loc(parr), host_parr) == host_parr) stop 22
+#endif
+      !$acc end host_data
+
+      !$acc host_data use_device(p, parr) if(cond)
+#if ACC_MEM_SHARED
+        if (transfer(c_loc(p), host_p) /= host_p) stop 23
+        if (transfer(c_loc(parr), host_parr) /= host_parr) stop 24
+#else
+        ! is the device pointer iff cond is true.
+        if ((transfer(c_loc(p), host_p) /= host_p) .neqv. cond) stop 25
+        if ((transfer(c_loc(parr), host_parr) /= host_parr) .neqv. cond) stop 26
+#endif
+      !$acc end host_data
+    !$acc end data
+  !$acc end data
+end subroutine foo
+end