Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)

[Thomas Schwinge <thomas@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 3871 bytes --]

Hi!

While reviewing
<http://mid.mail-archive.com/20191003163505.49997-2-julian@codesourcery.com>
"OpenACC reference count overhaul", I've just now stumbled over one thing
that originally was designed here:

On 2018-12-10T19:41:37+0000, Julian Brown <julian@codesourcery.com> wrote:
> On Fri, 7 Dec 2018 14:50:19 +0100
> Jakub Jelinek <jakub@redhat.com> wrote:
>
>> On Fri, Nov 30, 2018 at 03:41:09AM -0800, Julian Brown wrote:
>> > @@ -918,8 +920,13 @@ struct splay_tree_key_s {
>> >    uintptr_t tgt_offset;
>> >    /* Reference count.  */
>> >    uintptr_t refcount;
>> > -  /* Dynamic reference count.  */
>> > -  uintptr_t dynamic_refcount;
>> > +  /* Reference counts beyond those that represent genuine references in the
>> > +     linked splay tree key/target memory structures, e.g. for multiple OpenACC
>> > +     "present increment" operations (via "acc enter data") refering to the same
>> > +     host-memory block.  */
>> > +  uintptr_t virtual_refcount;
>> > +  /* For a block with attached pointers, the attachment counters for each.  */
>> > +  unsigned short *attach_count;
>> >    /* Pointer to the original mapping of "omp declare target link" object.  */
>> >    splay_tree_key link_key;
>> >  };  
>> 
>> This is something I'm worried about a lot, the nodes keep growing way
>> too much.

Is that just a would-be-nice-to-avoid, or is it an actual problem?

If the latter, can we maybe move some data into on-the-side data
structures, say an associative array keyed by [something suitable]?  I
would assume that compared to actual host to/from device data movement
(or even lookup etc.), lookup of values from such an associative array
should be relatively cheap?

I'm bringing this up, because:

>> Is there a way to reuse some other field if it is of
>> certain kind?
>
> How about this -- it seems that the link_key is only used for OpenMP,

So, is that actually correct?  Per my understanding, for the OpenACC
'link' clause we uses 'GOMP_MAP_LINK', which sets "omp declare target
link", and thus:

> and the attach count is only needed for OpenACC. So the obvious thing
> to do is probably to put those two together into a tagged union. The
> question is where to put the tag?
>
> Options are, I guess:
>
> 1. The high or low bits of the address.  Potentially non-portable, ugly.
>
> 2. Or, the virtual refcount is also only needed for OpenACC, so we can
>    reserve a magic value for that field to act as a tag.
>
> I've tried implementing the latter in the attached patch, and it seems
> to work OK.

... this is not actually feasible?

It's certainly possible that we're totally lacking sufficient testsuite
coverage, and that there are issues in the 'link' implementation
(<https://gcc.gnu.org/PR81689> "libgomp.c/target-link-1.c fails for
nvptx: #pragma omp target link not implemented" comes to mind
immediatelly, and certainly for OpenACC I used to be aware of additional
issues; I think I intended to use that mechanism for Fortran
'allocatable' with OpenACC 'declare'), but the libgomp handling to me
seems reasonable upon quick review -- just that we need to keep it alive
for OpenACC, too, unless I'm confused?


Simplifying the libgomp code to avoid the 'VREFCOUNT_LINK_KEY' toggle
flag, and not putting 'link_key' into an union together with
'attach_count', that should -- I hope -- resolve/obsolete some of the
questions raised in my late-night pre-Christmas 2018 review,
<https://gcc.gnu.org/ml/gcc-patches/2018-12/msg01620.html>, where I'm now
not sure yet whether all my questions have been addressed (or disputed,
but I didn't hear anything) in the recent -- split-out, thanks! --
version of this patch,
<http://mid.mail-archive.com/20191003163505.49997-2-julian@codesourcery.com>
"OpenACC reference count overhaul".


Grüße
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]