OpenMP 5.0: Allow multiple clauses mapping same variable

This is a merge of:
https://gcc.gnu.org/pipermail/gcc-patches/2020-December/562081.html

This patch now allows multiple clauses on the same construct to map
the same variable, which was not valid in OpenMP 4.5, but now allowed
in 5.0.

This may possibly reverted/updated when a final patch is approved
for mainline.

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

gcc/c/ChangeLog:

	* c-typeck.c (c_finish_omp_clauses): Adjust to allow duplicate
	mapped variables for OpenMP.

gcc/cp/ChangeLog:

	* semantics.c (finish_omp_clauses):  Adjust to allow duplicate
	mapped variables for OpenMP.

gcc/ChangeLog:

	* omp-low.c (install_parm_decl): Add new 'tree key_expr' parameter.
	Use key_expr as splay-tree key instead of var itself.
	(install_var_field): Add new 'tree key_expr = NULL_TREE'
	default parameter. Set splay-tree lookup key to key_expr instead of
	var if key_expr is non-NULL. Adjust call to install_parm_decl.
	Update comments.
	(scan_sharing_clauses): Use clause tree expression as splay-tree key
	for map/to/from and OpenACC firstprivate cases when installing the
	variable field into the send/receive record type.
	(maybe_lookup_field_in_outer_ctx): Add code to search through
	construct clauses instead of entirely based on splay-tree lookup.
	(lower_oacc_reductions): Adjust to find map-clause of reduction
	variable, then create receiver-ref.
	(lower_omp_target): Adjust to lookup var field using clause expression.

gcc/testsuite/ChangeLog:

	* c-c++-common/gomp/clauses-2.c: Adjust testcase.

diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 651f29bb21ba66ea70190c878649034896a1aa08..9c8e37fd1e9af241477248bfe4c5c06f873d10f1 100644 (file)
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14671,8 +14671,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
                bitmap_set_bit (&generic_head, DECL_UID (t));
            }
          else if (bitmap_bit_p (&map_head, DECL_UID (t))
-                  && (ort != C_ORT_OMP
-                      || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+                  && !bitmap_bit_p (&map_field_head, DECL_UID (t))
+                  && ort != C_ORT_OMP)
            {
              if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
                error_at (OMP_CLAUSE_LOCATION (c),

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index b4544f78aa7df0b803809689f63665b67efb1cb2..30d94aae960f23944699c4976cec6205b04dd696 100644 (file)
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7628,7 +7628,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
                bitmap_set_bit (&generic_head, DECL_UID (t));
            }
          else if (bitmap_bit_p (&map_head, DECL_UID (t))
-                  && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
+                  && !bitmap_bit_p (&map_field_head, DECL_UID (t))
+                  && ort != C_ORT_OMP)
            {
              if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
                error_at (OMP_CLAUSE_LOCATION (c),

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ddeceba2bc7024f98390f8d10d50fc0a336d8219..7888006495e4b4b02ddb1b7ff622e02b99a6a556 100644 (file)
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -788,12 +788,12 @@ build_sender_ref (tree var, omp_context *ctx)
 }
 
 static void
-install_parm_decl (tree var, tree type, omp_context *ctx)
+install_parm_decl (tree key_expr, tree var, tree type, omp_context *ctx)
 {
   if (!is_oacc_parallel_or_serial (ctx) || !targetm.goacc.explode_args ())
     return;
 
-  splay_tree_key key = (splay_tree_key) var;
+  splay_tree_key key = (splay_tree_key) key_expr;
   tree decl_name = NULL_TREE, t;
   location_t loc = UNKNOWN_LOCATION;
 
@@ -813,25 +813,34 @@ install_parm_decl (tree var, tree type, omp_context *ctx)
   splay_tree_insert (ctx->parm_map, key, (splay_tree_value) t);
 }
 
-/* Add a new field for VAR inside the structure CTX->SENDER_DECL.  If
-   BASE_POINTERS_RESTRICT, declare the field with restrict.  */
+/* Add a new field for VAR inside the structure CTX->SENDER_DECL.
+   If KEY_EXPR is passed a non-NULL_TREE, use that tree expression as the
+   splay_tree_key for the field, instead of VAR itself.
+   If BASE_POINTERS_RESTRICT, declare the field with restrict.  */
 
 static void
 install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
+                  tree key_expr = NULL_TREE,
                   bool base_pointers_restrict = false)
 {
   tree field, type, sfield = NULL_TREE;
   splay_tree_key key = (splay_tree_key) var;
 
-  if ((mask & 16) != 0)
-    {
-      key = (splay_tree_key) &DECL_NAME (var);
-      gcc_checking_assert (key != (splay_tree_key) var);
-    }
-  if ((mask & 8) != 0)
+  if (key_expr)
+    /* Allow user to explicitly set the expression used as the key.  */
+    key = (splay_tree_key) key_expr;
+  else
     {
-      key = (splay_tree_key) &DECL_UID (var);
-      gcc_checking_assert (key != (splay_tree_key) var);
+      if ((mask & 16) != 0)
+       {
+         key = (splay_tree_key) &DECL_NAME (var);
+         gcc_checking_assert (key != (splay_tree_key) var);
+       }
+      if ((mask & 8) != 0)
+       {
+         key = (splay_tree_key) &DECL_UID (var);
+         gcc_checking_assert (key != (splay_tree_key) var);
+       }
     }
   gcc_assert ((mask & 1) == 0
              || !splay_tree_lookup (ctx->field_map, key));
@@ -922,7 +931,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
   if (mask & 1)
     {
       splay_tree_insert (ctx->field_map, key, (splay_tree_value) field);
-      install_parm_decl (var, type, ctx);
+      install_parm_decl ((tree) key, var, type, ctx);
     }
   if ((mask & 2) && ctx->sfield_map)
     splay_tree_insert (ctx->sfield_map, key, (splay_tree_value) sfield);
@@ -1459,7 +1468,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
              && is_gimple_omp_offloaded (ctx->stmt))
            {
              if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
-               install_var_field (decl, !omp_is_reference (decl), 3, ctx);
+               {
+                 /* OpenACC firstprivate clauses are later processed with same
+                    code path as map clauses in lower_omp_target, so follow
+                    the same convention of using the whole clause expression
+                    as splay-tree key.  */
+                 tree k = (is_oacc_parallel_or_serial (ctx) ? c : NULL_TREE);
+                 install_var_field (decl, !omp_is_reference (decl), 3, ctx, k);
+               }
              else if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
                install_var_field (decl, true, 3, ctx);
              else
@@ -1677,7 +1693,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
                  t = TREE_TYPE (t);
                }
 
-             install_var_field (array_decl, by_ref, 3, ctx);
+             install_var_field (array_decl, by_ref, 3, ctx, c);
              install_var_local (array_decl, ctx);
              break;
            }
@@ -1691,7 +1707,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
                  gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
                  decl2 = TREE_OPERAND (decl2, 0);
                  gcc_assert (DECL_P (decl2));
-                 install_var_field (decl2, true, 3, ctx);
+                 install_var_field (decl2, true, 3, ctx, c);
                  install_var_local (decl2, ctx);
                  install_var_local (decl, ctx);
                }
@@ -1701,9 +1717,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
                      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
                      && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
                      && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
-                   install_var_field (decl, true, 7, ctx);
+                   install_var_field (decl, true, 7, ctx, c);
                  else
-                   install_var_field (decl, true, 3, ctx,
+                   install_var_field (decl, true, 3, ctx, c,
                                       base_pointers_restrict);
                  if (is_gimple_omp_offloaded (ctx->stmt)
                      && !OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -1738,9 +1754,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
                                  FIELD_DECL, NULL_TREE, ptr_type_node);
                  SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
                  insert_field_into_struct (ctx->record_type, field);
-                 splay_tree_insert (ctx->field_map, (splay_tree_key) decl,
+                 splay_tree_insert (ctx->field_map, (splay_tree_key) c,
                                     (splay_tree_value) field);
-                 install_parm_decl (decl, ptr_type_node, ctx);
+                 install_parm_decl (c, decl, ptr_type_node, ctx);
                }
            }
          break;
@@ -4479,8 +4495,19 @@ maybe_lookup_field_in_outer_ctx (tree decl, omp_context *ctx)
   omp_context *up;
 
   for (up = ctx->outer; up; up = up->outer)
-    if (maybe_lookup_field (decl, up))
-      return true;
+    {
+      for (tree c = gimple_omp_target_clauses (up->stmt);
+          c != NULL_TREE; c = OMP_CLAUSE_CHAIN (c))
+       if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+            || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+            || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM
+            || (is_oacc_parallel_or_serial (up)
+                && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE))
+           && OMP_CLAUSE_DECL (c) == decl)
+         return true;
+      if (maybe_lookup_field (decl, up))
+       return true;
+    }
 
   return false;
 }
@@ -7199,6 +7226,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
     if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
       {
        tree orig = OMP_CLAUSE_DECL (c);
+       tree orig_clause;
        tree var;
        tree ref_to_res = NULL_TREE;
        tree incoming, outgoing;
@@ -7284,8 +7312,18 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
          do_lookup:
            /* This is the outermost construct with this reduction,
               see if there's a mapping for it.  */
-           if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET
-               && (maybe_lookup_field (orig, outer) || is_fpp) && !is_private)
+           orig_clause = NULL_TREE;
+           if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET)
+             for (tree cls = gimple_omp_target_clauses (outer->stmt);
+                  cls; cls = OMP_CLAUSE_CHAIN (cls))
+               if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_MAP
+                   && orig == OMP_CLAUSE_DECL (cls)
+                   && maybe_lookup_field (cls, outer))
+                 {
+                   orig_clause = cls;
+                   break;
+                 }
+           if ((orig_clause != NULL_TREE || is_fpp) && !is_private)
              {
                tree type = TREE_TYPE (var);
 
@@ -7297,7 +7335,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
                  }
                else
                  {
-                   ref_to_res = build_receiver_ref (orig, false, outer);
+                   ref_to_res = build_receiver_ref (orig_clause, false, outer);
                    if (omp_is_reference (orig))
                      ref_to_res = build_simple_mem_ref (ref_to_res);
                  }
@@ -12188,7 +12226,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                continue;
              }
 
-           if (!maybe_lookup_field (var, ctx))
+           if (!maybe_lookup_field (c, ctx))
              continue;
 
            init_cnt++;
@@ -12330,7 +12368,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            continue;
          }
 
-       if (!maybe_lookup_field (var, ctx))
+       if (!maybe_lookup_field (c, ctx))
          continue;
 
        /* Don't remap compute constructs' reduction variables, because the
@@ -12348,7 +12386,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
               && TREE_CODE (var_type) != ARRAY_TYPE
               ? false : true);
 
-           x = build_receiver_ref (var, rcv_by_ref, ctx);
+           x = build_receiver_ref (c, rcv_by_ref, ctx);
 
            if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
                && (FLOAT_TYPE_P (inner_type)
@@ -12581,7 +12619,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                  }
                else
                  {
-                   tree x = build_sender_ref (ovar, ctx);
+                   tree x = build_sender_ref (c, ctx);
                    tree v
                      = build_fold_addr_expr_with_type (ovar, ptr_type_node);
                    gimplify_assign (x, v, &ilist);
@@ -12599,7 +12637,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                    gcc_assert (DECL_P (ovar2));
                    ovar = ovar2;
                  }
-               if (!maybe_lookup_field (ovar, ctx)
+               if (!maybe_lookup_field (c, ctx)
                    && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
                         && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
                             || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
@@ -12623,7 +12661,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
            else if (nc)
              {
                var = lookup_decl_in_outer_ctx (ovar, ctx);
-               x = build_sender_ref (ovar, ctx);
+               x = build_sender_ref (nc, ctx);
 
                if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
                    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -12768,7 +12806,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
              s = integer_one_node;
            s = fold_convert (size_type_node, s);
            if (oacc_explode_args)
-             decl_args = append_decl_arg (ovar, decl_args, ctx);
+             {
+               tree key = (nc != NULL_TREE
+                           && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+                           && OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
+                           ? nc : c);
+               decl_args = append_decl_arg (key, decl_args, ctx);
+             }
            purpose = size_int (map_idx++);
            CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
            if (TREE_CODE (s) != INTEGER_CST)
@@ -13399,7 +13443,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
                    new_var = decl2;
                    type = TREE_TYPE (new_var);
                  }
-               x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx);
+               x = build_receiver_ref (prev, false, ctx);
                x = fold_convert_loc (clause_loc, type, x);
                if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
                  {

diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index bbc8fb4e32bf81267293de248057c42558f88312..f1ce73993108b91096325546d8bc0db9781485f0 100644 (file)
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -15,7 +15,7 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
     bar (p);
   #pragma omp target map (p) , map (p[0])
     bar (p);
-  #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (q) map (q)
     bar (&q);
   #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
     bar (p);