This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

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


On Tue, Jun 30, 2015 at 06:42:01PM +0300, Ilya Verbin wrote:
> On Tue, Jun 30, 2015 at 14:57:02 +0200, Jakub Jelinek wrote:
> > On Tue, Jun 30, 2015 at 03:19:30PM +0300, Ilya Verbin wrote:
> > > --- a/libgomp/target.c
> > > +++ b/libgomp/target.c
> > > @@ -580,10 +581,16 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
> > >        bool do_unmap = false;
> > >        if (k->refcount > 1)
> > >  	k->refcount--;
> > > -      else if (k->async_refcount > 0)
> > > -	k->async_refcount--;
> > > -      else
> > > -	do_unmap = true;
> > > +      else if (k->refcount == 1)
> > > +	{
> > > +	  if (k->async_refcount > 0)
> > > +	    k->async_refcount--;
> > > +	  else
> > > +	    {
> > > +	      k->refcount--;
> > > +	      do_unmap = true;
> > > +	    }
> > > +	}
> > 
> > What is the rationale of this hunk change?
> 
> Without whis change, when k->refcount == 1, do_unmap is true, but refcount is
> not decremented.  So, if gomp_unmap_vars is called multiple times (now it's
> possible for 4.1), refcount will remain 1, and it will try to unmap k at each
> next call, that is wrong.  That's why I decrement refcount to zero, and do
> nothing when hit gomp_unmap_vars next time with k->refcount == 0.

Ok.

> > >    if (is_enter_data)
> > > -    {
> > > -      /* TODO  */
> > > -    }
> > > +    gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, true, false);
> > 
> > This will leak the return value.  Either we need to arrange not to allocate
> > it for enter data, or we need to assign it to some variable and free
> > immediately (we don't want to perform the release operations for it).
> 
> But we can't not allocate or free immediately it, since it's used later through
> splay_tree_key_s::tgt, e.g. here:
> 
>   if (is_target)
>     {
>       for (i = 0; i < mapnum; i++)
> 	{
> 	  if (tgt->list[i].key == NULL)
> 	    cur_node.tgt_offset = (uintptr_t) NULL;
> 	  else
> 	    cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
> 				  + tgt->list[i].key->tgt_offset;

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.

	Jakub


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]