Hi! While reviewing <20191003163505.49997-2-julian@codesourcery.com">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 <jul...@codesourcery.com> wrote:
> On Fri, 7 Dec 2018 14:50:19 +0100
> Jakub Jelinek <ja...@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,
<20191003163505.49997-2-julian@codesourcery.com">http://mid.mail-archive.com/20191003163505.49997-2-julian@codesourcery.com>
"OpenACC reference count overhaul".
Grüße
Thomas
signature.asc
Description: PGP signature
