diff options
author | Tobias Burnus <burnus@gcc.gnu.org> | 2020-01-10 16:08:41 +0100 |
---|---|---|
committer | Tobias Burnus <burnus@gcc.gnu.org> | 2020-01-10 16:08:41 +0100 |
commit | d5c23c6ceacf666f218676b648801379044e326a (patch) | |
tree | 27ff72e6195bc05973f17caeee04c66ed8f1db57 /libgomp | |
parent | 7cee96370cf624dbda81fcd3cd32ddb48a2fc3d3 (diff) | |
download | gcc-d5c23c6ceacf666f218676b648801379044e326a.zip gcc-d5c23c6ceacf666f218676b648801379044e326a.tar.gz gcc-d5c23c6ceacf666f218676b648801379044e326a.tar.bz2 |
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
Diffstat (limited to 'libgomp')
-rw-r--r-- | libgomp/oacc-parallel.c | 3 | ||||
-rw-r--r-- | libgomp/target.c | 14 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c | 66 | ||||
-rw-r--r-- | libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 | 92 |
4 files changed, 173 insertions, 2 deletions
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index edfc606..c7e46e3 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 617baec..522b69e 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 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 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 |