Add OpenACC 2.6 if and if_present clauses on host_data construct: GOACC_FLAG_HOST_DAT...

    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): Likewise.

    gcc/fortran/
    * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
    OMP_CLAUSE_IF_PRESENT.

    gcc/
    * omp-expand.c (expand_omp_target): Handle if_present flag on
    OpenACC host_data construct.

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

    include/
    * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.

    libgomp/
    * libgomp.h (enum gomp_map_vars_kind): Add
    GOMP_MAP_VARS_OPENACC_IF_PRESENT.
    * oacc-parallel.c (GOACC_data_start): Handle
    GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
    * target.c (gomp_map_vars_async): Handle
    GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
    * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
(cherry picked from openacc-gcc-9-branch commit
b52c8d006581b2a657ac3d81ab2527bb398b6615)

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 797a62e1af2934625b69b52d9df339f2fe084a9f..78f4042c5e24460d20f4cabfa2bfcfb27603da5b 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * omp-expand.c (expand_omp_target): Handle if_present flag on
+       OpenACC host_data construct.
+
 2018-12-20  Gergö Barany  <gergo@codesourcery.com>
 
        * omp-low.c (struct omp_context): New fields

diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp
index db92b2422816ac5dcb800ad888063898a555a6ff..e085e879f6731d069a9c3dec71291fcb599e4e7c 100644 (file)
--- a/gcc/c/ChangeLog.omp
+++ b/gcc/c/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF
+       and PRAGMA_OACC_CLAUSE_IF_PRESENT.
+
 2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
             Nathan Sidwell  <nathan@acm.org>
             Julian Brown  <julian@codesourcery.com>

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 0e49167381d3789fd9390db3786b6654db4fc0c1..d11216472d7bbc4bd72b37510427578f14574f13 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15616,7 +15616,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/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 370d895d2c5f5fc0633e834570a003bb0cfeec7b..b2c7cfc20dbd2e9cddd5a7a19a6743455e4a6c18 100644 (file)
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,7 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Likewise.
+
 2018-12-13  Cesar Philippidis  <cesar@codesourcery.com>
             Nathan Sidwell  <nathan@acm.org>
             Julian Brown  <julian@codesourcery.com>

diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index b8de1427650a266df9b5d42075474e4033d42637..3c42f43742ed1d9c456e0f7acacc0eea8c5381d9 100644 (file)
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -38767,7 +38767,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/ChangeLog.omp b/gcc/fortran/ChangeLog.omp
index b01f1ba83b5371fe0e06a1ef19a86374045499bc..00ccd26462c45b31c513e4359b02cb724ccbdcd7 100644 (file)
--- a/gcc/fortran/ChangeLog.omp
+++ b/gcc/fortran/ChangeLog.omp
@@ -1,3 +1,8 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and
+       OMP_CLAUSE_IF_PRESENT.
+
 2019-01-30  Kwok Cheung Yeung  <kcy@codesourcery.com>
 
        * trans-openmp.c (gfc_build_conditional_assign): New.

diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 81b69b9c7af740a6e5761eef3bb832742b6f1349..2a6dce7d65f841844922be1317a997e3b6da3b5a 100644 (file)
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -2001,7 +2001,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
   (OACC_LOOP_CLAUSES | OACC_PARALLEL_CLAUSES)
 #define OACC_KERNELS_LOOP_CLAUSES \
   (OACC_LOOP_CLAUSES | OACC_KERNELS_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-expand.c b/gcc/omp-expand.c
index 27de3be190d6373c03a14afc3aa029514b93a9ca..6dd87689e5ab8bba4455e01dd6cca75caa1d8c19 100644 (file)
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -7591,7 +7591,17 @@ expand_omp_target (struct omp_region *region)
   if (is_gimple_omp_oacc (entry_stmt))
     {
       /* By default, no GOACC_FLAGs are set.  */
-      goacc_flags = integer_zero_node;
+      int goacc_flags_i = 0;
+
+      if (start_ix != BUILT_IN_GOACC_UPDATE
+         && omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT))
+       {
+         gcc_checking_assert (gimple_omp_target_kind (entry_stmt)
+                              == GF_OMP_TARGET_KIND_OACC_HOST_DATA);
+         goacc_flags_i |= GOACC_FLAG_HOST_DATA_IF_PRESENT;
+       }
+
+      goacc_flags = build_int_cst (integer_type_node, goacc_flags_i);
     }
   else
     {

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 10b52bee15d0ba152ee3226175658c99b79c2332..c9917439f699426b089173be82a6a2a8529ea7a0 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,9 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * c-c++-common/goacc/host_data-1.c: Add tests of if and if_present
+       clauses on host_data.
+       * gfortran.dg/goacc/host_data-tree.f95: Likewise.
+
 2018-12-20  Gergö Barany  <gergo@codesourcery.com>
 
        * c-c++-common/goacc/nested-reductions-fail.c: New test.

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/ChangeLog.omp b/include/ChangeLog.omp
index 3cc2e9ace8dec527ce3a334216a1b295c7a96d9c..e632fdb4ebf9fbe648f7df3c51a18ec4c9b2f0d3 100644 (file)
--- a/include/ChangeLog.omp
+++ b/include/ChangeLog.omp
@@ -1,3 +1,7 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant.
+
 2018-10-04  Cesar Philippidis  <cesar@codesourcery.com>
             Julian Brown  <julian@codesourcery.com>
 

diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 2bbba6117123e3946a302fd3d10a928f870ae479..90b18091a13553daedcb719c994efd5848eab925 100644 (file)
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -240,6 +240,8 @@ enum gomp_map_kind
 
 /* Force host fallback execution.  */
 #define GOACC_FLAG_HOST_FALLBACK       (1 << 0)
+/* "if_present" semantics for OpenACC "host_data" constructs.  */
+#define GOACC_FLAG_HOST_DATA_IF_PRESENT        (1 << 1)
 
 /* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted
    bitmask.  */

diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 7f1642187843af36d0f14b3b489232f578b97435..5dc06978255ff6e21bf60e52456c8dab02e43429 100644 (file)
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,13 @@
+2018-12-21  Gergö Barany  <gergo@codesourcery.com>
+
+       * libgomp.h (enum gomp_map_vars_kind): Add
+       GOMP_MAP_VARS_OPENACC_IF_PRESENT.
+       * oacc-parallel.c (GOACC_data_start): Handle
+       GOACC_FLAG_HOST_DATA_IF_PRESENT flag.
+       * target.c (gomp_map_vars_async): Handle
+       GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind.
+       * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test.
+
 2018-12-20  Gergö Barany  <gergo@codesourcery.com>
 
        * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c:

diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index cf3682a09eee8e912a1b7b02c26f857f48547ad0..6bc39ce9aae4c16926f8d403b6a9f8af9509e92d 100644 (file)
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1082,6 +1082,9 @@ struct gomp_device_descr
 enum gomp_map_vars_kind
 {
   GOMP_MAP_VARS_OPENACC,
+  /* Like "GOMP_MAP_VARS_OPENACC", but with "GOACC_FLAG_HOST_DATA_IF_PRESENT"
+     semantics.  */
+  GOMP_MAP_VARS_OPENACC_IF_PRESENT,
   GOMP_MAP_VARS_OPENACC_ENTER_DATA,
   GOMP_MAP_VARS_TARGET,
   GOMP_MAP_VARS_DATA,

diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 3369a3656b7a3f6a9d09b96d00ae9c2f4a62c176..98302f1f7adb0a4c1256312727ce49952d6414e5 100644 (file)
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -331,12 +331,17 @@ GOACC_data_start (int flags_m, size_t mapnum,
 
   handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
 
+  enum gomp_map_vars_kind pragma_kind;
+  if (flags & GOACC_FLAG_HOST_DATA_IF_PRESENT)
+    pragma_kind = GOMP_MAP_VARS_OPENACC_IF_PRESENT;
+  else
+    pragma_kind = GOMP_MAP_VARS_OPENACC;
+
   /* Host fallback or 'do nothing'.  */
   if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
       || (flags & GOACC_FLAG_HOST_FALLBACK))
     {
-      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
-                          GOMP_MAP_VARS_OPENACC);
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, pragma_kind);
       tgt->prev = thr->mapped_data;
       thr->mapped_data = tgt;
 
@@ -345,7 +350,7 @@ GOACC_data_start (int flags_m, size_t mapnum,
 
   gomp_debug (0, "  %s: prepare mappings\n", __FUNCTION__);
   tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
-                      GOMP_MAP_VARS_OPENACC);
+                      pragma_kind);
   gomp_debug (0, "  %s: mappings prepared\n", __FUNCTION__);
   tgt->prev = thr->mapped_data;
   thr->mapped_data = tgt;

diff --git a/libgomp/target.c b/libgomp/target.c
index 320942742b356b6dcb49026d41027f22d289b67f..b34043f31423bff50cc3077df04d0fd0a553e6fa 100644 (file)
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1102,6 +1102,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
          splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
          if (n == NULL)
            {
+              if (pragma_kind == GOMP_MAP_VARS_OPENACC_IF_PRESENT)
+                /* No error, 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-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
new file mode 100644 (file)
index 0000000..c5744fe
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
@@ -0,0 +1,53 @@
+/* Test if, if_present clauses on host_data construct.  */
+
+#include <assert.h>
+#include <stdint.h>
+
+void
+foo (float *p, intptr_t host_p, int shared_mem_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);
+
+      if (!shared_mem_p)
+      {
+#pragma acc host_data use_device(p)
+        /* The device address is different from the host address.  */
+        assert (p != (float *) host_p);
+
+#pragma acc host_data use_device(p) if_present
+        /* p is present now, so this is the same as above.  */
+        assert (p != (float *) host_p);
+      }
+
+#pragma acc host_data use_device(p) if(cond)
+      /* p is the device pointer iff cond is true and device memory is
+         separate from host memory.  */
+      assert ((p != (float *) host_p) == (cond && !shared_mem_p));
+    }
+  }
+}
+
+int
+main (void)
+{
+  float arr[100];
+  int shared_mem_p = 0;
+#if ACC_MEM_SHARED
+  shared_mem_p = 1;
+#endif
+  foo (arr, (intptr_t) arr, shared_mem_p, 0);
+  foo (arr, (intptr_t) arr, shared_mem_p, 1);
+
+  return 0;
+}