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

Reply via email to