[gomp4.1] Support #pragma omp target {enter,exit} data

[Ilya Verbin]

On Thu, Jul 02, 2015 at 00:06:58 +0300, Ilya Verbin wrote:
> On Tue, Jun 30, 2015 at 18:10:44 +0200, Jakub Jelinek wrote:
> > The thing is whether it is actually a good idea to allocate the enter data
> > allocated objects together.
> > In OpenMP 4.0, generally objects would be allocated and deallocated at the
> > same times, except for multiple host threads trying to map the same variables
> > into the target.  In OpenMP 4.1, due to enter data/exit data, they can be
> > allocated and freed quite independently, and it is true that is the case
> > even for target data, one can either target data, then target enter data
> > to prevent something from being deallocated, then target data end freeing
> > only parts, etc.  So the question is if we think in real-world the
> > allocation or deallocation will be usually together or not.
> 
> IMHO, it's OK to allocate "target data" objects together and "target enter data"
> objects one by one.  I've implemented this approach in the patch bellow.
> 
> However, if someone writes a program like this:
> 
>   #pragma omp target data map(tofrom: small, arr[:big])
>     {
>       #pragma omp target enter data map(to: small)
>     }
>   do_a_lot_of_something ();
>   #pragma omp target exit data map(from: small)
> 
> Big array will be deallocated on target only with 'small' at the end.
> Is this acceptable?

Ping?

> The patch is not ready though, I don't know how to unmap GOMP_MAP_POINTER vars.
> In gomp_unmap_vars they're unmapped through tgt->list[], but in gomp_exit_data
> it's impossible to find such var in the splay tree, because hostaddr differs
> from the address, used at mapping.

I can keep a splay_tree_key of the GOMP_MAP_POINTER in the new field in
target_mem_desc of the previous var (i.e. corresponding memory block).
Or could you suggest a better approach?

Thanks,
  -- Ilya