On the front-end side, adding "if" and "if_present" to the "acc
host_data" directive is simple as other directives already support those
clauses.
The 'if_present' status has to be passed along the use_device_ptr flag;
for this a new flag has been introduced, using the gap in the
gomp_map_kind enum (16 was still free; now used for
GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT, use_device_ptr has 14).
There was some forth and back whether a new enum should be used or a
flag (passed in the device argument, as OpenACC only uses host and
device and not a device number). Initial version was with new enum, then
with flag and now again a new enum.
On the libgomp side, one simply needs to skip the if-not-present error
and otherwise handle it as use_device_ptr.
Build on x86-64-gnu-linux without offloading and with nvptx offloading.
OK for the trunk?
Cheers,
Tobias
PS: History: The initial version (not public) used an enum but Thomas
was wondering whether a flag is not the better solution, cf. discussion
at https://gcc.gnu.org/ml/gcc/2018-12/msg00118.html – Hence, the first
public version has switched to a flag and submitted for the OG8
(openacc-gcc-8-branch GIT) branch at
https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01559.html and has been
committed as Rev. b52c8d006581b2a657ac3d81ab2527bb398b6615 (this commit
is also part of OG9. — A minor fix was committed to OG8?/OG9 as Rev.
995f9680a46c3a7246fe465faa847f8009e47ed8.
Thomas now wonders whether the flag will be future safe, hence, it now
uses again an enum as in the original version.
PPS: Changes to OG9: Using again an enum; Fortran run-time test and rediffs.
2019-12-24 Gergö Barany <ge...@codesourcery.com>
Thomas Schwinge <tho...@codesourcery.com>
Julian Brown <jul...@codesourcery.com>
Tobias Burnus <tob...@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.
gcc/c/c-parser.c | 4 +-
gcc/cp/parser.c | 4 +-
gcc/fortran/openmp.c | 5 +-
gcc/omp-low.c | 3 +
gcc/testsuite/c-c++-common/goacc/host_data-1.c | 28 ++++++-
gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 12 ++-
include/gomp-constants.h | 6 +-
libgomp/oacc-parallel.c | 3 +-
libgomp/target.c | 14 +++-
.../libgomp.oacc-c-c++-common/host_data-7.c | 65 ++++++++++++++++
.../testsuite/libgomp.oacc-fortran/host_data-5.F90 | 91 ++++++++++++++++++++++
11 files changed, 227 insertions(+), 8 deletions(-)
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index b3763c2e888..a1041b03edb 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 c3c968d9ce4..931fc17fcad 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40461,7 +40461,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 01964f964d7..838c70da5d4 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_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 e697f3c620d..13ccd42966e 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 0c7a857d0bf..658b7a677bc 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 d44ca587051..2ac1c0d66d6 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" } }
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index d14e8b0394a..ae323291b01 100644
--- 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
@@ -93,6 +93,10 @@ enum gomp_map_kind
at the address. If not already mapped, do nothing (and pointer translate
to NULL). */
GOMP_MAP_ZERO_LEN_ARRAY_SECTION = (GOMP_MAP_FLAG_SPECIAL | 3),
+ /* Like GOMP_MAP_USE_DEVICE_PTR below, 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),
/* Allocate. */
GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC),
/* ..., and copy to device. */
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index c853f01b27f..485ca43797e 100644
--- 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 3562c3790f1..f9dd2112a17 100644
--- 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
index 00000000000..39d4700edad
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+
+/* Test if, if_present clauses on host_data construct. */
+
+#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
index 00000000000..46d8865074e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90
@@ -0,0 +1,91 @@
+! { dg-do run }
+!
+! Test if, if_present clauses on host_data construct.
+!
+
+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
