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.

diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 8b75e7aab6389858b47534adae8e1fe052971546..8c91822ef1e2a4440b6e9204592fe608a1089f6a 100644 (file)
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15765,7 +15765,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 efcd1e1620f71af9516f88d1ce3ad7b0f7646e21..d71d69c5b7f5ea85854ec0271b7cf5fc8090e72d 100644 (file)
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13093,6 +13093,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))
@@ -13954,6 +13960,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;
@@ -14707,6 +14714,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)))
@@ -14763,6 +14775,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_)
            {
@@ -14835,6 +14856,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 dc3e085c3025df9a589d22991345c32940ff1d9f..03fbe461bd5cdb6361cdc5b1ab1d4cf9e6f246f4 100644 (file)
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -9264,7 +9264,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
@@ -9365,7 +9367,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)
@@ -9408,6 +9421,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 (file)
index 0000000..3e79212
--- /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" } } */