[NFC][Clang][OpenMP] Add helper functions/utils for finding/comparing attach base-ptrs. by abhinavgaba · Pull Request #155625 · llvm/llvm-project (original) (raw)
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.Learn more about bidirectional Unicode characters
[ Show hidden characters]({{ revealButtonHref }})
…base-ptrs.
These have been pulled out of the codegen PR llvm#153683, to reduce the size of that PR.
…undant variable, add some asserts.
abhinavgaba deleted the introduce-attach-helper-utils-in-clang branch
abhinavgaba added a commit to abhinavgaba/llvm-project that referenced this pull request
abhinavgaba added a commit that referenced this pull request
mahesh-attarde pushed a commit to mahesh-attarde/llvm-project that referenced this pull request
… attach base-ptrs. (llvm#155625)
These have been pulled out of the codegen PR llvm#153683, to reduce the size of that PR.
mahesh-attarde pushed a commit to mahesh-attarde/llvm-project that referenced this pull request
abhinavgaba added a commit that referenced this pull request
…rs. (#153683)
This adds support for using ATTACH map-type for proper
pointer-attachment when mapping list-items that have base-pointers.
For example, for the following:
int *p;
#pragma omp target enter data map(p[1:10])The following maps are now emitted by clang:
(A)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
&p, &p[1], sizeof(p), ATTACHPreviously, the two possible maps emitted by clang were:
(B)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
(C)
&p, &p[1], 10 * sizeof(p[1]), TO | FROM | PTR_AND_OBJ(B) does not perform any pointer attachment, while (C) also maps the pointer p, both of which are incorrect.
With this change, we are using ATTACH-style maps, like (A), for cases
where the expression has a base-pointer. For example:
int *p, **pp;
S *ps, **pps;
... map(p[0])
... map(p[10:20])
... map(*p)
... map(([20])p)
... map(ps->a)
... map(pps->p->a)
... map(pp[0][0])
... map(*(pp + 10)[0])
Grouping of maps based on attach base-pointers
We also group mapping of clauses with the same base decl in the order of the increasing complexity of their base-pointers, e.g. for something like:
S **spp;
map(spp[0][0], spp[0][0].a), // attach-ptr: spp[0]
map(spp[0]), // attach-ptr: spp
map(spp), // attach-ptr: N/AWe first map spp, then spp[0] then spp[0][0] and spp[0][0].a.
This allows us to also group "struct" allocation based on their attach
pointers. This resolves the issues of us always mapping everything from
the beginning of the symbol spp. Each group is mapped independently,
and at the same level, like spp[0][0] and its member spp[0][0].a, we
still get map them together as part of the same contiguous struct
spp[0][0]. This resolves issue #141042.
use_device_ptr/addr fixes
The handling of use_device_ptr/addr was updated to use the attach-ptr
information, and works for many cases that were failing before. It has
to be done as part of this series because otherwise, the switch from
ptr_to_obj to attach-style mapping would have caused regressions in
existing use_device_ptr/addr tests.
Handling of attach-pointers that are members of implicitly mapped
structs:
- When a struct member-pointer, like
pbelow, is a base-pointer in amapclause on a target construct (likemap(p[0:1]), and the base of that struct is either thethispointer (implicitly or explicitly), or a struct that is implicitly mapped on that construct, we add an implicitmap(p)so that we don't implicitly map the full struct.
struct S { int *p;
void f1() {
#pragma omp target map(p[0:1]) // Implicitly map this->p, to ensure
// that the implicit map of `this[:]` does
// not map the full struct
printf("%p %p\n", &p, p);
}Scope for improvement:
- We may be able to compute attach-ptr expr while collecting component-lists in Sema.
- But we cache the computation results already, and
findAttachPtrExpris fairly simple, and fast. - There may be a better way to implement semantic expr comparison.
Needs future work:
- Attach-style maps not yet emitted for declare mappers.
- Mapping of class member references: We are still using PTR_AND_OBJ
maps for them. We will likely need to change that to handle
ref_ptr/ref_ptee, andattachmap-type-modifier on them. - Implicit capturing of "this" needs to map the full
this[0:1]unless there is an explicit map on one of the members, or a map with a member as its base-pointer. - Implicit map added for capturing a class member pointer needs to also add a zero-length-array-section map.
use_device_addron array-sections-on-pointers need further improvements (documented using FIXMEs)
Why a large PR
While it's unfortunate that this PR has gotten large and difficult to review, the issue is that all the functional changes have to be made together, to prevent regressions from partially implemented changes.
For example, the changes to capturing were previously done separately (#145454), but they would still cause stability issues in absence of full attach-mapping. And attach-mapping needs those changes to be able to launch kernels.
We extracted the utilities and functions, like those for finding attach-ptrs, or comparing exprs, out as a separate NFC PR that doesn't call those functions, just adds them (#155625). Maybe the change that adds a new error message for use_device_addr on array-sections with non-var base-pointers could have been extracted out too (but that would have had to be a follow-up change in that case, and we would get comp-fails with this PR when the erroneous case was not caught/diagnosed).
Co-authored-by: Alex Duran alejandro.duran@intel.com
ronlieb pushed a commit to ROCm/llvm-project that referenced this pull request
…rs. (llvm#153683)
This adds support for using ATTACH map-type for proper
pointer-attachment when mapping list-items that have base-pointers.
For example, for the following:
int *p;
#pragma omp target enter data map(p[1:10])The following maps are now emitted by clang:
(A)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
&p, &p[1], sizeof(p), ATTACHPreviously, the two possible maps emitted by clang were:
(B)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
(C)
&p, &p[1], 10 * sizeof(p[1]), TO | FROM | PTR_AND_OBJ(B) does not perform any pointer attachment, while (C) also maps the pointer p, both of which are incorrect.
With this change, we are using ATTACH-style maps, like (A), for cases
where the expression has a base-pointer. For example:
int *p, **pp;
S *ps, **pps;
... map(p[0])
... map(p[10:20])
... map(*p)
... map(([20])p)
... map(ps->a)
... map(pps->p->a)
... map(pp[0][0])
... map(*(pp + 10)[0])
We also group mapping of clauses with the same base decl in the order of the increasing complexity of their base-pointers, e.g. for something like:
S **spp;
map(spp[0][0], spp[0][0].a), // attach-ptr: spp[0]
map(spp[0]), // attach-ptr: spp
map(spp), // attach-ptr: N/AWe first map spp, then spp[0] then spp[0][0] and spp[0][0].a.
This allows us to also group "struct" allocation based on their attach
pointers. This resolves the issues of us always mapping everything from
the beginning of the symbol spp. Each group is mapped independently,
and at the same level, like spp[0][0] and its member spp[0][0].a, we
still get map them together as part of the same contiguous struct
spp[0][0]. This resolves issue llvm#141042.
The handling of use_device_ptr/addr was updated to use the attach-ptr
information, and works for many cases that were failing before. It has
to be done as part of this series because otherwise, the switch from
ptr_to_obj to attach-style mapping would have caused regressions in
existing use_device_ptr/addr tests.
structs:
- When a struct member-pointer, like
pbelow, is a base-pointer in amapclause on a target construct (likemap(p[0:1]), and the base of that struct is either thethispointer (implicitly or explicitly), or a struct that is implicitly mapped on that construct, we add an implicitmap(p)so that we don't implicitly map the full struct.
struct S { int *p;
void f1() {
#pragma omp target map(p[0:1]) // Implicitly map this->p, to ensure
// that the implicit map of `this[:]` does
// not map the full struct
printf("%p %p\n", &p, p);
}We may be able to compute attach-ptr expr while collecting component-lists in Sema.
But we cache the computation results already, and
findAttachPtrExpris fairly simple, and fast.There may be a better way to implement semantic expr comparison.
Attach-style maps not yet emitted for declare mappers.
Mapping of class member references: We are still using PTR_AND_OBJ maps for them. We will likely need to change that to handle
ref_ptr/ref_ptee, andattachmap-type-modifier on them.Implicit capturing of "this" needs to map the full
this[0:1]unless there is an explicit map on one of the members, or a map with a member as its base-pointer.Implicit map added for capturing a class member pointer needs to also add a zero-length-array-section map.
use_device_addron array-sections-on-pointers need further improvements (documented using FIXMEs)
While it's unfortunate that this PR has gotten large and difficult to review, the issue is that all the functional changes have to be made together, to prevent regressions from partially implemented changes.
For example, the changes to capturing were previously done separately (llvm#145454), but they would still cause stability issues in absence of full attach-mapping. And attach-mapping needs those changes to be able to launch kernels.
We extracted the utilities and functions, like those for finding attach-ptrs, or comparing exprs, out as a separate NFC PR that doesn't call those functions, just adds them (llvm#155625). Maybe the change that adds a new error message for use_device_addr on array-sections with non-var base-pointers could have been extracted out too (but that would have had to be a follow-up change in that case, and we would get comp-fails with this PR when the erroneous case was not caught/diagnosed).
Co-authored-by: Alex Duran alejandro.duran@intel.com