On Tue, Jun 23, 2015 at 02:40:43PM +0300, Ilya Verbin wrote:
> On Sat, Jun 20, 2015 at 00:35:14 +0300, Ilya Verbin wrote:
> > Given that a mapped variable in 4.1 can have different kinds across nested
> > data
> > regions, we need to store map-type not only for each var, but also for each
> > structured mapping. Here is my WIP patch, is it sane? :)
> > Attached testcase works OK on the device with non-shared memory.
>
> A bit updated version with a fix for GOMP_MAP_TO_PSET.
> make check-target-libgomp passed.
Ok, thanks.
> include/gcc/
> * gomp-constants.h (GOMP_MAP_ALWAYS_TO_P,
> GOMP_MAP_ALWAYS_FROM_P): Define.
> libgomp/
> * libgomp.h (struct target_var_desc): New.
> (struct target_mem_desc): Replace array of splay_tree_key with array of
> target_var_desc.
> (struct splay_tree_key_s): Move copy_from to target_var_desc.
> * oacc-mem.c (gomp_acc_remove_pointer): Use copy_from from
> target_var_desc.
> * oacc-parallel.c (GOACC_parallel): Use copy_from from target_var_desc.
> * target.c (gomp_map_vars_existing): Copy data to device if map-type is
> 'always to' or 'always tofrom'.
> (gomp_map_vars): Use key from target_var_desc. Set copy_from and
> always_copy_from.
> (gomp_copy_from_async): Use key and copy_from from target_var_desc.
> (gomp_unmap_vars): Copy data from device if always_copy_from is set.
> (gomp_offload_image_to_device): Do not use copy_from.
> * testsuite/libgomp.c/target-11.c: New test.
> + /* Set dd on target to 0 for the further check. */
> + #pragma omp target map(always to: dd)
> + { dd; }
This reminds me that:
if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
remove = true;
in gimplify.c is not what we want, if it is has GOMP_MAP_KIND_ALWAYS,
then we shouldn't remove it even when it is not mentioned inside of the
region's body, because it then has side-effects.
Jakub
