Message ID | a3be658301113143e5ff5efea74e46ea6efc3e5f.1663101299.git.julian@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | OpenMP 5.0: Struct & mapping clause expansion rework | expand |
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
diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 4949f001fed..a5411df4ad7 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -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); diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray.f90 new file mode 100644 index 00000000000..85f5af3a2a6 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 new file mode 100644 index 00000000000..c7f90131cba --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subcomponents.f90 @@ -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 diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 index 58550c79d69..f128ebcffc1 100644 --- a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 @@ -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