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
