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.
For Fortran, the discussion converged to
* array descriptor must be the same at beginning/end of target
region, i.e:
- if allocated before target region, allocation may not be changed
inside the target region. If unallocated, it can be allocated but
must be deallocated before the target region is left.
- if a pointer, pointer association must be the same at the end of
the target region as when it was entered (unmodified or reset)
* meta data (array bounds, alloc status, len type parameters like
char length) are accessible on the host - and are then made
accessible to the device (firstprivate, direct access on shared
memory or whatever means. Due to the conditions above, it is
indistinguishable).
However, the discussion has not yet fully settled and spec updates
are still needed.
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.)
I have not checked what Marcel's implementation does and whether
that's compatible with the incomplete OpenMP 5.{1,2} spec and
the on-going discussion of/in OpenMP lang-spec Issue 3180.
On 11.01.22 12:53, Jakub Jelinek via Fortran wrote:
+++ b/gcc/fortran/trans-openmp.c
...
+ /* For HAS_DEVICE_ADDR of an array descriptor, firstprivatize the
+ descriptor such that the bounds are available; its data component
+ is unmodified; it is handled as device address inside target. */
...
Not sure about the above,
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
...
but this looks weird.
If decl after stripping the ARRAY_REFs is a var with pointer type, sure,
firstprivatizing it is the way to go.
But it can be also a variable with ARRAY_TYPE, can't it? Something like:
int a[64];
#pragma omp target data map(a) use_device_addr(a)
{
#pragma omp target has_device_addr(a[3:16])
a[3] = 1;
}
and in this case firstprivatization of a looks wrong. use_device_addr
should replace (but only at omp-low.c time I think) a used in the block
with the remapped a (i.e. *device_address_of_a).
Or perhaps it could be a non-static data member with array type
inside of a C++ method.
Tobias
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
