Message ID | a721109c-6c43-4a92-2480-ea867617a5a2@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | OpenACC – support "if" + "if_present" clauses with "host_data" | expand |
Hi Tobias, > Build on x86-64-gnu-linux without offloading and with nvptx offloading. > OK for the trunk? I can't really say a lot about OpenACC, but the changes do look reasonable. So, OK for trunk. Regards Thomas
First *ping* (two weeks old but all during the holiday season). Side note, for what it is worth: Thomas K regarded the Fortran part as reasonable/OK – but those OpenACC changes also affect C/C++/omp-low.c/libgomp :-) – cf. https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01596.html Cheers, Tobias On 12/24/19 3:23 PM, Tobias Burnus wrote: > 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. >
Hi! Jakub, please see one question below. On 2019-12-24T15:23:56+0100, Tobias Burnus <tobias@codesourcery.com> wrote: > OK for the trunk? Tobias, thanks for taking over this patch. I have a few additional small changes that I'd like to do (have WIP patches already), but what we've got here already is OK for trunk with minor changes, see below. > PS: History: [...] > A minor fix was committed to OG8?/OG9 as Rev. > 995f9680a46c3a7246fe465faa847f8009e47ed8. That patch by Julian is on og9 only, and -- as far as I can tell -- has never been posted/discussed. It may be obvious enough; we shall look into it later. (I have not yet verified whether it is sufficiently contained in the patch you posted here.) > --- a/gcc/fortran/openmp.c > +++ b/gcc/fortran/openmp.c > -#define OACC_HOST_DATA_CLAUSES omp_mask (OMP_CLAUSE_USE_DEVICE) > +#define OACC_HOST_DATA_CLAUSES omp_mask \ Please remove the superfluous (benign) 'omp_mask' on the line above. > + (omp_mask (OMP_CLAUSE_USE_DEVICE) \ > + | OMP_CLAUSE_IF \ > + | OMP_CLAUSE_IF_PRESENT) > --- 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) | case OMP_CLAUSE_USE_DEVICE_PTR: | case OMP_CLAUSE_USE_DEVICE_ADDR: | case OMP_CLAUSE_IS_DEVICE_PTR: | ovar = OMP_CLAUSE_DECL (c); | var = lookup_decl_in_outer_ctx (ovar, ctx); | | if (lang_hooks.decls.omp_array_data (ovar, true)) | { | tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR | ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT); | x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx); | } | else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) | { | tkind = GOMP_MAP_USE_DEVICE_PTR; | x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx); | } | else | { > 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; (For context: OpenACC only has an 'use_device' clause, mapped to 'OMP_CLAUSE_USE_DEVICE_PTR'; the other ones are for OpenMP only.) Can you foresee any C/C++/Fortran "magic" ;-) such that for OpenACC 'host_data' construct with 'use_device' clause(s), we'd get something different from 'tkind == GOMP_MAP_USE_DEVICE_PTR' here? In that case, any 'if_present' clause would not be effective. (I have a WIP patch to 'gcc_assert' that.) > type = TREE_TYPE (ovar); > if (lang_hooks.decls.omp_array_data (ovar, true)) > var = lang_hooks.decls.omp_array_data (ovar, false); > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -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 'GOMP_MAP_USE_DEVICE_PTR' is "above", not "below". ;-) > + 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), Move 'GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT' after the 'GOMP_MAP_FORCE_*' ones following here: > /* Allocate. */ > GOMP_MAP_FORCE_ALLOC = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_ALLOC), > /* ..., and copy to device. */ | GOMP_MAP_FORCE_TO = (GOMP_MAP_FLAG_FORCE | GOMP_MAP_TO), | /* ..., and copy from device. */ | 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), ..., so right before the 'GOMP_MAP_ALWAYS_*' ones, which are the other 'GOMP_MAP_FLAG_SPECIAL_2' users: | /* If not already present, allocate. And unconditionally copy to | device. */ | GOMP_MAP_ALWAYS_TO = (GOMP_MAP_FLAG_SPECIAL_2 | GOMP_MAP_TO), | /* If not already present, allocate. And unconditionally copy from | device. */ | GOMP_MAP_ALWAYS_FROM = (GOMP_MAP_FLAG_SPECIAL_2 | | GOMP_MAP_FROM), | /* If not already present, allocate. And unconditionally copy to and from | device. */ | GOMP_MAP_ALWAYS_TOFROM = (GOMP_MAP_FLAG_SPECIAL_2 | | GOMP_MAP_TOFROM), Jakub, please speak up in case that it's not OK to occupy 'GOMP_MAP_FLAG_SPECIAL_2 | 0' for 'GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT', if you foresee any other/better use of that value. This value would correspond to (non-sensical) "'GOMP_MAP_ALWAYS_ALLOC'". > --- a/libgomp/target.c > +++ b/libgomp/target.c I'm have a WIP patch that should make this even more simpler. > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-7.c > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-fortran/host_data-5.F90 If I understand correctly that these two are testing the same things, then please cross-reference them: "C/C++ variant of 'libgomp.oacc-fortran/host_data-5.F90'", "Fortran variant of 'libgomp.oacc-c-c++-common/host_data-7.c'. If that makes sense to consider (now or later), as I see there's some overlap in the 'gcc/omp-low.c:lower_omp_target' code paths: do we need test cases to verify 'if_present' in combination with Fortran optional arguments? Grüße Thomas
Hi! On 2020-01-09T22:00:09+0100, I wrote: > On 2019-12-24T15:23:56+0100, Tobias Burnus <tobias@codesourcery.com> wrote: >> PS: History: [...] >> A minor fix was committed to OG8?/OG9 as Rev. >> 995f9680a46c3a7246fe465faa847f8009e47ed8. > > That patch by Julian is on og9 only, and -- as far as I can tell -- has > never been posted/discussed. It may be obvious enough; we shall look > into it later. (I have not yet verified whether it is sufficiently > contained in the patch you posted here.) Done. >> --- 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) > | case OMP_CLAUSE_USE_DEVICE_PTR: > | case OMP_CLAUSE_USE_DEVICE_ADDR: > | case OMP_CLAUSE_IS_DEVICE_PTR: > | ovar = OMP_CLAUSE_DECL (c); > | var = lookup_decl_in_outer_ctx (ovar, ctx); > | > | if (lang_hooks.decls.omp_array_data (ovar, true)) > | { > | tkind = (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR > | ? GOMP_MAP_USE_DEVICE_PTR : GOMP_MAP_FIRSTPRIVATE_INT); > | x = build_sender_ref ((splay_tree_key) &DECL_NAME (ovar), ctx); > | } > | else if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_IS_DEVICE_PTR) > | { > | tkind = GOMP_MAP_USE_DEVICE_PTR; > | x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar), ctx); > | } > | else > | { >> 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; > > (For context: OpenACC only has an 'use_device' clause, mapped to > 'OMP_CLAUSE_USE_DEVICE_PTR'; the other ones are for OpenMP only.) > Can you foresee any C/C++/Fortran "magic" ;-) such that for OpenACC > 'host_data' construct with 'use_device' clause(s), we'd get something > different from 'tkind == GOMP_MAP_USE_DEVICE_PTR' here? Tobias is looking into that. > In that case, > any 'if_present' clause would not be effective. (I have a WIP patch to > 'gcc_assert' that.) Done; doesn't trigger given the current test cases. Also, calling 'omp_find_clause' here in a loop (over all mappings) is something that Jakub would normally flag up ;-) -- I reworked this code to use the standard idiom of doing such a transformation (here: 'if_present' clause applies to all 'use_device' clauses) early, in the standard place in the gimplifier, by means of a new 'OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT' flag, and also add more testsuite coverage. >> --- a/libgomp/target.c >> +++ b/libgomp/target.c > > I'm have a WIP patch that should make this even more simpler. Done. This is mostly code restructuring, but also, in the first loop, "not present" case, we set 'tgt->list[i].offset = OFFSET_INLINED' already, to avoid again processing this in the second loop. > If that makes sense to consider (now or later), as I see there's some > overlap in the 'gcc/omp-low.c:lower_omp_target' code paths: do we need > test cases to verify 'if_present' in combination with Fortran optional > arguments? Tobias is looking into that. See attached "Further changes for the OpenACC 'if_present' clause on the 'host_data' construct"; committed to trunk in r280149. Grüße Thomas
2019-12-24 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. 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