Enable gimplify GOMP_MAP_STRUCT handling of (COMPONENT_REF (INDIRECT_REF ...)) map clauses.

This patch tries to allow map(A->ptr) to be properly handled the same way as
map(B.ptr) expressions. map(struct:*A) clauses are now produced during
gimplify.

This patch, as of time of commit, is only pushed to devel/omp/gcc-10, not yet
submitted as mainline patch to upstream.

2021-02-08  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/ChangeLog:

	* gimplify.c ("tree-hash-traits.h"): Add include.
	(gimplify_scan_omp_clauses): Change struct_map_to_clause to type
	hash_map<tree_operand, tree> *. Adjust struct map handling to handle
	cases of *A and A->B expressions.
	(gimplify_adjust_omp_clauses): Move GOMP_MAP_STRUCT removal code for
	exit data directives code to earlier position.

gcc/testsuite/ChangeLog:

	* g++.dg/gomp/target-3.C: Adjust testcase gimple scanning.
	* g++.dg/gomp/target-this-2.C: Likewise.
	* g++.dg/gomp/target-this-3.C: Likewise.
	* g++.dg/gomp/target-this-4.C: Likewise.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-23.C: New testcase.

5 files changed, 44 insertions, 17 deletions

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index b90ba5bbbf1..ba190177092 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -53,6 +53,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "langhooks.h"
 #include "tree-cfg.h"
 #include "tree-ssa.h"
+#include "tree-hash-traits.h"
 #include "omp-general.h"
 #include "omp-low.h"
 #include "gimple-low.h"
@@ -8514,7 +8515,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 {
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
-  hash_map<tree, tree> *struct_map_to_clause = NULL;
+  hash_map<tree_operand_hash, tree> *struct_map_to_clause = NULL;
   hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL, *orig_list_p = list_p;
   int handled_depend_iterators = -1;
@@ -9082,12 +9083,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  && TREE_CODE (decl) == INDIRECT_REF
 		  && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
 		  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
-		      == REFERENCE_TYPE))
+		      == REFERENCE_TYPE)
+		  && (OMP_CLAUSE_MAP_KIND (c)
+		      != GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION))
 		{
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
 	      bool indir_p = false;
+	      bool component_ref_p = false;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
 	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
@@ -9098,6 +9102,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    {
 		      decl = TREE_OPERAND (decl, 0);
+		      component_ref_p = true;
 		      if (((TREE_CODE (decl) == MEM_REF
 			    && integer_zerop (TREE_OPERAND (decl, 1)))
 			   || INDIRECT_REF_P (decl))
@@ -9117,8 +9122,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		    }
 		}
-	      else if (TREE_CODE (decl) == COMPONENT_REF)
+	      else if (TREE_CODE (decl) == COMPONENT_REF
+		       && (OMP_CLAUSE_MAP_KIND (c)
+			   != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
 		{
+		  component_ref_p = true;
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    decl = TREE_OPERAND (decl, 0);
 		  if (TREE_CODE (decl) == INDIRECT_REF
@@ -9188,7 +9196,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      if (code == OACC_UPDATE
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
-	      if (DECL_P (decl)
+	      if ((DECL_P (decl)
+		   || (component_ref_p && INDIRECT_REF_P (decl)))
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
@@ -9245,7 +9254,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  gcc_assert (base == decl);
 
 		  splay_tree_node n
-		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
+		    = (DECL_P (decl)
+		       ? splay_tree_lookup (ctx->variables,
+					    (splay_tree_key) decl)
+		       : NULL);
 		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)
 			      == GOMP_MAP_ALWAYS_POINTER);
 		  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
@@ -9271,7 +9283,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		      has_attachments = true;
 		    }
-		  if (n == NULL || (n->value & GOVD_MAP) == 0)
+		  if ((DECL_P (decl)
+		       && (n == NULL || (n->value & GOVD_MAP) == 0))
+		      || (!DECL_P (decl)
+			  && (!struct_map_to_clause
+			      || struct_map_to_clause->get (decl) == NULL)))
 		    {
 		      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 						 OMP_CLAUSE_MAP);
@@ -9290,7 +9306,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
 			   : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
 		      if (struct_map_to_clause == NULL)
-			struct_map_to_clause = new hash_map<tree, tree>;
+			struct_map_to_clause
+			  = new hash_map<tree_operand_hash, tree>;
 		      struct_map_to_clause->put (decl, l);
 		      if (ptr || attach_detach)
 			{
@@ -9324,7 +9341,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			flags |= GOVD_SEEN;
 		      if (has_attachments)
 			flags |= GOVD_MAP_HAS_ATTACHMENTS;
-		      goto do_add_decl;
+		      if (DECL_P (decl))
+			goto do_add_decl;
 		    }
 		  else if (struct_map_to_clause)
 		    {
@@ -9433,6 +9451,13 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		      else if (*sc != c)
 			{
+			  if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
+					     fb_lvalue)
+			      == GS_ERROR)
+			    {
+			      remove = true;
+			      break;
+			    }
 			  *list_p = OMP_CLAUSE_CHAIN (c);
 			  OMP_CLAUSE_CHAIN (c) = *sc;
 			  *sc = c;
@@ -10811,6 +10836,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		    }
 		}
 	    }
+	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+	      && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA))
+	    {
+	      remove = true;
+	      break;
+	    }
 	  if (!DECL_P (decl))
 	    {
 	      if ((ctx->region_type & ORT_TARGET) != 0
@@ -10857,10 +10888,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 		      = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
 		}
 	    }
-	  else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
-		   && (code == OMP_TARGET_EXIT_DATA
-		       || code == OACC_EXIT_DATA))
-	    remove = true;
 	  else if (DECL_SIZE (decl)
 		   && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
 		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
diff --git a/gcc/testsuite/g++.dg/gomp/target-3.C b/gcc/testsuite/g++.dg/gomp/target-3.C
index fe2e38b46a3..f4d40ec8e4b 100644
--- a/gcc/testsuite/g++.dg/gomp/target-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-3.C
@@ -33,4 +33,4 @@ T<N>::bar (int x)
 
 template struct T<0>;
 
-/* { dg-final { scan-tree-dump-times "map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*this \\\[len: 2\\\]\\) map\\(alloc:this->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:this->b \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C b/gcc/testsuite/g++.dg/gomp/target-this-2.C
index a5e832130fb..679c85a54dd 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-2.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -46,4 +46,4 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {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\]\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C b/gcc/testsuite/g++.dg/gomp/target-this-3.C
index 208ea079b95..a450b3723e5 100644
--- a/gcc/testsuite/g++.dg/gomp/target-this-3.C
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -100,6 +100,6 @@ int main (void)
   return 0;
 }
 
-/* { 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:\*this->refptr \[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\(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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C b/gcc/testsuite/g++.dg/gomp/target-this-4.C
index f42cf384541..af23cbb9023 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\(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: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\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias: 0\]\) map\(attach:_3->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" } } */
+/* { 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" } } */