Hi Gergő! On Fri, 21 Dec 2018 13:29:09 +0100, Gergö Barany <ge...@codesourcery.com> wrote: > OpenACC 2.6 specifies `if' and `if_present' clauses on the `host_data' > construct. These patches add support for these clauses. The first patch, > by Thomas, reorganizes libgomp internals to turn a "device" argument > into "flags" that can provide more information to the runtime. The > second patch adds support for the `if' and `if_present' clauses, using > the new flag mechanism. > > OK for openacc-gcc-8-branch?
Yes, thanks. To record the review effort, please include "Reviewed-by: Thomas Schwinge <tho...@codesourcery.com>" in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>. (Not for my own commit, of course.) Again, just the commit message of the second commit needs to be adjusted, from: > [...] > gcc/testsuite/c-c++-common/goacc/ > * host_data-1.c: Add tests of if and if_present clauses on host_data. > gcc/testsuite/gfortran.dg/goacc/ > * host_data-tree.f95: Likewise. > [...] > libgomp/ > * libgomp.h (enum gomp_map_vars_kind): Add > GOMP_MAP_VARS_OPENACC_IF_PRESENT. > > libgomp/ > * 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. > > libgomp/testsuite/libgomp.oacc-c-c++-common/ > * host_data-6.c: New test. ... to: > [...] > gcc/testsuite/ > * c-c++-common/goacc/host_data-1.c: Add tests of if and if_present > clauses on host_data. [add suitable line break some where] > * gfortran.dg/goacc/host_data-tree.f95: Likewise. > [...] > 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. Grüße Thomas > From 6d719cc2bcfa8f7ed8cb59e753e44aab6bf634fb Mon Sep 17 00:00:00 2001 > From: Thomas Schwinge <tho...@codesourcery.com> > Date: Wed, 19 Dec 2018 20:04:18 +0100 > Subject: [PATCH 1/2] For libgomp OpenACC entry points, redefine the "device" > argument to "flags" > > ... so that we're then able to use this for other flags in addition to > "GOACC_FLAG_HOST_FALLBACK". > > gcc/ > * omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP > code paths. Update for libgomp OpenACC entry points change. > include/ > * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK) > (GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define. > libgomp/ > * oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel) > (GOACC_data_start, GOACC_enter_exit_data, GOACC_update) > (GOACC_declare): Redefine the "device" argument to "flags". > --- > gcc/ChangeLog.openacc | 5 ++ > gcc/omp-expand.c | 111 > +++++++++++++++++++++++++++++---------------- > gcc/tree-ssa-structalias.c | 4 +- > include/ChangeLog.openacc | 5 ++ > include/gomp-constants.h | 12 +++++ > libgomp/ChangeLog.openacc | 6 +++ > libgomp/oacc-parallel.c | 60 ++++++++++++++---------- > 7 files changed, 139 insertions(+), 64 deletions(-) > > diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc > index 718044c..6a51b1e 100644 > --- a/gcc/ChangeLog.openacc > +++ b/gcc/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Thomas Schwinge <tho...@codesourcery.com> > + > + * omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP > + code paths. Update for libgomp OpenACC entry points change. > + > 2018-12-21 Gergö Barany <ge...@codesourcery.com> > > * omp-low.c (scan_sharing_clauses): Fix call to renamed function > diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c > index 988b1bb..ea264da 100644 > --- a/gcc/omp-expand.c > +++ b/gcc/omp-expand.c > @@ -7204,7 +7204,7 @@ expand_omp_target (struct omp_region *region) > transfers. */ > tree t1, t2, t3, t4, device, cond, depend, c, clauses; > enum built_in_function start_ix; > - location_t clause_loc; > + location_t clause_loc = UNKNOWN_LOCATION; > unsigned int flags_i = 0; > > switch (gimple_omp_target_kind (entry_stmt)) > @@ -7249,49 +7249,62 @@ expand_omp_target (struct omp_region *region) > > clauses = gimple_omp_target_clauses (entry_stmt); > > - /* By default, the value of DEVICE is GOMP_DEVICE_ICV (let runtime > - library choose) and there is no conditional. */ > - cond = NULL_TREE; > - device = build_int_cst (integer_type_node, GOMP_DEVICE_ICV); > - > - c = omp_find_clause (clauses, OMP_CLAUSE_IF); > - if (c) > - cond = OMP_CLAUSE_IF_EXPR (c); > - > - c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE); > - if (c) > + device = NULL_TREE; > + tree goacc_flags = NULL_TREE; > + if (is_gimple_omp_oacc (entry_stmt)) > { > - /* Even if we pass it to all library function calls, it is currently > only > - defined/used for the OpenMP target ones. */ > - gcc_checking_assert (start_ix == BUILT_IN_GOMP_TARGET > - || start_ix == BUILT_IN_GOMP_TARGET_DATA > - || start_ix == BUILT_IN_GOMP_TARGET_UPDATE > - || start_ix == BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA); > - > - device = OMP_CLAUSE_DEVICE_ID (c); > - clause_loc = OMP_CLAUSE_LOCATION (c); > + /* By default, no GOACC_FLAGs are set. */ > + goacc_flags = integer_zero_node; > } > else > - clause_loc = gimple_location (entry_stmt); > - > - c = omp_find_clause (clauses, OMP_CLAUSE_NOWAIT); > - if (c) > - flags_i |= GOMP_TARGET_FLAG_NOWAIT; > + { > + c = omp_find_clause (clauses, OMP_CLAUSE_DEVICE); > + if (c) > + { > + device = OMP_CLAUSE_DEVICE_ID (c); > + clause_loc = OMP_CLAUSE_LOCATION (c); > + } > + else > + { > + /* By default, the value of DEVICE is GOMP_DEVICE_ICV (let runtime > + library choose). */ > + device = build_int_cst (integer_type_node, GOMP_DEVICE_ICV); > + clause_loc = gimple_location (entry_stmt); > + } > > - /* Ensure 'device' is of the correct type. */ > - device = fold_convert_loc (clause_loc, integer_type_node, device); > + c = omp_find_clause (clauses, OMP_CLAUSE_NOWAIT); > + if (c) > + flags_i |= GOMP_TARGET_FLAG_NOWAIT; > + } > > - /* If we found the clause 'if (cond)', build > - (cond ? device : GOMP_DEVICE_HOST_FALLBACK). */ > + /* By default, there is no conditional. */ > + cond = NULL_TREE; > + c = omp_find_clause (clauses, OMP_CLAUSE_IF); > + if (c) > + cond = OMP_CLAUSE_IF_EXPR (c); > + /* If we found the clause 'if (cond)', build: > + OpenACC: goacc_flags = (cond ? goacc_flags : flags | > GOACC_FLAG_HOST_FALLBACK) > + OpenMP: device = (cond ? device : GOMP_DEVICE_HOST_FALLBACK) */ > if (cond) > { > + tree *tp; > + if (is_gimple_omp_oacc (entry_stmt)) > + tp = &goacc_flags; > + else > + { > + /* Ensure 'device' is of the correct type. */ > + device = fold_convert_loc (clause_loc, integer_type_node, device); > + > + tp = &device; > + } > + > cond = gimple_boolify (cond); > > basic_block cond_bb, then_bb, else_bb; > edge e; > tree tmp_var; > > - tmp_var = create_tmp_var (TREE_TYPE (device)); > + tmp_var = create_tmp_var (TREE_TYPE (*tp)); > if (offloaded) > e = split_block_after_labels (new_bb); > else > @@ -7314,13 +7327,20 @@ expand_omp_target (struct omp_region *region) > gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); > > gsi = gsi_start_bb (then_bb); > - stmt = gimple_build_assign (tmp_var, device); > + stmt = gimple_build_assign (tmp_var, *tp); > gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); > > gsi = gsi_start_bb (else_bb); > - stmt = gimple_build_assign (tmp_var, > - build_int_cst (integer_type_node, > - GOMP_DEVICE_HOST_FALLBACK)); > + if (is_gimple_omp_oacc (entry_stmt)) > + stmt = gimple_build_assign (tmp_var, > + BIT_IOR_EXPR, > + *tp, > + build_int_cst (integer_type_node, > + GOACC_FLAG_HOST_FALLBACK)); > + else > + stmt = gimple_build_assign (tmp_var, > + build_int_cst (integer_type_node, > + GOMP_DEVICE_HOST_FALLBACK)); > gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); > > make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE); > @@ -7330,14 +7350,17 @@ expand_omp_target (struct omp_region *region) > make_edge (then_bb, new_bb, EDGE_FALLTHRU); > make_edge (else_bb, new_bb, EDGE_FALLTHRU); > > - device = tmp_var; > + *tp = tmp_var; > + > gsi = gsi_last_nondebug_bb (new_bb); > } > else > { > gsi = gsi_last_nondebug_bb (new_bb); > - device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE, > - true, GSI_SAME_STMT); > + > + if (device != NULL_TREE) > + device = force_gimple_operand_gsi (&gsi, device, true, NULL_TREE, > + true, GSI_SAME_STMT); > } > > t = gimple_omp_target_data_arg (entry_stmt); > @@ -7361,7 +7384,17 @@ expand_omp_target (struct omp_region *region) > bool tagging = false; > /* The maximum number used by any start_ix, without varargs. */ > auto_vec<tree, 11> args; > - args.quick_push (device); > + if (is_gimple_omp_oacc (entry_stmt)) > + { > + tree goacc_flags_m = fold_build1 (GOACC_FLAGS_MARSHAL_OP, > + TREE_TYPE (goacc_flags), goacc_flags); > + goacc_flags_m = force_gimple_operand_gsi (&gsi, goacc_flags_m, true, > + NULL_TREE, true, > + GSI_SAME_STMT); > + args.quick_push (goacc_flags_m); > + } > + else > + args.quick_push (device); > if (start_ix == BUILT_IN_GOACC_PARALLEL) > { > tree use_params = oacc_parallel ? integer_one_node : integer_zero_node; > diff --git a/gcc/tree-ssa-structalias.c b/gcc/tree-ssa-structalias.c > index a4f7251..bcf3fd3 100644 > --- a/gcc/tree-ssa-structalias.c > +++ b/gcc/tree-ssa-structalias.c > @@ -4684,7 +4684,7 @@ find_func_aliases_for_builtin_call (struct function > *fn, gcall *t) > argpos = 1; > break; > case BUILT_IN_GOACC_PARALLEL: > - /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs, > + /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs, > sizes, kinds, ...). */ > fnpos = 2; > argpos = 4; > @@ -5263,7 +5263,7 @@ find_func_clobbers (struct function *fn, gimple *origt) > argpos = 1; > break; > case BUILT_IN_GOACC_PARALLEL: > - /* __builtin_GOACC_parallel (device, fn, mapnum, hostaddrs, > + /* __builtin_GOACC_parallel (flags_m, fn, mapnum, hostaddrs, > sizes, kinds, ...). */ > fnpos = 2; > argpos = 4; > diff --git a/include/ChangeLog.openacc b/include/ChangeLog.openacc > index 20ed27f..aa583ea 100644 > --- a/include/ChangeLog.openacc > +++ b/include/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Thomas Schwinge <tho...@codesourcery.com> > + > + * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK) > + (GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define. > + > 2018-12-20 Julian Brown <jul...@codesourcery.com> > Maciej W. Rozycki <ma...@codesourcery.com> > > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index 27de5bc..b5d8441 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -252,6 +252,18 @@ enum gomp_map_kind > /* Internal to libgomp. */ > #define GOMP_TARGET_FLAG_UPDATE (1U << 31) > > + > +/* OpenACC construct flags. */ > + > +/* Force host fallback execution. */ > +#define GOACC_FLAG_HOST_FALLBACK (1 << 0) > + > +/* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted > + bitmask. */ > +#define GOACC_FLAGS_MARSHAL_OP BIT_NOT_EXPR > +#define GOACC_FLAGS_UNMARSHAL(X) (~(X)) > + > + > /* Versions of libgomp and device-specific plugins. GOMP_VERSION > should be incremented whenever an ABI-incompatible change is introduced > to the plugin interface defined in libgomp/libgomp.h. */ > diff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc > index b48453b..04cea5f 100644 > --- a/libgomp/ChangeLog.openacc > +++ b/libgomp/ChangeLog.openacc > @@ -1,3 +1,9 @@ > +2018-12-21 Thomas Schwinge <tho...@codesourcery.com> > + > + * oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel) > + (GOACC_data_start, GOACC_enter_exit_data, GOACC_update) > + (GOACC_declare): Redefine the "device" argument to "flags". > + > 2018-12-20 Gergö Barany <ge...@codesourcery.com> > Thomas Schwinge <tho...@codesourcery.com> > > diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c > index c74221f..0b5f41a 100644 > --- a/libgomp/oacc-parallel.c > +++ b/libgomp/oacc-parallel.c > @@ -41,6 +41,16 @@ > #include <stdarg.h> > #include <assert.h> > > + > +/* In the ABI, the GOACC_FLAGs are encoded as an inverted bitmask, so that we > + continue to support the following two legacy values. */ > +_Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_ICV) == 0, > + "legacy GOMP_DEVICE_ICV broken"); > +_Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_HOST_FALLBACK) > + == GOACC_FLAG_HOST_FALLBACK, > + "legacy GOMP_DEVICE_HOST_FALLBACK broken"); > + > + > /* Returns the number of mappings associated with the pointer or pset. PSET > have three mappings, whereas pointer have two. */ > > @@ -159,17 +169,18 @@ goacc_call_host_fn (void (*fn) (void *), size_t mapnum, > void **hostaddrs, > fn (hostaddrs); > } > > -/* Launch a possibly offloaded function on DEVICE. FN is the host fn > +/* Launch a possibly offloaded function with FLAGS. FN is the host fn > address. MAPNUM, HOSTADDRS, SIZES & KINDS describe the memory > blocks to be copied to/from the device. Varadic arguments are > keyed optional parameters terminated with a zero. */ > > static void > -GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *), > +GOACC_parallel_keyed_internal (int flags_m, int params, void (*fn) (void *), > size_t mapnum, void **hostaddrs, size_t *sizes, > unsigned short *kinds, va_list *ap) > { > - bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; > + int flags = GOACC_FLAGS_UNMARSHAL (flags_m); > + > struct goacc_thread *thr; > struct gomp_device_descr *acc_dev; > struct target_mem_desc *tgt; > @@ -252,7 +263,7 @@ GOACC_parallel_keyed_internal (int device, int params, > void (*fn) (void *), > > /* Host fallback if "if" clause is false or if the current device is set to > the host. */ > - if (host_fallback) > + if (flags & GOACC_FLAG_HOST_FALLBACK) > { > //TODO > prof_info.device_type = acc_device_host; > @@ -448,25 +459,25 @@ GOACC_parallel_keyed_internal (int device, int params, > void (*fn) (void *), > } > > void > -GOACC_parallel_keyed (int device, void (*fn) (void *), > +GOACC_parallel_keyed (int flags_m, void (*fn) (void *), > size_t mapnum, void **hostaddrs, size_t *sizes, > unsigned short *kinds, ...) > { > va_list ap; > va_start (ap, kinds); > - GOACC_parallel_keyed_internal (device, 0, fn, mapnum, hostaddrs, sizes, > + GOACC_parallel_keyed_internal (flags_m, 0, fn, mapnum, hostaddrs, sizes, > kinds, &ap); > va_end (ap); > } > > void > -GOACC_parallel_keyed_v2 (int device, int args, void (*fn) (void *), > +GOACC_parallel_keyed_v2 (int flags_m, int args, void (*fn) (void *), > size_t mapnum, void **hostaddrs, size_t *sizes, > unsigned short *kinds, ...) > { > va_list ap; > va_start (ap, kinds); > - GOACC_parallel_keyed_internal (device, args, fn, mapnum, hostaddrs, sizes, > + GOACC_parallel_keyed_internal (flags_m, args, fn, mapnum, hostaddrs, sizes, > kinds, &ap); > va_end (ap); > } > @@ -474,7 +485,7 @@ GOACC_parallel_keyed_v2 (int device, int args, void (*fn) > (void *), > /* Legacy entry point, only provide host execution. */ > > void > -GOACC_parallel (int device, void (*fn) (void *), > +GOACC_parallel (int flags_m, void (*fn) (void *), > size_t mapnum, void **hostaddrs, size_t *sizes, > unsigned short *kinds, > int num_gangs, int num_workers, int vector_length, > @@ -486,10 +497,11 @@ GOACC_parallel (int device, void (*fn) (void *), > } > > void > -GOACC_data_start (int device, size_t mapnum, > +GOACC_data_start (int flags_m, size_t mapnum, > void **hostaddrs, size_t *sizes, unsigned short *kinds) > { > - bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; > + int flags = GOACC_FLAGS_UNMARSHAL (flags_m); > + > struct target_mem_desc *tgt; > > #ifdef HAVE_INTTYPES_H > @@ -575,7 +587,7 @@ GOACC_data_start (int device, size_t mapnum, > > /* Host fallback or 'do nothing'. */ > if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) > - || host_fallback) > + || (flags & GOACC_FLAG_HOST_FALLBACK)) > { > //TODO > prof_info.device_type = acc_device_host; > @@ -694,13 +706,14 @@ GOACC_data_end (void) > } > > void > -GOACC_enter_exit_data (int device, size_t mapnum, > +GOACC_enter_exit_data (int flags_m, size_t mapnum, > void **hostaddrs, size_t *sizes, unsigned short *kinds, > int async, int num_waits, ...) > { > + int flags = GOACC_FLAGS_UNMARSHAL (flags_m); > + > struct goacc_thread *thr; > struct gomp_device_descr *acc_dev; > - bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; > bool data_enter = false; > size_t i; > > @@ -815,7 +828,7 @@ GOACC_enter_exit_data (int device, size_t mapnum, > &api_info); > > if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) > - || host_fallback) > + || (flags & GOACC_FLAG_HOST_FALLBACK)) > { > //TODO > prof_info.device_type = acc_device_host; > @@ -1098,11 +1111,12 @@ goacc_wait (int async, int num_waits, va_list *ap) > } > > void > -GOACC_update (int device, size_t mapnum, > +GOACC_update (int flags_m, size_t mapnum, > void **hostaddrs, size_t *sizes, unsigned short *kinds, > int async, int num_waits, ...) > { > - bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK; > + int flags = GOACC_FLAGS_UNMARSHAL (flags_m); > + > size_t i; > > goacc_lazy_initialize (); > @@ -1163,7 +1177,7 @@ GOACC_update (int device, size_t mapnum, > goacc_profiling_dispatch (&prof_info, &update_event_info, &api_info); > > if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) > - || host_fallback) > + || (flags & GOACC_FLAG_HOST_FALLBACK)) > { > //TODO > prof_info.device_type = acc_device_host; > @@ -1309,7 +1323,7 @@ GOACC_get_thread_num (void) > } > > void > -GOACC_declare (int device, size_t mapnum, > +GOACC_declare (int flags_m, size_t mapnum, > void **hostaddrs, size_t *sizes, unsigned short *kinds) > { > int i; > @@ -1329,7 +1343,7 @@ GOACC_declare (int device, size_t mapnum, > case GOMP_MAP_POINTER: > case GOMP_MAP_RELEASE: > case GOMP_MAP_DELETE: > - GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], > + GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > &kinds[i], 0, 0); > break; > > @@ -1338,18 +1352,18 @@ GOACC_declare (int device, size_t mapnum, > > case GOMP_MAP_ALLOC: > if (!acc_is_present (hostaddrs[i], sizes[i])) > - GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], > + GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > &kinds[i], 0, 0); > break; > > case GOMP_MAP_TO: > - GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], > + GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > &kinds[i], 0, 0); > > break; > > case GOMP_MAP_FROM: > - GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i], > + GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > &kinds[i], 0, 0); > break; > > -- > 2.8.1 > > From cbd9efcd4ebb6c73a14ead01d85e452d63b7c937 Mon Sep 17 00:00:00 2001 > From: =?UTF-8?q?Gerg=C3=B6=20Barany?= <ge...@codesourcery.com> > Date: Fri, 21 Dec 2018 01:12:44 -0800 > Subject: [PATCH 2/2] [og8] Add OpenACC 2.6 if and if_present clauses on > host_data construct: GOACC_FLAG_HOST_DATA_IF_PRESENT > > 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. > gcc/testsuite/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. > > libgomp/ > * 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. > > libgomp/testsuite/libgomp.oacc-c-c++-common/ > * host_data-6.c: New test. > --- > gcc/ChangeLog.openacc | 5 ++ > gcc/c/ChangeLog.openacc | 5 ++ > gcc/c/c-parser.c | 4 +- > gcc/cp/ChangeLog.openacc | 5 ++ > gcc/cp/parser.c | 4 +- > gcc/fortran/ChangeLog.openacc | 5 ++ > gcc/fortran/openmp.c | 4 +- > gcc/omp-expand.c | 12 ++++- > gcc/testsuite/ChangeLog.openacc | 6 +++ > gcc/testsuite/c-c++-common/goacc/host_data-1.c | 28 +++++++++++- > gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 12 ++++- > include/ChangeLog.openacc | 4 ++ > include/gomp-constants.h | 2 + > libgomp/ChangeLog.openacc | 10 ++++ > libgomp/libgomp.h | 3 ++ > libgomp/oacc-parallel.c | 11 +++-- > libgomp/target.c | 3 ++ > .../libgomp.oacc-c-c++-common/host_data-6.c | 53 > ++++++++++++++++++++++ > 18 files changed, 167 insertions(+), 9 deletions(-) > create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c > > diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc > index 6a51b1e..66eba7b 100644 > --- a/gcc/ChangeLog.openacc > +++ b/gcc/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Gergö Barany <ge...@codesourcery.com> > + > + * omp-expand.c (expand_omp_target): Handle if_present flag on > + OpenACC host_data construct. > + > 2018-12-21 Thomas Schwinge <tho...@codesourcery.com> > > * omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP > diff --git a/gcc/c/ChangeLog.openacc b/gcc/c/ChangeLog.openacc > index 10c00e5..e607ea8 100644 > --- a/gcc/c/ChangeLog.openacc > +++ b/gcc/c/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Gergö Barany <ge...@codesourcery.com> > + > + * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF > + and PRAGMA_OACC_CLAUSE_IF_PRESENT. > + > 2018-12-20 Julian Brown <jul...@codesourcery.com> > Maciej W. Rozycki <ma...@codesourcery.com> > > diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c > index a352d54..2bc4f45 100644 > --- a/gcc/c/c-parser.c > +++ b/gcc/c/c-parser.c > @@ -14876,7 +14876,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.openacc b/gcc/cp/ChangeLog.openacc > index 37b0028..76889c7 100644 > --- a/gcc/cp/ChangeLog.openacc > +++ b/gcc/cp/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Gergö Barany <ge...@codesourcery.com> > + > + * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF > + and PRAGMA_OACC_CLAUSE_IF_PRESENT. > + > 2018-12-20 Julian Brown <jul...@codesourcery.com> > Maciej W. Rozycki <ma...@codesourcery.com> > > diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c > index 083700b..38b0a6d 100644 > --- a/gcc/cp/parser.c > +++ b/gcc/cp/parser.c > @@ -36973,7 +36973,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.openacc b/gcc/fortran/ChangeLog.openacc > index a369219..306871e 100644 > --- a/gcc/fortran/ChangeLog.openacc > +++ b/gcc/fortran/ChangeLog.openacc > @@ -1,3 +1,8 @@ > +2018-12-21 Gergö Barany <ge...@codesourcery.com> > + > + * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and > + OMP_CLAUSE_IF_PRESENT. > + > 2018-12-20 Julian Brown <jul...@codesourcery.com> > Maciej W. Rozycki <ma...@codesourcery.com> > > diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c > index 350f4b1..4273dee 100644 > --- a/gcc/fortran/openmp.c > +++ b/gcc/fortran/openmp.c > @@ -2107,7 +2107,9 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask > mask, > | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ > | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_ATTACH) > #define OACC_HOST_DATA_CLAUSES \ > - (omp_mask (OMP_CLAUSE_USE_DEVICE)) > + (omp_mask (OMP_CLAUSE_USE_DEVICE) \ > + | OMP_CLAUSE_IF \ > + | OMP_CLAUSE_IF_PRESENT) > #define OACC_LOOP_CLAUSES \ > (omp_mask (OMP_CLAUSE_COLLAPSE) \ > | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \ > diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c > index ea264da..42c4910 100644 > --- a/gcc/omp-expand.c > +++ b/gcc/omp-expand.c > @@ -7254,7 +7254,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.openacc b/gcc/testsuite/ChangeLog.openacc > index 473eb9d..2e4bd3d 100644 > --- a/gcc/testsuite/ChangeLog.openacc > +++ b/gcc/testsuite/ChangeLog.openacc > @@ -1,5 +1,11 @@ > 2018-12-21 Gergö Barany <ge...@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-21 Gergö Barany <ge...@codesourcery.com> > + > * c-c++-common/goacc/nested-reductions-fail.c: Renamed to... > * c-c++-common/goacc/nested-reductions-parallel-fail.c: ...this file, > with kernels tests... > 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 0c7a857..658b7a6 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 d44ca58..2ac1c0d 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/ChangeLog.openacc b/include/ChangeLog.openacc > index aa583ea..82058e7 100644 > --- a/include/ChangeLog.openacc > +++ b/include/ChangeLog.openacc > @@ -1,3 +1,7 @@ > +2018-12-21 Gergö Barany <ge...@codesourcery.com> > + > + * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant. > + > 2018-12-21 Thomas Schwinge <tho...@codesourcery.com> > > * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK) > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index b5d8441..953df8f 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -257,6 +257,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.openacc b/libgomp/ChangeLog.openacc > index 04cea5f..7b9e2c5 100644 > --- a/libgomp/ChangeLog.openacc > +++ b/libgomp/ChangeLog.openacc > @@ -1,3 +1,13 @@ > +2018-12-21 Gergö Barany <ge...@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-21 Thomas Schwinge <tho...@codesourcery.com> > > * oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel) > diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h > index 64895c5..11948d5 100644 > --- a/libgomp/libgomp.h > +++ b/libgomp/libgomp.h > @@ -1024,6 +1024,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 0b5f41a..3da87a1 100644 > --- a/libgomp/oacc-parallel.c > +++ b/libgomp/oacc-parallel.c > @@ -585,6 +585,12 @@ 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)) > @@ -592,8 +598,7 @@ GOACC_data_start (int flags_m, size_t mapnum, > //TODO > prof_info.device_type = acc_device_host; > api_info.device_type = prof_info.device_type; > - 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; > goto out; > @@ -601,7 +606,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 0594405..bdfd640 100644 > --- a/libgomp/target.c > +++ b/libgomp/target.c > @@ -1137,6 +1137,9 @@ gomp_map_vars_async (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 > 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; > +} > -- > 2.8.1 >
