[v3,06/11] OpenMP: Pointers and member mappings
Commit Message
This patch was previously posted as part of the series supporting
"declare mapper" for Fortran, here:
https://gcc.gnu.org/pipermail/gcc-patches/2022-June/596041.html
Implementing the "omp declare mapper" functionality, I noticed some
cases where handling of derived type members that are pointers doesn't
seem to be quite right. At present, a type such as this:
type T
integer, pointer, dimension(:) :: arrptr
end type T
type(T) :: tvar
[...]
!$omp target map(tofrom: tvar%arrptr)
will be mapped using three mapping nodes:
GOMP_MAP_TO tvar%arrptr (the descriptor)
GOMP_MAP_TOFROM *tvar%arrptr%data (the actual array data)
GOMP_MAP_ALWAYS_POINTER tvar%arrptr%data (a pointer to the array data)
This follows OMP 5.0, 2.19.7.1 "map Clause":
"If a list item in a map clause is an associated pointer and the
pointer is not the base pointer of another list item in a map clause
on the same construct, then it is treated as if its pointer target
is implicitly mapped in the same clause. For the purposes of the map
clause, the mapped pointer target is treated as if its base pointer
is the associated pointer."
However, we can also write this:
map(to: tvar%arrptr) map(tofrom: tvar%arrptr(3:8))
and then instead we should follow:
"If the structure sibling list item is a pointer then it is treated
as if its association status is undefined, unless it appears as
the base pointer of another list item in a map clause on the same
construct."
But, that's not implemented quite right at the moment (and completely
breaks once we introduce declare mappers), because we still map the "to:
tvar%arrptr" as the descriptor and the entire array, then we map the
"tvar%arrptr(3:8)" part using the descriptor (again!) and the array slice.
The solution is to detect when we're mapping a smaller part of the array
(or a subcomponent) on the same directive, and only map the descriptor
in that case. So we get mappings like this instead:
map(to: tvar%arrptr) -->
GOMP_MAP_ALLOC tvar%arrptr (the descriptor)
map(tofrom: tvar%arrptr(3:8) -->
GOMP_MAP_TOFROM tvar%arrptr%data(3) (size 8-3+1, etc.)
GOMP_MAP_ALWAYS_POINTER tvar%arrptr%data (bias 3, etc.)
2022-09-13 Julian Brown <julian@codesourcery.com>
gcc/fortran/
* trans-openmp.cc (dependency.h): Include.
(gfc_trans_omp_array_section): Do not map descriptors here for OpenMP.
(gfc_trans_omp_clauses): Check subcomponent and subarray/element
accesses elsewhere in the clause list for pointers to derived types or
array descriptors, and map just the pointer/descriptor if we have any.
libgomp/
* testsuite/libgomp.fortran/map-subarray.f90: New test.
* testsuite/libgomp.fortran/map-subcomponents.f90: New test.
* testsuite/libgomp.fortran/struct-elem-map-1.f90: Adjust for
descriptor-mapping changes.
---
gcc/fortran/trans-openmp.cc | 104 ++++++++++++++++--
.../libgomp.fortran/map-subarray.f90 | 33 ++++++
.../libgomp.fortran/map-subcomponents.f90 | 32 ++++++
.../libgomp.fortran/struct-elem-map-1.f90 | 10 +-
4 files changed, 163 insertions(+), 16 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/map-subcomponents.f90
Comments
On Tue, Sep 13, 2022 at 02:03:16PM -0700, Julian Brown wrote:
> @@ -3440,6 +3437,50 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> {
> if (pointer || (openacc && allocatable))
> {
> + gfc_omp_namelist *n2
> + = openacc ? NULL : clauses->lists[OMP_LIST_MAP];
> +
> + /* If the last reference is a pointer to a derived
> + type ("foo%dt_ptr"), check if any subcomponents
> + of the same derived type member are being mapped
> + elsewhere in the clause list ("foo%dt_ptr%x",
> + etc.). If we have such subcomponent mappings,
> + we only create an ALLOC node for the pointer
> + itself, and inhibit mapping the whole derived
> + type. */
> +
> + for (; n2 != NULL; n2 = n2->next)
> + {
> + if (n == n2 || !n2->expr)
> + continue;
> +
> + int dep
> + = gfc_dep_resolver (n->expr->ref, n2->expr->ref,
> + NULL, true);
> + if (dep == 0)
> + continue;
Isn't this and the other loop quadratic compile time in number of clauses?
Could it be done linearly through some perhaps lazily built hash table or
something similar?
Jakub
On Wed, 14 Sep 2022 14:53:54 +0200
Jakub Jelinek <jakub@redhat.com> wrote:
> On Tue, Sep 13, 2022 at 02:03:16PM -0700, Julian Brown wrote:
> > @@ -3440,6 +3437,50 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> > gfc_omp_clauses *clauses, {
> > if (pointer || (openacc && allocatable))
> > {
> > + gfc_omp_namelist *n2
> > + = openacc ? NULL :
> > clauses->lists[OMP_LIST_MAP]; +
> > + /* If the last reference is a pointer to
> > a derived
> > + type ("foo%dt_ptr"), check if any
> > subcomponents
> > + of the same derived type member are
> > being mapped
> > + elsewhere in the clause list
> > ("foo%dt_ptr%x",
> > + etc.). If we have such subcomponent
> > mappings,
> > + we only create an ALLOC node for the
> > pointer
> > + itself, and inhibit mapping the whole
> > derived
> > + type. */
> > +
> > + for (; n2 != NULL; n2 = n2->next)
> > + {
> > + if (n == n2 || !n2->expr)
> > + continue;
> > +
> > + int dep
> > + = gfc_dep_resolver (n->expr->ref,
> > n2->expr->ref,
> > + NULL, true);
> > + if (dep == 0)
> > + continue;
>
> Isn't this and the other loop quadratic compile time in number of
> clauses? Could it be done linearly through some perhaps lazily built
> hash table or something similar?
How about this -- a hash table is used to split up the list by the root
symbol the first time we try to resolve dependencies. This arguably
doesn't get rid of the potential for quadratic behaviour *completely*,
but I think the input where it would still be troublesome would be
rather pathological.
I'm not sure how to do better without reimplementing gfc_dep_resolver
or implementing general hashing for Fortran FE expression/symbol nodes,
which isn't there at the moment, so far as I can tell.
I noticed during testing that one of the new added tests for this patch
doesn't actually work until the next patch is in, i.e.:
https://gcc.gnu.org/pipermail/gcc-patches/2022-September/601558.html
So possibly the patches should be committed squashed together. I
XFAILed the test for now.
Re-tested with offloading to NVPTX. OK?
Thanks,
Julian
On Sun, Sep 18, 2022 at 08:19:29PM +0100, Julian Brown wrote:
> @@ -2609,6 +2672,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
> if (clauses == NULL)
> return NULL_TREE;
>
> + hash_map<gfc_symbol *, gfc_omp_namelist *> sym_rooted_nl;
Isn't hash_map ctor pretty costly (allocates memory etc.)?
And gfc_trans_omp_clauses is called for all OpenMP constructs, in many
cases they are never going to have any map clauses or even if they do,
they might not trigger this code.
> + bool built_sym_hash = false;
So, I think usually we don't construct such hash_maps right away,
but have just pointer to the hash map initialized to NULL (then you
don't need to built_sym_hash next to it) and you simply new the hash_map
when needed the first time and delete it at the end (which does nothing
if it is NULL).
Jakub
On Thu, 22 Sep 2022 15:17:08 +0200
Jakub Jelinek <jakub@redhat.com> wrote:
> > + bool built_sym_hash = false;
>
> So, I think usually we don't construct such hash_maps right away,
> but have just pointer to the hash map initialized to NULL (then you
> don't need to built_sym_hash next to it) and you simply new the
> hash_map when needed the first time and delete it at the end (which
> does nothing if it is NULL).
How about this version? (Re-tested.)
Thanks,
Julian
On Fri, Sep 23, 2022 at 08:29:46AM +0100, Julian Brown wrote:
> On Thu, 22 Sep 2022 15:17:08 +0200
> Jakub Jelinek <jakub@redhat.com> wrote:
>
> > > + bool built_sym_hash = false;
> >
> > So, I think usually we don't construct such hash_maps right away,
> > but have just pointer to the hash map initialized to NULL (then you
> > don't need to built_sym_hash next to it) and you simply new the
> > hash_map when needed the first time and delete it at the end (which
> > does nothing if it is NULL).
>
> How about this version? (Re-tested.)
I'd appreciate if Tobias could have a second look, I'm getting less and
less familiar with Fortran, from my POV LGTM.
The patch is ok if Tobias doesn't spot anything today.
Jakub
Hi Julian and Jakub, hi all,
On 23.09.22 09:29, Julian Brown wrote:
How about this version? (Re-tested.)
[...]
* * *
Some more generic (pre)remarks – not affecting the patch code,
but possibly the commit log message:
This follows OMP 5.0, 2.19.7.1 "map Clause":
which is also in "OMP 5.2, 5.8.3 map Clause [152:1-4]". It might
make sense to add this ref in addition (or instead):
"If a list item in a map clause is an associated pointer and the
pointer is not the base pointer of another list item in a map clause
on the same construct, then it is treated as if its pointer target
is implicitly mapped in the same clause. For the purposes of the map
clause, the mapped pointer target is treated as if its base pointer
is the associated pointer."
Pre-remark 1: Issue with the wording in the 5.x spec. Namely, assume
integer, pointer :: p(:)
! allocate(p(1024)); deallocate(p) ! < (un)commented lined or not
!$omp target map(p)
p => null()
!$omp end target
Here, 'p' is neither associated nor unassociated (NULL pointer),
but it is undefined. However, the 5.x spec does require an implicit mapping of
the associated pointer target – but the compiler has no idea whether
the pointer address is valid (associated) or dangling (undefined) - and
the p.data address might be either invalid to access and/or the p.dim[...]
garbage data could yield a size(p) that is huge.
Thus, the following restriction was proposed for OpenMP 6.0 (TR11):
"The association status of a list item that is a pointer must not be
undefined unless it is a structure component and it results from a
predefined default mapper."
which makes my example invalid. (Add some caveat here about TR11 not
yet being released and also TRs being not final named-version releases.)
This also affects the quote you show, which now reads (2nd line is new):
"If a list item in a map clause is an associated pointer
that is not attach-ineligible
and the pointer is not the base pointer [...]".
with 'attach-ineligible' being defined in the previous bullet list
(i.e. in the preceding paragraph (>= 5.1) about derived-type components;
first bullet point there: pointer implies attach-ineligible).
I think the 5.x wording has issues, but on the other hand, the wording
above is not in an OpenMP release and not even in a TR. As the spirit
has not changed, it probably makes sense to keep the 5.0 (5.x) wording.
Cf. https://github.com/OpenMP/spec/pull/3319/files and
https://github.com/OpenMP/spec/issues/3290
(And apologies to non-OpenMP members as those are non-publicly accessible.)
* * *
However, we can also write this:
map(to: tvar%arrptr) map(tofrom: tvar%arrptr(3:8))
and then instead we should follow:
"If the structure sibling list item is a pointer then it is treated
as if its association status is undefined, unless it appears as
the base pointer of another list item in a map clause on the same
construct."
This wording disappeared in 5.1 due to some cleanup (cf. Issue 2152, which has multiple changes; this one is Pull Req. 2379). I think the matching current / OpenMP 5.2 wording (5.8.3 map Clause [152:5-8, 11-13 (,14-16)]) is "For map clauses on map-entering constructs, if any list item has a base pointer for which a corresponding pointer exists in the data environment upon entry to the region and either a new list item or the corresponding pointer is created in the device data environment on entry to the region, then: (Fortran) 1. The corresponding pointer variable is associated with a pointer target that has the same rank and bounds as the pointer target of the original pointer, such that the corresponding list item can be accessed through the pointer in a target region. ..." I think here 'a new list item ... is created ... on entry' applies. However, this should not affect what you wrote later on. Still, I wonder whether 5.2 instead of (I think it makes sense to also read 5.2 when implementing this for bug-fix changes (and trivial feature changes), but not for larger-effort changes, which can be implemented later on.)
But, that's not implemented quite right at the moment [...]
The solution is to detect when we're mapping a smaller part of the array
(or a subcomponent) on the same directive, and only map the descriptor
in that case. So we get mappings like this instead:
map(to: tvar%arrptr) -->
GOMP_MAP_ALLOC tvar%arrptr (the descriptor)
map(tofrom: tvar%arrptr(3:8) -->
GOMP_MAP_TOFROM tvar%arrptr%data(3) (size 8-3+1, etc.)
GOMP_MAP_ALWAYS_POINTER tvar%arrptr%data (bias 3, etc.)
(I concur.)
* * *
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
...
@@ -2470,22 +2471,18 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
...
+ if (ptr_kind != GOMP_MAP_ALWAYS_POINTER)
{
...
+ /* For OpenMP, the descriptor must be mapped with its own explicit
+ map clause (e.g. both "map(foo%arr)" and "map(foo%arr(:))" must
+ be present in the clause list if "foo%arr" is a pointer to an
+ array). So, we don't create a GOMP_MAP_TO_PSET node here. */
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
I found the last sentence of the comment and the set_map_kind confusing:
The comment says no MAP_TO_PSET and the SET_MAP_KIND use it.
I wonder whether that should be something like 'if (openacc)' instead,
which kind of matches the first way gfc_trans_omp_array_section is called:
else if (lastref->type == REF_ARRAY)
{
/* An array element or section. */
bool element = lastref->u.ar.type == AR_ELEMENT;
gomp_map_kind kind = (openacc ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
gfc_trans_omp_array_section (block, n, inner, element,
kind, node, node2, node3,
node4);
However, there is also a second call to it:
/* An array element or array section which is not part of a
derived type, etc. */
...
gomp_map_kind k = GOMP_MAP_POINTER;
...
gfc_trans_omp_array_section (block, n, decl, element, k,
node, node2, node3, node4);
And without following all 'if' conditions through, I don't see why that should
be handled differently.
+ /* We're only interested in cases where we have an expression, e.g. a
+ component access. */
+ if (n->expr && n->expr->symtree)
+ use_sym = n->expr->symtree->n.sym;
The second part looks unsafe in light of 'lvalues'. The next OpenMP version
will likely permit:
"A locator list item is a variable list item, a function reference with data
pointer result, or a reserved locator."
Thus, permitting 'map( f(cnt=4))' for a function that returns a data pointer like
interface
function f(cnt) result(res)
integer :: cnt
real, pointer :: res(:,:)
end
end interface
(In Fortran, referencing 'f' counts as variable. However, contrary to C/C++,
'f()%comp or 'f()(1:4)', i.e. component and array reverences, are not permitted.)
Thus, it seems to make more sense to have n->expr->expr_type == EXPR_VARIABLE
as the symtree is also set for EXPR_FUNCTION and EXPR_COMPCALL.
The later is something like dt%f(5) where 'dt' is a variable of derived type
and 'f' is a variable bound to the derived type. – I think it might also be
set for PARAMETER, but I am not sure until which point.
+ if (!n2->expr || !n2->expr->symtree)
+ continue;
Likewise.
+ /* If the last reference is a pointer to a derived
+ type ("foo%dt_ptr"), check if any subcomponents
+ of the same derived type member are being mapped
+ elsewhere in the clause list ("foo%dt_ptr%x",
+ etc.). If we have such subcomponent mappings,
+ we only create an ALLOC node for the pointer
+ itself, and inhibit mapping the whole derived
+ type. */
Does the current code handle also the following?
i = 1; j = 2
map (foo(i)%dt_ptr(1:3), foo(j)%dt_ptr)
Note: I have not thought about validity nor checked your code, but it does
not seem to be completely odd code to write.
A similar mean way to write code would be:
integer, target :: A(5)
integer, pointer :: p(:), p2(:)
type(t) :: var
allocate(p2(1:20))
p => A
var%p2 => p2
!$omp target map(A(3:4), p2(4:8), p, var%p2)
....
!$omp end target
which has a similar issue – it is not clear from the syntax whether
p's or var%p2's pointer target has been mapped or not.
I don't currently see whether that's handled or not - but I fear it is not.
All this seems to points to first handle all non-pointer variables,
then all with tailing array refs and only then the rest - and implicit mapping
last (followed by use_device_{ptr,addr}).
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
On Fri, 23 Sep 2022 14:10:51 +0200
Tobias Burnus <tobias@codesourcery.com> wrote:
> Hi Julian and Jakub, hi all,
>
> On 23.09.22 09:29, Julian Brown wrote:
> > How about this version? (Re-tested.)
>
> Some more generic (pre)remarks – not affecting the patch code,
> but possibly the commit log message:
>
> > This follows OMP 5.0, 2.19.7.1 "map Clause":
>
> which is also in "OMP 5.2, 5.8.3 map Clause [152:1-4]". It might
> make sense to add this ref in addition (or instead):
>
> > "If a list item in a map clause is an associated pointer and the
> > pointer is not the base pointer of another list item in a map
> > clause on the same construct, then it is treated as if its pointer
> > target is implicitly mapped in the same clause. For the purposes of
> > the map clause, the mapped pointer target is treated as if its base
> > pointer is the associated pointer."
I've changed the wording in the commit log text...
> Thus, the following restriction was proposed for OpenMP 6.0 (TR11):
>
> "The association status of a list item that is a pointer must not be
> undefined unless it is a structure component and it results from a
> predefined default mapper."
>
> which makes my example invalid. (Add some caveat here about TR11 not
> yet being released and also TRs being not final named-version
> releases.)
(But not this bit, for now.)
> > and then instead we should follow:
> >
> > "If the structure sibling list item is a pointer then it is
> > treated as if its association status is undefined, unless it
> > appears as the base pointer of another list item in a map clause on
> > the same construct."
>
>
> This wording disappeared in 5.1 due to some cleanup (cf. Issue 2152,
> which has multiple changes; this one is Pull Req. 2379).
>
> I think the matching current / OpenMP 5.2 wording (5.8.3 map Clause
> [152:5-8, 11-13 (,14-16)]) is
>
> "For map clauses on map-entering constructs, if any list item has a
> base pointer for which a corresponding pointer exists in the data
> environment upon entry to the region and either a new list item or
> the corresponding pointer is created in the device data environment
> on entry to the region, then: (Fortran)
> 1. The corresponding pointer variable is associated with a pointer
> target that has the same rank and bounds as the pointer target of the
> original pointer, such that the corresponding list item can be
> accessed through the pointer in a target region. ..."
>
> I think here 'a new list item ... is created ... on entry' applies.
> However, this should not affect what you wrote later on.
I changed the text here too.
> > But, that's not implemented quite right at the moment [...]
> > The solution is to detect when we're mapping a smaller part of the
> > array (or a subcomponent) on the same directive, and only map the
> > descriptor in that case. So we get mappings like this instead:
> >
> > map(to: tvar%arrptr) -->
> > GOMP_MAP_ALLOC tvar%arrptr (the descriptor)
> >
> > map(tofrom: tvar%arrptr(3:8) -->
> > GOMP_MAP_TOFROM tvar%arrptr%data(3) (size 8-3+1, etc.)
> > GOMP_MAP_ALWAYS_POINTER tvar%arrptr%data (bias 3, etc.)
>
> (I concur.)
Thank you!
> > --- a/gcc/fortran/trans-openmp.cc
> > +++ b/gcc/fortran/trans-openmp.cc
> > ...
> > @@ -2470,22 +2471,18 @@ gfc_trans_omp_array_section (stmtblock_t
> > *block, gfc_omp_namelist *n, }
> > if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
> > {
> > ...
> > + if (ptr_kind != GOMP_MAP_ALWAYS_POINTER)
> > {
> > ...
> > + /* For OpenMP, the descriptor must be mapped with its
> > own explicit
> > + map clause (e.g. both "map(foo%arr)" and
> > "map(foo%arr(:))" must
> > + be present in the clause list if "foo%arr" is a
> > pointer to an
> > + array). So, we don't create a GOMP_MAP_TO_PSET node
> > here. */
> > + node2 = build_omp_clause (input_location,
> > OMP_CLAUSE_MAP);
> > + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
>
> I found the last sentence of the comment and the set_map_kind
> confusing: The comment says no MAP_TO_PSET and the SET_MAP_KIND use
> it.
>
> I wonder whether that should be something like 'if (openacc)' instead,
> which kind of matches the first way gfc_trans_omp_array_section is
> called:
I agree it was confusing -- I've tweaked the wording of the comment.
The condition changes in the "address tokenization" follow-up patch
anyway.
> inner, element, kind, node, node2, node3,
> node4);
> However, there is also a second call to it:
>
> /* An array element or array section which is not
> part of a derived type, etc. */
> ...
> gomp_map_kind k = GOMP_MAP_POINTER;
> ...
> gfc_trans_omp_array_section (block, n, decl,
> element, k, node, node2, node3, node4);
> And without following all 'if' conditions through, I don't see why
> that should be handled differently.
GOMP_MAP_POINTER is used for non-component accesses -- those get the
GOMP_MAP_TO_PSET mapping in gfc_trans_omp_array_section. (Some of
these conditions are based on perhaps-not-quite obvious implications,
but again, they change in the follow-up patch mentioned above anyway.)
> > + /* We're only interested in cases where we have an expression,
> > e.g. a
> > + component access. */
> > + if (n->expr && n->expr->symtree)
> > + use_sym = n->expr->symtree->n.sym;
>
> The second part looks unsafe in light of 'lvalues'. The next OpenMP
> version will likely permit:
> "A locator list item is a variable list item, a function reference
> with data pointer result, or a reserved locator."
>
> Thus, permitting 'map( f(cnt=4))' for a function that returns a data
> pointer like interface
> function f(cnt) result(res)
> integer :: cnt
> real, pointer :: res(:,:)
> end
> end interface
>
> (In Fortran, referencing 'f' counts as variable. However, contrary to
> C/C++, 'f()%comp or 'f()(1:4)', i.e. component and array reverences,
> are not permitted.)
>
> Thus, it seems to make more sense to have n->expr->expr_type ==
> EXPR_VARIABLE as the symtree is also set for EXPR_FUNCTION and
> EXPR_COMPCALL. The later is something like dt%f(5) where 'dt' is a
> variable of derived type and 'f' is a variable bound to the derived
> type. – I think it might also be set for PARAMETER, but I am not sure
> until which point.
I added n->expr->expr_type == EXPR_VARIABLE to the condition -- I think
that should suffice for now?
> > + if (!n2->expr || !n2->expr->symtree)
> > + continue;
> Likewise.
> > + /* If the last reference is a pointer to
> > a derived
> > + type ("foo%dt_ptr"), check if any
> > subcomponents
> > + of the same derived type member are
> > being mapped
> > + elsewhere in the clause list
> > ("foo%dt_ptr%x",
> > + etc.). If we have such subcomponent
> > mappings,
> > + we only create an ALLOC node for the
> > pointer
> > + itself, and inhibit mapping the whole
> > derived
> > + type. */
>
> Does the current code handle also the following?
>
> i = 1; j = 2
> map (foo(i)%dt_ptr(1:3), foo(j)%dt_ptr)
Good catch! In that gfc_dep_resolver considers those terms to have a
dependency, and that triggers the mapping node transformation. But I
don't think OpenMP allows you to write this: IIUC if "foo" is an array,
you're not allowed to separately map two parts of the array because of
(OpenMP 5.2, "5.8.3 map Clause"):
"Two list items of the map clauses on the same construct must not
share original storage unless they are the same list item or unless
one is the containing structure of the other."
and,
"If a list item is an array section, it must specify contiguous
storage."
and maybe also (for different directive arrangements),
"If an array appears as a list item in a map clause, multiple parts
of the array have corresponding storage in the device data
environment prior to a task encountering the construct associated
with the map clause, and the corresponding storage for those parts
was created by maps from more than one earlier construct, the
behavior is unspecified."
One thing that *does* work is a similar test to yours but with "i" and
"j" pointing to the *same* location. That needs the "address
tokenization" patch to be applied too, though. (I added a new test to
this version of the patch.)
(If multiple variable indices to the same struct array component *were*
allowed, that'd cause serious problems for the struct sibling list
building code -- it'd no longer be possible to sort members statically.
And if different "names" for the same blocks of host memory were
allowed, e.g. allowing a pointer and a pointed-to block to be mapped
using different variables, that'd imply a requirement to compare each
clause against each other clause at runtime -- a quadratic number of
tests, probably.)
> Note: I have not thought about validity nor checked your code, but it
> does not seem to be completely odd code to write.
>
> A similar mean way to write code would be:
>
> integer, target :: A(5)
> integer, pointer :: p(:), p2(:)
> type(t) :: var
>
> allocate(p2(1:20))
> p => A
> var%p2 => p2
> !$omp target map(A(3:4), p2(4:8), p, var%p2)
> ....
> !$omp end target
>
> which has a similar issue – it is not clear from the syntax whether
> p's or var%p2's pointer target has been mapped or not.
Again, I don't think you're allowed to write that: that's "different
list items" sharing the same "original storage", IIUC. (It'd be nice to
diagnose it at compile time, but that's probably not that easy...)
> I don't currently see whether that's handled or not - but I fear it
> is not. All this seems to points to first handle all non-pointer
> variables, then all with tailing array refs and only then the rest -
> and implicit mapping last (followed by use_device_{ptr,addr}).
I think I'd need that plan explained a bit more verbosely! But anyway,
hopefully it's not necessary.
Attached patch re-tested with offloading to NVPTX. OK?
Thanks,
Julian
Hi Julian,
On 30.09.22 15:30, Julian Brown wrote:
i = 1; j = 2
map (foo(i)%dt_ptr(1:3), foo(j)%dt_ptr)
Good catch! In that gfc_dep_resolver considers those terms to have a
dependency, and that triggers the mapping node transformation. But I
don't think OpenMP allows you to write this: IIUC if "foo" is an array,
you're not allowed to separately map two parts of the array because of
(OpenMP 5.2, "5.8.3 map Clause"):
"Two list items of the map clauses on the same construct must not
share original storage unless they are the same list item or unless
one is the containing structure of the other."
I was thinking of something like:
type t
integer, pointer :: p(:)
end t
type(t2) :: var(2)
allocate (var(1)%p, source=[1,2,3,5])
allocate (var(2)%p, source=[2,3,5])
!$omp target map ( <somehow> )
var(1)%p(1) = 5
var(2)%p(2) = 7
!$omp end target
Similarly for C:
struct st {
int *p;
};
struct st s[2];
s[0].p = (int*)__malloc(sizeof(int)*5);
s[1].p = (int*)__malloc(sizeof(int)*3);
#pragma omp target map( <somehow> )
{
s[0].p[0] = 5;
s[1].p[1] = 7;
}
And now I somehow need to map "p" in both the C and the Fortran case.
And I believe that should be possible...
and my
i = 1; j = 2
map (var(i)%p(1:3), var(j)%p)
or
map (s[0].p[:3], s[1].p[:7])
does not look to far fetched to get this result ...
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
On 30.09.22 15:30, Julian Brown wrote:
On Fri, 23 Sep 2022 14:10:51 +0200 Tobias Burnus <tobias@codesourcery.com><mailto:tobias@codesourcery.com> wrote:
...
I added n->expr->expr_type == EXPR_VARIABLE to the condition -- I think
that should suffice for now?
Yes.
A similar mean way to write code would be:
integer, target :: A(5)
integer, pointer :: p(:), p2(:)
type(t) :: var
allocate(p2(1:20))
p => A
var%p2 => p2
!$omp target map(A(3:4), p2(4:8), p, var%p2)
....
!$omp end target
which has a similar issue – it is not clear from the syntax whether
p's or var%p2's pointer target has been mapped or not.
Again, I don't think you're allowed to write that: that's "different
list items" sharing the same "original storage", IIUC. (It'd be nice to
diagnose it at compile time, but that's probably not that easy...)
Hmm, isn't that implied to be valid per:
"[Fortran] If a list item in a map clause is an associated pointer and the pointer
is not the base pointer of another list item in a map clause on the same
construct, then it is treated as if its pointer target is implicitly
mapped in the same clause."
(OpenMP 5.2, 5.8.3 map Clause [152:1-4]; seems to be identical to
OpenMP 5.1, 2.21.7.1 map Clause [349:9-12])
Admittedly, in 5.0 I only see the wording (OpenMP 5.0, [316:22-25]),
which is also not handled - but different:
"[Fortran] Each pointer component that is a list item that results from
a mapped derived type variable is treated as if its association status
is undefined, unless the pointer component appears as another list
item or as the base pointer of another list item in a map clause on
the same construct."
I think the that's for the following case (which is also covered by the
more general 5.1/5.2 wording):
type t
integer, pointer :: p(:)
integer, pointer :: p2(:)
end type t
type(t) :: var
integer, target :: tgt(5), tgt2(1000)
var%p => tgt
var%p => tgt2
!$omp target map(tgt, tgt2(4:6), var)
var%p(1) = 5
var%p2(5) = 7
!$omp end target
and I think GCC does not handled this, but I might be wrong.
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
@@ -40,6 +40,7 @@ along with GCC; see the file COPYING3. If not see
#include "omp-general.h"
#include "omp-low.h"
#include "memmodel.h" /* For MEMMODEL_ enums. */
+#include "dependency.h"
#undef GCC_DIAG_STYLE
#define GCC_DIAG_STYLE __gcc_tdiag__
@@ -2470,22 +2471,18 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
- tree desc_node;
tree type = TREE_TYPE (decl);
ptr2 = gfc_conv_descriptor_data_get (decl);
- desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
- OMP_CLAUSE_DECL (desc_node) = decl;
- OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
{
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
- node2 = node;
- node = desc_node; /* Needs to come first. */
- }
- else
- {
- OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
- node2 = desc_node;
+ /* For OpenMP, the descriptor must be mapped with its own explicit
+ map clause (e.g. both "map(foo%arr)" and "map(foo%arr(:))" must
+ be present in the clause list if "foo%arr" is a pointer to an
+ array). So, we don't create a GOMP_MAP_TO_PSET node here. */
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
@@ -3440,6 +3437,50 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
{
if (pointer || (openacc && allocatable))
{
+ gfc_omp_namelist *n2
+ = openacc ? NULL : clauses->lists[OMP_LIST_MAP];
+
+ /* If the last reference is a pointer to a derived
+ type ("foo%dt_ptr"), check if any subcomponents
+ of the same derived type member are being mapped
+ elsewhere in the clause list ("foo%dt_ptr%x",
+ etc.). If we have such subcomponent mappings,
+ we only create an ALLOC node for the pointer
+ itself, and inhibit mapping the whole derived
+ type. */
+
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if (n == n2 || !n2->expr)
+ continue;
+
+ int dep
+ = gfc_dep_resolver (n->expr->ref, n2->expr->ref,
+ NULL, true);
+ if (dep == 0)
+ continue;
+
+ gfc_ref *ref1 = n->expr->ref;
+ gfc_ref *ref2 = n2->expr->ref;
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next)
+ {
+ inner = build_fold_addr_expr (inner);
+ OMP_CLAUSE_SET_MAP_KIND (node,
+ GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (node) = inner;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+ goto finalize_map_clause;
+ }
+ }
+
tree data, size;
if (lastref->u.c.component->ts.type == BT_CLASS)
@@ -3541,8 +3582,47 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
node2 = desc_node;
else
{
+ gfc_omp_namelist *n2
+ = clauses->lists[OMP_LIST_MAP];
node2 = node;
node = desc_node; /* Put first. */
+ for (; n2 != NULL; n2 = n2->next)
+ {
+ if (n == n2 || !n2->expr)
+ continue;
+
+ int dep
+ = gfc_dep_resolver (n->expr->ref,
+ n2->expr->ref,
+ NULL, true);
+ if (dep == 0)
+ continue;
+
+ gfc_ref *ref1 = n->expr->ref;
+ gfc_ref *ref2 = n2->expr->ref;
+
+ /* We know ref1 and ref2 overlap. We're
+ interested in whether ref2 describes a
+ smaller part of the array than ref1, which
+ we already know refers to the full
+ array. */
+
+ while (ref1->next && ref2->next)
+ {
+ ref1 = ref1->next;
+ ref2 = ref2->next;
+ }
+
+ if (ref2->next
+ || (ref2->type == REF_ARRAY
+ && (ref2->u.ar.type == AR_ELEMENT
+ || (ref2->u.ar.type
+ == AR_SECTION))))
+ {
+ node2 = NULL_TREE;
+ goto finalize_map_clause;
+ }
+ }
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
new file mode 100644
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+program myprog
+type u
+ integer, dimension (:), pointer :: tarr
+end type u
+
+type(u) :: myu
+integer, dimension (12), target :: myarray
+
+myu%tarr => myarray
+
+myu%tarr = 0
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(:))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1:2))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(to:myu%tarr) map(tofrom:myu%tarr(1))
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+!$omp target map(tofrom:myu%tarr)
+myu%tarr(1) = myu%tarr(1) + 1
+!$omp end target
+
+if (myu%tarr(1).ne.4) stop 1
+
+end program myprog
new file mode 100644
@@ -0,0 +1,32 @@
+! { dg-do run }
+
+module mymod
+type F
+integer :: a, b, c
+integer, dimension(10) :: d
+end type F
+
+type G
+integer :: x, y
+type(F), pointer :: myf
+integer :: z
+end type G
+end module mymod
+
+program myprog
+use mymod
+
+type(F), target :: ftmp
+type(G) :: gvar
+
+gvar%myf => ftmp
+
+gvar%myf%d = 0
+
+!$omp target map(to:gvar%myf) map(tofrom: gvar%myf%b, gvar%myf%d)
+gvar%myf%d(1) = gvar%myf%d(1) + 1
+!$omp end target
+
+if (gvar%myf%d(1).ne.1) stop 1
+
+end program myprog
@@ -229,7 +229,8 @@ contains
! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3)) &
! !$omp& map(tofrom: var%str4(2:2), var%uni2(2:3), var%uni4(2:2))
- !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3), var%uni2(2:3))
+ !$omp target map(to: var%f) map(tofrom: var%d(4:7), var%f(2:3), &
+ !$omp& var%str2(2:3), var%uni2(2:3))
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
@@ -274,7 +275,7 @@ contains
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
!$omp end target
- !$omp target map(tofrom: var%f(2:3))
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f(2:3) /= [33, 44])) stop 11
@@ -314,7 +315,8 @@ contains
! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), &
! !$omp var%str4(2), var%uni2(3), var%uni4(2))
- !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), var%uni2(3))
+ !$omp target map(to: var%f) map(tofrom: var%d(5), var%f(3), &
+ !$omp& var%str2(3), var%uni2(3))
if (var%d(5) /= -3*5) stop 4
if (var%str2(3) /= "ABCDE") stop 6
if (var%uni2(3) /= 4_"ABCDE") stop 7
@@ -362,7 +364,7 @@ contains
if (any (var%uni2(2:3) /= [4_"67890", 4_"ABCDE"])) stop 7
!$omp end target
- !$omp target map(tofrom: var%f(2:3))
+ !$omp target map(to: var%f) map(tofrom: var%f(2:3))
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f(2:3) /= [33, 44])) stop 11