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