@Thomas – this code also affects OpenACC. Thus, you might want
to have a look as well.
I think we have two options: Make this OpenMP only or
also use this for OpenACC code. (See also second comment
below.) - WDYT?
* * *
Paul-Antoine Arras wrote:
This patch allows multiple clauses on the same construct to map the same
variable, which was not valid in OpenMP 4.5, but allowed in 5.0.
Thanks for the patch!
The following example - related to pointer attachment - shows why
it makes sense to permit this (cf. iterator example in OpenMP_VV
+ PA's own follow up email.)
#pragma omp target map (iterator(it = 0:DIM), tofrom: x[it][:1]) map (to: x)
(→ libgomp.c-c++-common/target-map-iterators-6.c )
* * *
[BUG]
I think we need to file a PR for the following issue (unchanged by your patch);
that's C and C++ (Fortran is fine):
void f() {
int x, y;
#pragma omp target map(tofrom:x,y) private(y)
;
#pragma omp target map(tofrom:x,y) firstprivate(x)
;
}
gives:
test.c:5:31: error: ‘x’ appears both in data and map clauses
5 | #pragma omp target map(tofrom:x,y) firstprivate(x)
That's correct – but IMHO it should give the very same error also for PRIVATE.
* * *
[OpenACC]
void f() {
int X;
#pragma acc data copyin(X) copyout(X)
X += 5;
}
Before this patch:
foo.c:3:26: error: ‘X’ appears more than once in data clauses
3 | #pragma acc data copyin(X) copyout(X)
After the patch:
a-foo.c.006t.original: #pragma acc data map(tofrom:X)
a-foo.c.007t.gimple: #pragma omp target oacc_data map(tofrom:X [len: 4])
a-foo.c.010t.omplower: #pragma omp target oacc_data map(tofrom:X [len:
4])
I believe that's fine - also light of the still open OpenACC Issue #552 (+ pull req. draft #554). It seems as if OpenACC indented to permit
specifying the same variable multiple times (at least it does not
disallow it). With this patch, the following option is handled: Comment
from August last year (listening multiple options): "If we only want to handle perfect aliasing, then sorting clauses at
compile time is probably sufficient"
If more should be supported, doing the simple cases at compile time is
fine - leaving the rest to the runtime. Otherwise, it is invalid, but
likely usually not detected because it is unlikely that one implements
an expensive check in the runtime.
Thus, we should be fine.
HOWEVER: I guess, we need some some testcases for this OpenACC case?
Alternatively, we have to/could exclude it from applying to OpenACC.
NOTE: Contrary to OpenMP, there is currently no testcase in this
regard at all - neither one checking that merging workes by this
patch nor any testcase that checks for 'Error: Symbol ‘x’ present
on multiple clauses' (which existed before for OpenMP and are
now removed.)
* * *
[FYI only]
For general information - but no action required for this comment and the
next two.
I note that in OpenMP 6.x, 'delete' has become a modifier and
omp target map(delete,always,present : x)
now works as combination as well. That's rather separate from
this implementation, however,
* * *
[FYI only]
I was wondering whether the the following should be merged:
void f() {
int x;
#pragma omp target exit data map(release: x) map(delete: x)
#pragma omp target exit data map(release: x) map(from: x)
}
Namely, the first one to 'delete' and the second one to 'from',
but while it would be possible, there is no issue when not done
so - as not present does not cause an issue.
The following requires that 'from' comes first - but the sorting
ensures this:
#pragma omp target exit data map(delete: x) map(from: x)
#pragma omp target exit data map(from: x) map(delete: x)
Thus, there is no issue, either.
→ Those aren't handled, but they also do not need to handled.
* * *
[FYI only]
The one is not handled in the FE - but the ME ensures it works:
#pragma omp target enter data map(present,alloc: x) map(always,to: x)
becomes in gimplify
#pragma omp target enter data map(force_present:x [len: 4]) map(always,to:x
[len: 4])
* * *
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -17872,7 +17874,7 @@ c_finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
}
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
- && ort != C_ORT_OMP
+ && ort != C_ORT_OMP && ort != C_ORT_OMP_TARGET
&& ort != C_ORT_OMP_EXIT_DATA)
{
I think the coding style at least in that FE is to have the '&&'
aligned, i.e. C_ORT_OMP_TARGET on a line of its own.
* * *
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -9588,7 +9589,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type
ort)
bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
else if (bitmap_bit_p (&map_head, DECL_UID (t))
&& !bitmap_bit_p (&map_field_head, DECL_UID (t))
- && ort != C_ORT_OMP
+ && ort != C_ORT_OMP && ort != C_ORT_OMP_TARGET
&& ort != C_ORT_OMP_EXIT_DATA)
{
Likewise.
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
...
+tree
+omp_remove_duplicate_maps (tree clauses, bool first)
+ bool maybe_dup_found
+ = (OMP_CLAUSE_CODE (c1) == OMP_CLAUSE_CODE (c2)
+ && ((OMP_CLAUSE_DECL (c1) == NULL_TREE
+ && OMP_CLAUSE_DECL (c2) == NULL_TREE)
+ || operand_equal_p (OMP_CLAUSE_DECL (c1),
+ OMP_CLAUSE_DECL (c2)))
I believe this cannot occur; we have sometimes NULL_POINTER for 'omp_…'
special identifiers and we might have some error conditions, but I doubt
we generate any map entry for those – and having an error_mark_node
seems to be more likely if something goes wrong later.
However regarding:
+ && ((OMP_CLAUSE_SIZE (c1) == NULL_TREE
+ && OMP_CLAUSE_SIZE (c2) == NULL_TREE)
+ || (OMP_CLAUSE_SIZE (c1) != NULL_TREE
+ && OMP_CLAUSE_SIZE (c2) != NULL_TREE
+ && operand_equal_p (OMP_CLAUSE_SIZE (c1),
+ OMP_CLAUSE_SIZE (c2))))
The NULL_TREE seems to occur mostly when called from the front end as
the size is not yet filled in. (It might also be the case during the
gimplification, I have not checked what comes first.)
+ && ((OMP_CLAUSE_ITERATORS (c1) == NULL_TREE
+ && OMP_CLAUSE_ITERATORS (c2) == NULL_TREE)
+ || operand_equal_p (OMP_CLAUSE_ITERATORS (c1),
+ OMP_CLAUSE_ITERATORS (c2))));
while this one is obvious.
Maybe a comment makes sense that the OMP_CLAUSE_DECL is just to be prudent but
it shouldn't occur while OMP_CLAUSE_SIZE often gets filled in only later?
(In any case, it is now documented here.)
* * *
+ if (first)
…
+ /* When called from the gimplifier, only compare bits up to
+ GOMP_MAP_FLAG_SPECIAL_2. Higher flags may need to be present
+ multiple times.*/
This might be clearer if worded as
When called from the gimplifier, remove duplicate map clauses with
identical kind only when the bits above GOMP_MAP_FLAG_SPECIAL_2 are
unset - as clauses with those flags set may need to be present
multiple times.
The current wording it not really wrong, but I first read it backward.
+ else if (OMP_CLAUSE_MAP_KIND (c1) == OMP_CLAUSE_MAP_KIND (c2)
+ && (OMP_CLAUSE_MAP_KIND (c1) & ~0b11111) == 0)
This means For the FE handling merge all clauses that are identical plus
those that only differ in the TO/FROM flag (alloc = 00, to = 01,
from = 10, tofrom = 11) by taking the superset.
But during the ME/gimplify merge, only merge those that are fully identical
and even if they are fully identical, don't remove duplicated occurrences if
any of the following special flags are set:
DEEP_COPY = 010|100|00
IMPLICIT = 011|000|00
FORCE = 100|000|00
PRESENT = 100|001|00
ALWAYS_PRESENT = 100|101|00
GOMP_MAP_LAST =1000|000|00
However, still remove any that aren't those or combined with those
(i.e. ALWAYS_PRESENT_{TO,FROM,TOFROM}, (all with map_last:)
firstprivate(_pointer), attach/detach, present+(alloc,to,from,tofrom),
push/pop mapper.)
The ones inside b111'11 are:
00001 is TO
00010 is FROM
00100 = GOMP_MAP_FLAG_SPECIAL_0 with:
001xx is POINTER (00), PSET (01), (FORCE_)PRESENT (10), DELETE(11)
010xx is FORCE_DEVICE_PTR (00), DEVICE_RESIDENT (10), LINK (10), IF_PRESENT (11)
...
* * *
For the following: Thanks for actually documenting the MASK flags;
this was long overdue!
(As a next step, we probably should move to named constants as
currently the values 1, 2, 3, 7, 8|3, 11, 16|3, 19, 32|1 are
passed as numeric constants. The table helps, but still. It is
also interesting to have both 11 and 8|3, which could have been 8|2|1
instead ...)
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -784,24 +784,40 @@ 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. */
+/* Add a new field for VAR inside the structure CTX. If BY_REF is true,
+ use a pointer to the VAR rather than VAR itself.
+ MASK is a bit mask of other options. Bits are interpreted as:
+ 1: Install VAR in ctx->field_map.
+ 2: Install VAR in ctx->sfield_map.
+ 4: VAR is an array, convert it to a pointer.
+ 8: Use DECL_UID (VAR) instead of VAR as key.
+ 16: Use DECL_NAME (VAR) instead of VAR as key.
+ 32: Don't dereference omp_is_reference types.
+ KEY_EXPR allows specifying something other than VAR as the lookup key.
+ If specified, it overrides the 8 and 16 MASK bits. */
I would write 'If specified, it also overrides ...' - I think the 'also'
makes it clearer as by default the tree 'VAR' is used as key.
(The addition of 'key_expr' is the only other change besides
the comment.)
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)
+ if (key_expr)
+ /* Allow caller to explicitly set the expression used as the key. */
+ key = (splay_tree_key) key_expr;
+ else
{
- 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);
* * *
On the caller side, the change is that not the decl but the
clause is used as key – such that 'map(to: x)' and 'map(from: x)'
would give separate entries (except that this special case would be
merged). (As entries with different modifiers (always, present, ...)
aren't merged, this matters.)
@@ -1389,8 +1405,13 @@ scan_sharing_clauses (tree clauses, omp_context *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);
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)
{
@@ -1683,7 +1704,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);
}
@@ -1693,9 +1714,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)))
@@ -1730,7 +1751,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);
}
}
@@ -7453,6 +7474,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 = maybe_lookup_decl (orig, ctx);
tree ref_to_res = NULL_TREE;
tree incoming, outgoing, v1, v2, v3;
* * *
The downside of this approach show up for the following, where
the lookup changes from does the decl exist as the map-lookup
in the enclosing compute region - to: check all map clauses to
see whether they map that decl, for all that do, check whether
that clause is found in the lookup:
@@ -7523,10 +7545,20 @@ 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_privatize_by_reference (orig))
ref_to_res = build_simple_mem_ref (ref_to_res);
@@ -12949,7 +12981,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
@@ -12958,7 +12990,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
&& is_gimple_omp_oacc (ctx->stmt)
&& 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
@@ -13353,7 +13385,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
@@ -13402,7 +13434,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)))
@@ -13452,7 +13484,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
@@ -14416,7 +14448,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)))
{
* * *
LGTM – with the style/comments remarks addressed or at least considered.
Please give Thomas a few days (like early next week) to comment on OpenACC.
And please handle the BUG case, e.g., by creating a problem report by CCing
me.
Thanks again!
Tobias
