* [PATCH, OpenMP 5.0] Allow duplicate mapping of variables
@ 2020-12-16 15:05 Chung-Lin Tang
0 siblings, 0 replies; only message in thread
From: Chung-Lin Tang @ 2020-12-16 15:05 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Catherine Moore, Andrew Stubbs
[-- Attachment #1: Type: text/plain, Size: 1790 bytes --]
Hi Jakub,
when running sollve_vv tests (which we're trying to make work out-of-the-box),
some of the 5.0 tests which were testing map clause ordering also required
supporting mapping the same variable multiple times.
This was not allowed before, so besides the small changes needed in the
front-ends, the omp-lowering conventions for field lookup had to be adjusted.
For map clauses, we now use the whole OMP_CLAUSE expression to lookup
fields, instead of the variable DECL, which will clash and ICE when
multiple occurrences happen.
The modifications OpenACC firstprivate variables exist because they seem
to share code with normal map handling in many ways, so they were adjusted
to use the clause-as-key convention too (including the adjustments in
lower_oacc_reductions).
This has been tested for no regressions for gcc, g++, gfortran, and libgomp.
Is this okay for trunk after stage1 reopens?
Thanks,
Chung-Lin
2020-12-16 Chung-Lin Tang <cltang@codesourcery.com>
gcc/c/
* c-typeck.c (c_finish_omp_clauses): Adjust to allow duplicate
mapped variables for OpenMP.
gcc/cp/
* semantics.c (finish_omp_clauses): Adjust to allow duplicate
mapped variables for OpenMP.
gcc/
* omp-low.c (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.
(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.
(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/
* c-c++-common/gomp/clauses-2.c: Adjust testcase.
[-- Attachment #2: dup-map-clause.patch --]
[-- Type: text/plain, Size: 10395 bytes --]
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 7d58e8de342..6bf57ae2ba1 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -14776,12 +14776,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
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),
"%qD appears more than once in motion clauses", t);
else if (ort == C_ORT_ACC)
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 92e32c8e0ad..4159b770613 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7643,11 +7643,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
else
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),
"%qD appears more than once in motion clauses", t);
else if (ort == C_ORT_ACC)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 91a5e3d1431..f8109d9752a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -761,24 +761,31 @@ build_sender_ref (tree var, omp_context *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));
gcc_assert ((mask & 2) == 0 || !ctx->sfield_map
|| !splay_tree_lookup (ctx->sfield_map, key));
@@ -1306,11 +1313,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IS_DEVICE_PTR)
&& 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_gimple_omp_oacc (ctx->stmt) ? 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
install_var_field (decl, false, 3, ctx);
}
@@ -1517,23 +1531,23 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree decl2 = DECL_VALUE_EXPR (decl);
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);
}
else
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& 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)
&& !OMP_CLAUSE_MAP_IN_REDUCTION (c))
install_var_local (decl, ctx);
}
}
@@ -1563,11 +1577,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
tree field
= build_decl (OMP_CLAUSE_LOCATION (c),
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);
}
}
break;
@@ -7019,10 +7033,11 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,
gcc_checking_assert (!is_oacc_kernels (ctx));
/* Likewise, on OpenACC 'kernels' decomposed parts. */
gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx));
tree orig = OMP_CLAUSE_DECL (c);
+ tree orig_clause;
tree var = maybe_lookup_decl (orig, ctx);
tree ref_to_res = NULL_TREE;
tree incoming, outgoing, v1, v2, v3;
bool is_private = false;
@@ -7089,14 +7104,24 @@ 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_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_private)
{
- 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);
tree type = TREE_TYPE (var);
if (POINTER_TYPE_P (type))
@@ -11907,19 +11932,19 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (maybe_lookup_field (c, ctx));
map_cnt++;
continue;
}
- if (!maybe_lookup_field (var, ctx))
+ if (!maybe_lookup_field (c, ctx))
continue;
/* Don't remap compute constructs' reduction variables, because the
intermediate result must be local to each gang. */
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
{
- x = build_receiver_ref (var, true, ctx);
+ x = build_receiver_ref (c, true, ctx);
tree new_var = lookup_decl (var, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -12137,11 +12162,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
nc = OMP_CLAUSE_CHAIN (c);
ovar = OMP_CLAUSE_DECL (nc);
}
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);
nc = NULL_TREE;
}
@@ -12155,11 +12180,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
ovar2 = TREE_OPERAND (ovar2, 0);
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)))
continue;
}
@@ -12179,11 +12204,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimplify_assign (x, build_fold_addr_expr (var), &ilist);
}
else if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
- x = build_sender_ref (ovar, ctx);
+ x = build_sender_ref (c, ctx);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
@@ -12877,11 +12902,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
decl2 = TREE_OPERAND (decl2, 0);
gcc_assert (DECL_P (decl2));
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)))
{
tree bias = OMP_CLAUSE_SIZE (c);
if (DECL_P (bias))
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index bbc8fb4e32b..f1ce7399310 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -13,11 +13,11 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (p);
#pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
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);
#pragma omp target map (t) map (t.r)
bar (&t.r);
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2020-12-16 15:05 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-12-16 15:05 [PATCH, OpenMP 5.0] Allow duplicate mapping of variables Chung-Lin Tang
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).