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
This commit is contained in:
Tobias Burnus
parent
7cee96370c
commit
d5c23c6cea
|
|
@ -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)
|
||||
|
|
|
|||
|
|
@ -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)
|
||||
|
|
|
|||
|
|
@ -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 \
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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" } }
|
||||
|
|
|
|||
|
|
@ -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),
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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");
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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
|
||||
Loading…
Reference in New Issue