public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Fix template case of non-static member access inside member functions
@ 2021-05-13 16:20 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:20 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:56314e382995ebbf35f77b9a2542b5411b9ae755
commit 56314e382995ebbf35f77b9a2542b5411b9ae755
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Thu Mar 11 00:31:08 2021 -0800
Fix template case of non-static member access inside member functions
Prior patches for C++ non-static member access had problems under template
classes, due to re-calling of finish_omp_clauses after finish_omp_target
created the implicit maps required, but not of allowed form in finish_omp_clauses.
This patch solves this by slightly relaxing the allowed expressions in
finish_omp_clauses.
2021-03-11 Chung-Lin Tang <cltang@codesourcery.com>
gcc/cp/ChangeLog:
* semantics.c (finish_omp_clauses): Adjustments to allow '*ptr' and
'ptr->member' cases in map clausess.
(finish_omp_target): Use INDIRECT_REF instead of MEM_REF in created
clauses, add processing_template_decl handling.
gcc/ChangeLog:
* gimplify.c (gimplify_scan_omp_clauses): Under !DECL_P case of
GOMP_CLAUSE_MAP handling, add STRIP_NOPS for indir_p case, add to
struct_deref_set for map(*ptr_to_struct) cases.
gcc/testsuite/ChangeLog:
* g++.dg/gomp/target-this-3.C: Adjust scan test.
* g++.dg/gomp/target-this-4.C: Likewise.
* g++.dg/gomp/target-this-5.C: New test.
libgomp/ChangeLog:
* testsuite/libgomp.c++/target-this-5.C: New test.
Diff:
---
gcc/cp/semantics.c | 45 +++++++++++++++++++++------
gcc/gimplify.c | 19 +++++++++++
gcc/testsuite/g++.dg/gomp/target-this-3.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-4.C | 2 +-
gcc/testsuite/g++.dg/gomp/target-this-5.C | 34 ++++++++++++++++++++
libgomp/testsuite/libgomp.c++/target-this-5.C | 30 ++++++++++++++++++
6 files changed, 120 insertions(+), 12 deletions(-)
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index d37e231641c..46eb5807eb9 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6600,6 +6600,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
bool order_seen = false;
bool schedule_seen = false;
bool oacc_async = false;
+ bool indirect_ref_p = false;
bool indir_component_ref_p = false;
tree last_iterators = NULL_TREE;
bool last_iterators_remove = false;
@@ -7784,6 +7785,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
indir_component_ref_p = true;
STRIP_NOPS (t);
}
+ indirect_ref_p = false;
+ if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+ && INDIRECT_REF_P (t))
+ {
+ t = TREE_OPERAND (t, 0);
+ indirect_ref_p = true;
+ STRIP_NOPS (t);
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
@@ -7819,6 +7828,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
t = TREE_OPERAND (t, 0);
+ if (INDIRECT_REF_P (t))
+ {
+ t = TREE_OPERAND (t, 0);
+ indir_component_ref_p = true;
+ STRIP_NOPS (t);
+ }
}
if (remove)
break;
@@ -7882,6 +7897,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|| (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_POINTER))
&& !indir_component_ref_p
+ && !indirect_ref_p
&& !cxx_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -7966,7 +7982,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
{
- bitmap_set_bit (&map_head, DECL_UID (t));
+ if (!indirect_ref_p && !indir_component_ref_p)
+ bitmap_set_bit (&map_head, DECL_UID (t));
if (t != OMP_CLAUSE_DECL (c)
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
bitmap_set_bit (&map_field_head, DECL_UID (t));
@@ -9027,9 +9044,12 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
tree closure = DECL_ARGUMENTS (current_function_decl);
tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
- OMP_CLAUSE_DECL (c) = build_simple_mem_ref (closure);
+ OMP_CLAUSE_DECL (c)
+ = build_indirect_ref (loc, closure, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure)));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure))));
tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
@@ -9049,7 +9069,8 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
/* Transform *this into *__closure->this in maps. */
tree this_map = *explicit_this_deref_map;
OMP_CLAUSE_DECL (this_map)
- = build_simple_mem_ref (omp_target_this_expr);
+ = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
+
tree nc = OMP_CLAUSE_CHAIN (this_map);
gcc_assert (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (nc)
@@ -9069,9 +9090,11 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM);
OMP_CLAUSE_DECL (c3)
- = build_simple_mem_ref (omp_target_this_expr);
+ = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c3)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
@@ -9093,9 +9116,12 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
{
tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
- OMP_CLAUSE_DECL (c) = build_simple_mem_ref (omp_target_this_expr);
+ OMP_CLAUSE_DECL (c)
+ = build_indirect_ref (loc, omp_target_this_expr, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c)
- = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+ = (processing_template_decl
+ ? NULL_TREE
+ : TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr))));
tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
@@ -9177,8 +9203,7 @@ finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
OMP_CLAUSE_DECL (c)
- = build2 (MEM_REF, char_type_node, ptr_member,
- build_int_cst (build_pointer_type (char_type_node), 0));
+ = build_indirect_ref (loc, ptr_member, RO_UNARY_STAR);
OMP_CLAUSE_SIZE (c) = size_zero_node;
OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 03fbe461bd5..0dd8e7d71b4 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9178,6 +9178,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
+ STRIP_NOPS (decl);
}
if (TREE_CODE (decl) == INDIRECT_REF
&& DECL_P (TREE_OPERAND (decl, 0))
@@ -9589,6 +9590,24 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
+ /* If this was of the form map(*pointer_to_struct), then the
+ 'pointer_to_struct' DECL should be considered deref'ed. */
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ || GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (c))
+ || GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (c)))
+ && INDIRECT_REF_P (orig_decl)
+ && DECL_P (TREE_OPERAND (orig_decl, 0))
+ && TREE_CODE (TREE_TYPE (orig_decl)) == RECORD_TYPE)
+ {
+ tree ptr = TREE_OPERAND (orig_decl, 0);
+ if (!struct_deref_set || !struct_deref_set->contains (ptr))
+ {
+ if (!struct_deref_set)
+ struct_deref_set = new hash_set<tree> ();
+ struct_deref_set->add (ptr);
+ }
+ }
+
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index a450b3723e5..08568f9284c 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -102,4 +102,4 @@ int main (void)
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index af23cbb9023..3b2d5811350 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-4.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -102,6 +102,6 @@ int main (void)
return 0;
}
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\) map\(from:mapped \[len: 1\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(struct:\*__closure \[len: 1\]\) map\(alloc:__closure->__this \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+ \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_[0-9]+->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-5.C b/gcc/testsuite/g++.dg/gomp/target-this-5.C
new file mode 100644
index 00000000000..a9ac74bcf1f
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-5.C
@@ -0,0 +1,34 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+ T a, b, c, d;
+
+ T sum (void)
+ {
+ T val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ T sum_offload (void)
+ {
+ T val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S<int> s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-this-5.C b/libgomp/testsuite/libgomp.c++/target-this-5.C
new file mode 100644
index 00000000000..e71c566687d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-5.C
@@ -0,0 +1,30 @@
+extern "C" void abort ();
+
+template<typename T>
+struct S
+{
+ T a, b, c, d;
+
+ T sum (void)
+ {
+ T val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ T sum_offload (void)
+ {
+ T val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S<int> s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:20 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:20 [gcc/devel/omp/gcc-11] Fix template case of non-static member access inside member functions Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).