[gcc/devel/omp/gcc-14] OpenMP 5.0: Allow multiple clauses mapping same variable
Paul-Antoine Arras
parras@gcc.gnu.org
Fri Jun 28 09:49:23 GMT 2024
https://gcc.gnu.org/g:6bb82ce8294046ec11d9f9a3bd84ad5cb0a7e01c
commit 6bb82ce8294046ec11d9f9a3bd84ad5cb0a7e01c
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date: Mon Feb 1 03:16:47 2021 -0800
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/cp/ChangeLog:
* semantics.cc (finish_omp_clauses): Adjust to allow duplicate
mapped variables for OpenMP.
gcc/ChangeLog:
* omp-low.cc (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:
---
gcc/ChangeLog.omp | 15 +++++
gcc/cp/ChangeLog.omp | 5 ++
gcc/omp-low.cc | 89 +++++++++++++++++++----------
gcc/testsuite/ChangeLog.omp | 4 ++
gcc/testsuite/c-c++-common/gomp/clauses-2.c | 2 +-
gcc/testsuite/c-c++-common/gomp/map-6.c | 4 +-
6 files changed, 87 insertions(+), 32 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index e35ced18751..d000c3614a8 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,18 @@
+2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * omp-low.cc (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.
+
2021-01-15 Andrew Stubbs <ams@codesourcery.com>
* dwarf2out.cc (add_location_or_const_value_attribute): Set
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index a87d118e4bb..bfdc3be2c4e 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * semantics.cc (finish_omp_clauses): Adjust to allow duplicate
+ mapped variables for OpenMP.
+
2022-02-03 Kwok Cheung Yeung <kcy@codesourcery.com>
* parser.ccc (cp_parser_omp_clause_map): Update call to
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 2db80e2498b..fe6e20f444a 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -776,24 +776,28 @@ build_sender_ref (tree var, omp_context *ctx)
return build_sender_ref ((splay_tree_key) var, ctx);
}
-/* Add a new field for VAR inside the structure CTX->SENDER_DECL. If
- BASE_POINTERS_RESTRICT, declare the field with restrict. */
-
static void
-install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
+ tree key_expr = NULL_TREE)
{
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));
@@ -1495,8 +1499,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
|| (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR
&& lang_hooks.decls.omp_array_data (decl, true)))
{
+ /* 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);
by_ref = !omp_privatize_by_reference (decl);
- install_var_field (decl, by_ref, 3, ctx);
+ install_var_field (decl, by_ref, 3, ctx, k);
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
@@ -1806,7 +1815,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;
}
@@ -1820,7 +1829,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
gcc_assert (INDIRECT_REF_P (decl2));
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);
}
@@ -1830,9 +1839,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);
if (is_gimple_omp_offloaded (ctx->stmt)
&& !(is_gimple_omp_oacc (ctx->stmt)
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
@@ -1867,7 +1876,7 @@ 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);
}
}
@@ -4623,8 +4632,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;
}
@@ -7620,6 +7640,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
tree orig = OMP_CLAUSE_DECL (c);
+ tree orig_clause;
tree var;
tree ref_to_res = NULL_TREE;
tree incoming, outgoing;
@@ -7702,8 +7723,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);
@@ -7715,7 +7746,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_privatize_by_reference (orig))
ref_to_res = build_simple_mem_ref (ref_to_res);
}
@@ -13102,7 +13133,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
@@ -13122,7 +13153,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)
@@ -13383,7 +13414,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 = ovar;
if (in_reduction_clauses
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13431,7 +13462,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)))
@@ -13481,7 +13512,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
else if (nc)
{
- 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
@@ -14396,7 +14427,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
type = TREE_TYPE (type);
ref_to_ptr = true;
}
- 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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 1ac91cf1b01..6ff9a06aab7 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,7 @@
+2021-02-01 Chung-Lin Tang <cltang@codesourcery.com>
+
+ * c-c++-common/gomp/clauses-2.c: Adjust testcase.
+
2020-08-21 Tobias Burnus <tobias@codesourcery.com>
* gfortran.dg/gomp/pr67500.f90: Change dg-warning to
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index 8f98d57a312..b4b50045cae 100644
--- 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);
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 014ed35ab41..1172cdbe818 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -157,10 +157,10 @@ foo (void)
#pragma omp target map (always, close)
;
- #pragma omp target map (always, always) /* { dg-error "'always' appears more than once in map clauses" } */
+ #pragma omp target map (always, always)
;
- #pragma omp target map (always, always, close) /* { dg-error "'always' appears more than once in map clauses" } */
+ #pragma omp target map (always, always, close)
;
#pragma omp target map (always, close, to: always, close, b7)
More information about the Gcc-cvs
mailing list