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 --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index d37e231641c3fafd14f9224a8ee00101f9d0d75f..46eb5807eb988807fe055f4c655b3c84554535da 100644 (file)
--- 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 03fbe461bd5cdb6361cdc5b1ab1d4cf9e6f246f4..0dd8e7d71b4df701cdae645995c15d9a2c017bb5 100644 (file)
--- 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 a450b3723e5c5b0fe739a677abe23f35159074dd..08568f9284c1cb481fd25648512198672fb9ee0d 100644 (file)
--- 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 af23cbb90231ea05dd749704806f1280db9cf60b..3b2d5811350bfe1efb36e1f3f75e9b10c13b2565 100644 (file)
--- 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 (file)
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 (file)
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;
+}