[PATCH 07/13] OpenACC 2.6 deep copy: libgomp parts

[Thomas Schwinge]

Hi!

On 2019-12-17T22:03:47-0800, Julian Brown <julian@codesourcery.com> wrote:
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c

>  static int
> -find_group_last (int pos, size_t mapnum, unsigned short *kinds)
> +find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
>  {
>    unsigned char kind0 = kinds[pos] & 0xff;
> -  int first_pos = pos, last_pos = pos;
> +  int first_pos = pos;
>
> -  if (kind0 == GOMP_MAP_TO_PSET)
> +  switch (kind0)
>      {
> +    case GOMP_MAP_TO_PSET:
>        while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> -     last_pos = ++pos;
> +     pos++;
>        /* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET.  */
> -      assert (last_pos > first_pos);
> -    }
> -  else
> -    {
> +      assert (pos > first_pos);
> +      break;
> +
> +    case GOMP_MAP_STRUCT:
> +      pos += sizes[pos];
> +      break;
> +
> +    case GOMP_MAP_POINTER:
> +    case GOMP_MAP_ALWAYS_POINTER:
> +      /* These mappings are only expected after some other mapping.  If we
> +      see one by itself, something has gone wrong.  */
> +      gomp_fatal ("unexpected mapping");
> +      break;
> +
> +    default:
>        /* 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;
> +      if (pos + 1 < mapnum)
> +     {
> +       unsigned char kind1 = kinds[pos + 1] & 0xff;
> +       if (kind1 == GOMP_MAP_ALWAYS_POINTER)
> +         return pos + 1;
> +     }
>
> -      /* We can have one or several GOMP_MAP_POINTER mappings after a to/from
> +      /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from
>        (etc.) mapping.  */
>        while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
> -     last_pos = ++pos;
> +     pos++;
>      }
>
> -  return last_pos;
> +  return pos;
>  }

So this now causes grouped (!) mapping of all of 'GOMP_MAP_STRUCT', that
is, all its "members" at once.

This, I suppose, mandated the removal of (some of) the 'is_tgt_unmapped'
checking (unfortunately committed not here, but as part of r279621
"OpenACC reference count overhaul"), where we had unmapping code
(conceptually) similar to:

    bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
    assert (is_tgt_unmapped);

I'd introduced this a little bit earlier, finding this a simple yet
effective run-time, low-overhead consistency checking of (certain aspects
of) reference counting -- so just noting here that it's somewhat bad that
we can't have this anymore "just" because of 'GOMP_MAP_STRUCT'.  (Maybe
there is a way to get it back; that's for later?)

Anyway, the code changes were incomplete, consider:

    #include <assert.h>
    #include <openacc.h>

    struct s
    {
      char a;
      char b;
    };

    int main ()
    {
      struct s s;

    #pragma acc enter data create(s.a, s.b)
      assert (acc_is_present (&s.a, sizeof s.a));
      assert (acc_is_present (&s.b, sizeof s.b));

    #if 0
      // works
    # pragma acc exit data delete(s.a)
    # pragma acc exit data delete(s.b)
    #else
      acc_delete (&s.a, sizeof s.a); // fails
      acc_delete (&s.b, sizeof s.b);
    #endif
      assert (!acc_is_present (&s.a, sizeof s.a));
      assert (!acc_is_present (&s.b, sizeof s.b));

      return 0;
    }

The 'acc_delete' variant exercises a code path that still contains the
'is_tgt_unmapped' checking, and that triggers for 'acc_delete ([s.a])'
then, as 's.a' has been mapped in one group together with 's.b', which
remains mapped until 'acc_delete ([s.b])', and thus
"[...]/libgomp/oacc-mem.c:726: goacc_exit_datum: Assertion
`is_tgt_unmapped' failed".  Unless anybody has a better plan for this
whole topic, I'll commit the obvious (removal), later (next week).

    --- libgomp/oacc-mem.c
    +++ libgomp/oacc-mem.c
    @@ -724,7 +724,11 @@ goacc_exit_datum (void *h, size_t s, unsigned short kind, int async)
           else
        {
          bool is_tgt_unmapped = gomp_remove_var (acc_dev, n);
    +#if 0
          assert (is_tgt_unmapped);
    +#else
    +     (void) is_tgt_unmapped;
    +#endif
        }
         }


Grüße
 Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter