'find_group_last' (was: [PATCH] OpenACC reference count overhaul)

[Julian Brown]

On Wed, 18 Dec 2019 10:18:14 +0100
Thomas Schwinge <thomas@codesourcery.com> wrote:

> Hi Julian!
> 
> Thanks for walking me through this.
> 
> On 2019-12-14T00:19:04+0000, Julian Brown <julian@codesourcery.com>
> wrote:
> > On Fri, 13 Dec 2019 16:25:25 +0100
> > Thomas Schwinge <thomas@codesourcery.com> wrote:  
> >> On 2019-10-29T12:15:01+0000, Julian Brown <julian@codesourcery.com>
> >> wrote:  
> >> >  static int
> >> > -find_pointer (int pos, size_t mapnum, unsigned short *kinds)
> >> > +find_group_last (int pos, size_t mapnum, unsigned short *kinds)
> >> >  {
> >> > -  if (pos + 1 >= mapnum)
> >> > -    return 0;
> >> > +  unsigned char kind0 = kinds[pos] & 0xff;
> >> > +  int first_pos = pos, last_pos = pos;
> >> >  
> >> > -  unsigned char kind = kinds[pos+1] & 0xff;
> >> > -
> >> > -  if (kind == GOMP_MAP_TO_PSET)
> >> > -    return 3;
> >> > -  else if (kind == GOMP_MAP_POINTER)
> >> > -    return 2;
> >> > +  if (kind0 == GOMP_MAP_TO_PSET)
> >> > +    {
> >> > +      while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) ==
> >> > GOMP_MAP_POINTER)
> >> > +	last_pos = ++pos;
> >> > +      /* We expect at least one GOMP_MAP_POINTER after a
> >> > GOMP_MAP_TO_PSET.  */
> >> > +      assert (last_pos > first_pos);
> >> > +    }
> >> > +  else
> >> > +    {
> >> > +      /* GOMP_MAP_ALWAYS_POINTER can only appear directly after
> >> > some other
> >> > +	 mapping.  */
> >> > +      if (pos + 1 < mapnum
> >> > +	  && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)
> >> > +	return pos + 1;
> >> > +
> >> > +      /* We can have one or several GOMP_MAP_POINTER mappings
> >> > after a to/from
> >> > +	 (etc.) mapping.  */
> >> > +      while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) ==
> >> > GOMP_MAP_POINTER)
> >> > +	last_pos = ++pos;
> >> > +    }
> >> >  
> >> > -  return 0;
> >> > +  return last_pos;
> >> >  }    
> 
> Given:
> 
>     program test
>       implicit none
>     
>       integer, parameter :: n = 64
>       integer :: a(n)
>     
>       call test_array(a)
>     
>     contains
>       subroutine test_array(a)
>         implicit none
>     
>         integer :: a(n)
>     
>         !$acc enter data copyin(a)
>     
>         !$acc exit data delete(a)
>       end subroutine test_array
>     end program test
> 
> ..., we get a 'GOMP_MAP_TO' followed by a 'GOMP_MAP_POINTER'.  That
> got us 'find_pointer () == 2', and now we get 'find_group_last (i) ==
> i + 1' (so, the same).
> 
> > In a previous iteration of the refcount overhaul patch, we had the
> > "magic" code fragment:
> >  
> >> +	      for (int j = 0; j < 2; j++)  
> >> +		gomp_map_vars_async (acc_dev, aq,
> >> +				     (j == 0 || pointer == 2) ?
> >> 1 : 2,
> >> +				     &hostaddrs[i + j], NULL,
> >> +				     &sizes[i + j], &kinds[i +
> >> j], true,
> >> +
> >> GOMP_MAP_VARS_OPENACC_ENTER_DATA);    
> 
> > The "pointer == 2" case (i.e. with a GOMP_MAP_TO and a
> > GOMP_MAP_POINTER)  
> 
> So, that's the example given above.
> 
> > will also handle the mappings separately in both the
> > earlier patch iteration  
> 
> ACK, given the "previous iteration" code presented above.
> 
> > and this one.  
> 
> NACK?  Given 'find_group_last (i) == i + 1', that means that
> 'GOMP_MAP_TO' and 'GOMP_MAP_POINTER' get mapped as one group?
> 
> On the other hand, it still does match the current 'find_pointer'
> behavior?
> 
> But what should the behavior here be: 'GOMP_MAP_TO',
> 'GOMP_MAP_POINTER' each separate, or as one group?
> 
> Confusing stuff.  :-|

Hmm.

I think that GOMP_MAP_POINTER is only intended to be used after some
other mapping (TO/TOFROM/TO_PSET/etc.). In the follow-up patch
supporting deep copy, this code is extended and refactored a little
more:

https://gcc.gnu.org/ml/gcc-patches/2019-12/msg01256.html

One of the changes made there is to disallow GOMP_MAP{,_ALWAYS}_POINTER
from appearing by itself. By my reading, that must be the case for
GOMP_MAP_ALWAYS_POINTER because it has a hard-wired dependency on the
previous mapping. GOMP_MAP_POINTER is slightly more questionable: at
least according to the comment in gomp-constants.h, these are "an
internal only map kind, used for pointer based array sections" -- so
it's a little surprising they now reach the libgomp runtime at all.
Maybe it was a mistake?

The GOMP_MAP_ATTACH mapping (as in the example upthread) is different --
that one *can* appear by itself. Perhaps the difference (wrt. reference
counting here) is that GOMP_MAP_POINTER refers to the same
target_mem_desc as the previous (grouped-together) mapping, but
GOMP_MAP_ATTACH does not (rather, referring to the location of the
*pointer* to the data of a previous mapping, rather than the data
itself).

For GOMP_MAP_TO_PSET, a subsequent GOMP_MAP_POINTER will refer to the
pointer set itself. So, same thing, and it's not problematic to group
the mappings together.

Anyway: thinking about it some more, I don't think any of the ways
these types of mappings get grouped together should really be causing
refcount-checking failures, so maybe something's wrong (at least
academically) in goacc_exit_data_internal. The "real" problem with
parasitical groupings is if we have multiple "enter data" mappings that
get bound together in a single target_mem_desc, and are unmapped at
different times:

#pragma acc enter data copyin(arr1) copyin(arr2)
...
#pragma acc exit data copyout(arr1)
#pragma acc exit data copyout(arr2)

That's clearly not what's happening here though.

I will investigate further.

Thanks,

Julian