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.

6 files changed, 120 insertions, 12 deletions

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;
+}