Re: [PATCH] C, C++, Fortran, OpenMP: Add 'has_device_addr' clause to 'target' construct

[Jakub Jelinek <jakub@redhat.com>]

On Tue, Jan 11, 2022 at 02:27:57PM +0100, Tobias Burnus wrote:
> Hi Jakub, hi all,
> 
> let me quickly comment on 'has_device_addr' with Fortran arrays
> and with an array section (i.e. regarding your comment quoted
> at the very bottom of this email).
> 
> Unfortunately, the OpenMP specification is rather unclear
> what has_device_addr means for C/C++ array sections and in general
> for Fortran, especially when arrays, allocatables/pointers, and
> type parameters like nonconst string lengths are involved. Thus,
> I opened a spec issue – after some discussions (lang-spec meeting,
> C++/affinity (→ Fortran) meeting), it starts to converge:
> https://github.com/OpenMP/spec/issues/3180
> 
> If I understood it correctly, for C/C++, using has_device_addr with
> an array section implies firstprivate, while it does not without
> array section.

That seems just wrong for arrays, that will just crash, see below.

Cases like:
  struct S { whatever; } s;
  #pragma omp target data map (s) use_device_addr (s)
  {
    // At this point it is invalid to use s.field etc. because
    // &s is a device address
    #pragma omp target has_device_addr (s)
    {
      access (&s);
    }
  }
and s/struct S { whatever; } s/int s[16];/
are similar, in all the cases use_device_addr will replace
s in the body with *device_addr_of_s and has_device_addr
needs to firstprivatize the artificial address of s and ensure
that even in the target body s is *some_addr_of_s.
And IMHO array sections should be treated the same, just the
target data could map only parts of the array and not the whole
array, but still device_addr_of_s will be something pointing to the
start of the array, perhaps before the actual object allocated on the
device.
So, I think the gimplifier should strip the ARRAY_REFs and if that yields
something with ARRAY_TYPE, should just treat that var as what appeared
in the has_device_addr clause.  Only if it the array section has
a base pointer that base pointer needs to be copied to the device as is
and so the artificial firstprivate on that pointer that copies the pointer
to the device code.

> Side remark: I note that use_device_addr permits array sections,
> but GCC does not support them yet. (Useful when doing a partial
> map of an array + 'omp data use_device_addr()' on the partially
> mapped array.)

Yes, we should implement that.  But even without that supported, one can
have:
  int a[32];
  #pragma omp target data map (a) use_device_addr (a)
  {
    // So, &a[0] is now a device pointer, whole a is mapped
    #pragma omp target has_device_addr (a[3:17])
    {
      ++a[3];
    }
  }
When whole a[0:32] is mapped, obviously a[3:17] is mapped too
and for has_device-addr it IMHO should act like has_device_addr (a)
under the hood, except the user doesn't guarantee that the whole
array is mapped, just that a has a device address.

Now, if we treat the has_device_addr (a[3:17]) in the above testcase
as firstprivate (a), that will mean we try to copy the whole array from
host to the device.  But &a[0] etc. aren't host addresses, they are device
addresses, so unless the host can access the device addresses, that will
segfault.

	Jakub