Arrow operator handling for C front-end in OpenMP map clauses

This patch merges some of the equivalent changes already done for the C++
front-end to the C parts.

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

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_clause_map): Set 'allow_deref' argument in
	call to c_parser_omp_variable_list to 'true'.
	* c-typeck.c (handle_omp_array_sections_1): Add strip of MEM_REF in
	array base handling.
	(c_finish_omp_clauses): Handle 'A->member' case in map clauses.

gcc/ChangeLog:

	* gimplify.c (gimplify_scan_omp_clauses): Add MEM_REF case when
	handling component_ref_p case. Add unshare_expr and gimplification
	when created GOMP_MAP_STRUCT is not a DECL. Add code to add
	firstprivate pointer for *pointer-to-struct case.

gcc/testsuite/ChangeLog:

	* gcc.dg/gomp/target-3.c: New test.

4 files changed, 79 insertions, 3 deletions

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index fae597128e9..0a6aee439f6 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15700,7 +15700,8 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
 	}
     }
 
-  nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list);
+  nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
+				   C_ORT_OMP, true);
 
   for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
     OMP_CLAUSE_SET_MAP_KIND (c, kind);
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 6af19766324..7c887a80ce9 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12917,6 +12917,12 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 		  return error_mark_node;
 		}
 	      t = TREE_OPERAND (t, 0);
+	      if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+		  && TREE_CODE (t) == MEM_REF)
+		{
+		  t = TREE_OPERAND (t, 0);
+		  STRIP_NOPS (t);
+		}
 	      if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
 		{
 		  if (maybe_ne (mem_ref_offset (t), 0))
@@ -13778,6 +13784,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   tree ordered_clause = NULL_TREE;
   tree schedule_clause = NULL_TREE;
   bool oacc_async = false;
+  bool indir_component_ref_p = false;
   tree last_iterators = NULL_TREE;
   bool last_iterators_remove = false;
   tree *nogroup_seen = NULL;
@@ -14505,6 +14512,11 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    {
 		      while (TREE_CODE (t) == COMPONENT_REF)
 			t = TREE_OPERAND (t, 0);
+		      if (TREE_CODE (t) == MEM_REF)
+			{
+			  t = TREE_OPERAND (t, 0);
+			  STRIP_NOPS (t);
+			}
 		      if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14561,6 +14573,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	       bias) to zero here, so it is not set erroneously to the pointer
 	       size later on in gimplify.c.  */
 	    OMP_CLAUSE_SIZE (c) = size_zero_node;
+	  indir_component_ref_p = false;
+	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
+	      && TREE_CODE (t) == COMPONENT_REF
+	      && TREE_CODE (TREE_OPERAND (t, 0)) == MEM_REF)
+	    {
+	      t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
+	      indir_component_ref_p = true;
+	      STRIP_NOPS (t);
+	    }
 	  if (TREE_CODE (t) == COMPONENT_REF
 	      && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
 	    {
@@ -14633,6 +14654,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  else if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
 		    || (OMP_CLAUSE_MAP_KIND (c)
 			!= GOMP_MAP_FIRSTPRIVATE_POINTER))
+		   && !indir_component_ref_p
 		   && !c_mark_addressable (t))
 	    remove = true;
 	  else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0a6ceb584b1..ff44d26c98e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9197,7 +9197,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
 	      if ((DECL_P (decl)
-		   || (component_ref_p && INDIRECT_REF_P (decl)))
+		   || (component_ref_p
+		       && (INDIRECT_REF_P (decl)
+			   || TREE_CODE (decl) == MEM_REF)))
 		  && 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
@@ -9298,7 +9300,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      if (base_ref)
 			OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
 		      else
-			OMP_CLAUSE_DECL (l) = decl;
+			{
+			  OMP_CLAUSE_DECL (l) = unshare_expr (decl);
+			  if (!DECL_P (OMP_CLAUSE_DECL (l))
+			      && (gimplify_expr (&OMP_CLAUSE_DECL (l),
+						 pre_p, NULL, is_gimple_lvalue,
+						 fb_lvalue)
+				  == GS_ERROR))
+			    {
+			      remove = true;
+			      break;
+			    }
+			}
 		      OMP_CLAUSE_SIZE (l)
 			= (!attach
 			   ? size_int (1)
@@ -9341,6 +9354,30 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			flags |= GOVD_SEEN;
 		      if (has_attachments)
 			flags |= GOVD_MAP_HAS_ATTACHMENTS;
+
+		      /* If this is a *pointer-to-struct expression, make sure a
+			 firstprivate map of the base-pointer exists.  */
+		      if (component_ref_p
+			  && ((TREE_CODE (decl) == MEM_REF
+			       && integer_zerop (TREE_OPERAND (decl, 1)))
+			      || INDIRECT_REF_P (decl))
+			  && DECL_P (TREE_OPERAND (decl, 0))
+			  && !splay_tree_lookup (ctx->variables,
+						 ((splay_tree_key)
+						  TREE_OPERAND (decl, 0))))
+			{
+			  decl = TREE_OPERAND (decl, 0);
+			  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						      OMP_CLAUSE_MAP);
+			  enum gomp_map_kind mkind
+			    = GOMP_MAP_FIRSTPRIVATE_POINTER;
+			  OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+			  OMP_CLAUSE_DECL (c2) = decl;
+			  OMP_CLAUSE_SIZE (c2) = size_zero_node;
+			  OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+			  OMP_CLAUSE_CHAIN (c) = c2;
+			}
+
 		      if (DECL_P (decl))
 			goto do_add_decl;
 		    }
diff --git a/gcc/testsuite/gcc.dg/gomp/target-3.c b/gcc/testsuite/gcc.dg/gomp/target-3.c
new file mode 100644
index 00000000000..3e7921270c9
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/gomp/target-3.c
@@ -0,0 +1,16 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+struct S
+{
+  int a, b;
+};
+
+void foo (struct S *s)
+{
+  #pragma omp target map (alloc: s->a, s->b)
+    ;
+  #pragma omp target enter data map (alloc: s->a, s->b)
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(struct:\\*s \\\[len: 2\\\]\\) map\\(alloc:s->a \\\[len: \[0-9\]+\\\]\\) map\\(alloc:s->b \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */