Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
OpenMP: Handle same-directive mapped vars with pointer predefined fir…
…stprivate [PR110639]
This patch fixes the issue:
int a[100];
p = &a[0];
#pragma omp target map(a)
p[0] = p[99] = 3;
where 'p' is predetermined firstprivate, i.e. it is firstprivatized
but its address gets updated to the device address of 'a' as there is
associated storage for the value of 'p', i.e. its pointee.
[This is a C/C++-only feature that cannot be replicated by using a single clause.
('target data map(a) use_device_ptr(p)' + 'target is_device_ptr(p)' would do
so in two steps. - or 'p2 = omp_get_mapped_ptr(p, devnum)' + 'target is_device_ptr(p2)'.)]
Before this only worked when that storage was mapped before and not on the same
directive.
The gimplify_scan_omp_clauses change was done when I saw some runtime fails; I think
those were due to a bug in libgomp (now fixed) and not due to having two pointer
privatisations in a now different order. Still, they at least prevent mapping
'this' multiple times when 'this' is not 'this' but __closure->this which is at least
a missed optimization. And also for libgomp.c++/pr108286.C which has a normal
'this' and map(tofrom:*this [len: 16]).
Build and tested without offloading and with nvptx offloading.
Comments, remarks, suggestions?
* * *
(I wonder whether our current approach of removing explicit MAP if its
DECL is unsued is the right one if there is any GOVD_MAP_0LEN_ARRAY around
- or even any OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION.)
(See new libgomp.c-c++-common/target-implicit-map-6.c; BTW, I tried:
before '(void) a;' but that only worked with C and not with C++.)
* * *
The other issue in the PR (still to be done) is for code like:
int a[100];
p = &a[0];
#pragma omp target map(a[20:20]) // Map only a[20] to a[40], but p points to &a[0]
p[20] = p[30] = 3;
where 'p' points to the base address of 'a' but p[0] == a[0] it not actually
mapped. As we currently do not keep track of base pointer, this won't work.
I have not (yet) explored how to best implement this.
* * *
OpenMP Spec:
The first feature is not new, but I have not checked the wording in 4.5 or 5.0;
it might be that older versions only required it to work for storage mapped before
the current taget directive. But at least TR12 is very explicit in permitting it
and the (nonpublic) issue which lead to the 5.1 change also uses this. (See PR.)
(The second feature is definitely new in OpenMP 5.1.)
TR12 states in "14.8 target Construct" [379:8-10]:
"[C/C++] If a list item in a map clause has a base pointer that is predetermined firstprivate
(see Section 6.1.1) and on entry to the target region the list item is mapped, the firstprivate
pointer is updated via corresponding base pointer initialization."
(For OpenMP 5.1, read its Section 2.21.7.2.)
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
OpenMP: Handle same-directive mapped vars with pointer predefined firstprivate [PR110639]
Predefined 'firstprivate' for pointer variables firstprivatizes the pointer
but if it is associated with a mapped target, its address is updated to the
corresponding target. (If not, the host value remains.)
This commit extends this handling to also update the pointer address for
storaged mapped on the same directive.
The 'gimplify_scan_omp_clauses' change avoids adding an additional
map(alloc:this) (+ptr assignment)
when there is already a
map(tofrom:*this) (+ptr assignment)
This shows up for libgomp.c++/pr108286.C and also when 'this' is
actually '__closure->this' (-> g++.dg/gomp/target-{this-{2,4},lambda-1}.C).
PR middle-end/110639
gcc/ChangeLog:
* gimplify.cc (struct gimplify_adjust_omp_clauses_data): Add
append_list.
(gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Add
GOVD_MAP_0LEN_ARRAY clauses at the end.
(gimplify_scan_omp_clauses): Mark also '*var' as found not only
'var'.
libgomp/ChangeLog:
* target.c (gomp_map_vars_internal): Handle also variables
mapped in the same directive for GOVD_MAP_0LEN_ARRAY.
* testsuite/libgomp.c++/pr108286.C: Add gimple tree-scan test.
* testsuite/libgomp.c-c++-common/target-implicit-map-6.c: New test.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-this-2.C: Remove 'this' pointer mapping alreay
mapped via __closure->this.
* g++.dg/gomp/target-this-4.C: Likewise.
* g++.dg/gomp/target-lambda-1.C: Likewise. Move 'iptr' pointer
mapping to the end in scan-tree-dump.
gcc/gimplify.cc | 45 ++++-
gcc/testsuite/g++.dg/gomp/target-lambda-1.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-2.C | 4 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 6 +-
libgomp/target.c | 11 +-
libgomp/testsuite/libgomp.c++/pr108286.C | 4 +
.../libgomp.c-c++-common/target-implicit-map-6.c | 212 +++++++++++++++++++++
7 files changed, 276 insertions(+), 10 deletions(-)- Loading branch information
1 parent
c2d62cd
commit 2e3c468