[Bug middle-end/95236] New: OMP 'GOMP_MAP_STRUCT': a structure is more than the sum of all its fields?

[Bug middle-end/95236] New: OMP 'GOMP_MAP_STRUCT': a structure is more than the sum of all its fields?

[tschwinge at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=95236

            Bug ID: 95236
           Summary: OMP 'GOMP_MAP_STRUCT': a structure is more than the
                    sum of all its fields?
           Product: gcc
           Version: 10.0
            Status: UNCONFIRMED
          Keywords: openacc, openmp
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tschwinge at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org, jules at gcc dot gnu.org
  Target Milestone: ---

Created attachment 48571
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=48571&action=edit
'itu_.c'

..., and the answer (technically) may easily be "yes", considering ABI-mandated
padding, for example.

But: is that user friendly?  Consider the test case I'm attaching, where
'struct s' consists of fields 'a', 'b', which both get mapped -- yet, 's'
itself isn't considered to be.  At the user level, should a structure equal the
sum of all its fields, or "do we have to be complicated"?

I have not fully thought that through, just noticed this when looking into
something else.  I also have not verified OpenMP behavior.

[Bug middle-end/95236] OMP 'GOMP_MAP_STRUCT': a structure is more than the sum of all its fields?

[julian at codesourcery dot com]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=95236

--- Comment #1 from julian at codesourcery dot com ---
On Wed, 20 May 2020 10:42:23 +0000
"tschwinge at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org> wrote:

> https://gcc.gnu.org/bugzilla/show_bug.cgi?id=95236
> 
>             Bug ID: 95236
>            Summary: OMP 'GOMP_MAP_STRUCT': a structure is more than
> the sum of all its fields?
>            Product: gcc
>            Version: 10.0
>             Status: UNCONFIRMED
>           Keywords: openacc, openmp
>           Severity: normal
>           Priority: P3
>          Component: middle-end
>           Assignee: unassigned at gcc dot gnu.org
>           Reporter: tschwinge at gcc dot gnu.org
>                 CC: jakub at gcc dot gnu.org, jules at gcc dot gnu.org
>   Target Milestone: ---
> 
> Created attachment 48571
>   -->
> https://gcc.gnu.org/bugzilla/attachment.cgi?id=48571&action=edit
> 'itu_.c'
> 
> ..., and the answer (technically) may easily be "yes", considering
> ABI-mandated padding, for example.
> 
> But: is that user friendly?  Consider the test case I'm attaching,
> where 'struct s' consists of fields 'a', 'b', which both get mapped
> -- yet, 's' itself isn't considered to be.  At the user level, should
> a structure equal the sum of all its fields, or "do we have to be
> complicated"?
> 
> I have not fully thought that through, just noticed this when looking
> into something else.  I also have not verified OpenMP behavior.

The ability to partially copy the struct to the target is deliberate
(and inherited from the OpenMP implementation) since the struct could
potentially be large, and the user might not need all of it on the
target. That's hand-wavingly similar to the ability to copy just a
slice of an array to the target, rather than the whole thing (except an
array slice covering the whole array *is* in fact the array, so will
show as present as such). I'm not sure we'd want to lose the ability to
do such partial mappings for structs.

I suppose the ideal -- which we don't entirely meet, because sometimes
"present(...)" mappings are needed -- is for arrays in structs to behave
the same as arrays which are not in structs, from the user perspective,
without extra cognitive load. So these examples should work the same:

int main(void)
{
  int *a = malloc (sizeof (int) * N));
  int *b = malloc (sizeof (int) * N));

  #pragma acc parallel copy(a[0:N], b[0:N])
  {
    a[x] = ...;
    b[y] = ...;
  }
}

vs:

int main(void)
{
  struct {
    int *a;
    int *b;
  } mystruct;

  mystruct s;

  s.a = malloc (sizeof (int) * N);
  s.b = malloc (sizeof (int) * N):

  #pragma acc parallel copy(s.a[0:N], s.b[0:N])
  {
    s.a[x] = ...;
    s.b[y] = ...;
  }
}

The problem (from an implementation perspective) of an implicit
"copy(s)" not working in your example, or if the above code is
modified like this:

  #pragma acc data copy(s.a[0:N], s.b[0:N])
  {
    #pragma acc parallel
    {
      s.a[x] = ...;
      s.b[y] = ...;
    }
  }

is that s.a and s.b are mapped as separate splay tree keys, and
the implicit (or an explicit) "present(s)" has no way of union'ing those
together to test if the whole struct is mapped. Maybe that could be
solved as part of support for allowing multiple slices of the same
array to be mapped at once, if or when we want to support that.

Thanks,

Julian