On Wed, 18 Dec 2019 10:18:14 +0100 Thomas Schwinge <tho...@codesourcery.com> wrote:
> Hi Julian! > > Thanks for walking me through this. > > On 2019-12-14T00:19:04+0000, Julian Brown <jul...@codesourcery.com> > wrote: > > On Fri, 13 Dec 2019 16:25:25 +0100 > > Thomas Schwinge <tho...@codesourcery.com> wrote: > >> On 2019-10-29T12:15:01+0000, Julian Brown <jul...@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