This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [gomp4.1] Support #pragma omp target {enter,exit} data
- From: Ilya Verbin <iverbin at gmail dot com>
- To: Jakub Jelinek <jakub at redhat dot com>
- Cc: gcc-patches at gcc dot gnu dot org, Kirill Yukhin <kirill dot yukhin at gmail dot com>
- Date: Mon, 6 Jul 2015 21:45:30 +0300
- Subject: Re: [gomp4.1] Support #pragma omp target {enter,exit} data
- Authentication-results: sourceware.org; auth=none
- References: <20150630121930 dot GA27446 at msticlxl57 dot ims dot intel dot com> <20150630125702 dot GI10247 at tucnak dot redhat dot com> <20150630154201 dot GB27446 at msticlxl57 dot ims dot intel dot com> <20150630161044 dot GM10247 at tucnak dot redhat dot com> <20150701210658 dot GA51887 at msticlxl57 dot ims dot intel dot com> <20150706153425 dot GA52133 at msticlxl57 dot ims dot intel dot com> <20150706172509 dot GY10247 at tucnak dot redhat dot com>
On Mon, Jul 06, 2015 at 19:25:09 +0200, Jakub Jelinek wrote:
> On Mon, Jul 06, 2015 at 06:34:25PM +0300, Ilya Verbin wrote:
> > On Thu, Jul 02, 2015 at 00:06:58 +0300, Ilya Verbin wrote:
> > > 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?
>
> What exactly do you have in mind here?
>
> void foo (int *p)
> {
> #pragma omp enter data (to:p[10])
> ...
> #pragma omp exit data (from:p[10])
> }
>
> where the latter will only deallocate &p[0] ... &p[9], but not &p?
> I've asked for clarification in that case, but if it should deallocate (or
> decrease the counter) for &p too, then I think this is something for the
> frontends to handle during handling of array sections in map clause, or
> during gimplification or omp lowering.
I mean, in enter data map(to:p[10]):
1. Map GOMP_MAP_TO var as usual, and save returned target_mem_desc *tgt_var into
last_tgt_var.
2. Map GOMP_MAP_POINTER var, and save returned tgt_var->list[0].key into
last_tgt_var->new_special_field_for_pointer.
And in exit data map(from:p[10]):
1. Unmap GOMP_MAP_FROM var as usual, *and* deallocate (or decrease refcount) of
k->tgt->new_special_field_for_pointer.
2. Do nothing for GOMP_MAP_POINTER var.
But I don't like this plan, there may be corner cases.
-- Ilya