There was a case of the implicit non-static pointer member mapping
not working properly with templates.
What happened was that the code in finish_omp_target() created the
map clauses (which normally runs after finish_omp_clauses), but being
a template class it was put through all the tsubst_* stuff and at the
end thrown into finish_omp_clauses a 2nd time. And because finish_omp_clauses
didn't handle some of the implicitly created map clauses, things didn't
work...
This patch slightly fixes many handled cases in these parts, plus some
adjustments in gimplify.c.
Tested without regressions, and pushed to devel/omp/gcc-10.
Chung-Lin
From 4e714eaad985f68533f267b8df2026e5c14d084a Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <clt...@codesourcery.com>
Date: Thu, 11 Mar 2021 00:31:08 -0800
Subject: [PATCH] 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 <clt...@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.
---
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(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/target-this-5.C
create mode 100644 libgomp/testsuite/libgomp.c++/target-this-5.C
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 55a5983..5b62fa3 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6407,6 +6407,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;
@@ -7516,6 +7517,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)
@@ -7551,6 +7560,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;
@@ -7614,6 +7629,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
@@ -7698,7 +7714,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));
@@ -8702,9 +8719,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);
@@ -8724,7 +8744,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)
@@ -8744,9 +8765,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);
@@ -8768,9 +8791,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);
@@ -8852,8 +8878,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 ff44d26..91aa15d 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9111,6 +9111,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))
@@ -9522,6 +9523,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 a450b37..08568f9 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 af23cbb..3b2d581 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 0000000..a9ac74b
--- /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 0000000..e71c566
--- /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;
+}
--
2.8.1
