Hi Kwok,
thanks for the patch; looks quite good, but I have a couple of remarks:
Kwok Cheung Yeung wrote:
Date: Sat, 3 May 2025 20:24:26 +0000
Subject: [PATCH 02/11] openmp: Add support for iterators in map clauses
(C/C++)
This adds preliminary support for iterators in map clauses within OpenMP
'target' constructs (which includes constructs such as 'target enter data').
Iterators with non-constant loop bounds are not currently supported.
(but will be later in the series)
(Also unsupported for now is Fortran deep mapping, which will be extra fun!)
(For either of the two, there is "sorry" in the code.)
* * *
Regarding gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c:
target-map-iterators-2.c:6:25: warning: iterator variable ‘i’ not used in
clause expression
and
+ warning_at (OMP_CLAUSE_LOCATION (c), 0,
+ "iterator variable %qE not used in clause "
+ "expression", DECL_NAME (var));
Can you replace '0' by some OPT_… such as OPT_Wopenmp?
Additionally, while I think it is okay to have a warning for:
#pragma omp target map(iterator(i2=0:10, j2=0:20), from: x[i2])
/* { dg-warning "iterator variable .j2. not used in clause expression" } */
I find the following warning odd:
#pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30), to: x[i3+j3],
y[j3+k3], z[k3+i3])
;
/* { dg-warning "iterator variable .i3. not used in clause expression" "" {
target *-*-* } .-2 } */
/* { dg-warning "iterator variable .j3. not used in clause expression" "" {
target *-*-* } .-3 } */
/* { dg-warning "iterator variable .k3. not used in clause expression" "" {
target *-*-* } .-4 } */
as all variables appears in expressions!
In the original dump, there is:
map(iterator(int i3=0:10:1, int j3=0:20:1, int k3=0:30:1):
to:*(z + (sizetype) ((long unsigned int) (k3 + i3) * 8)) [len: 8])
I understand that the iterator can be removed as x/y/z are unused, but this
does not make the warning more sensible, either.
* * *
BTW: I wonder whether we should special case:
(a) begin:end with begin == end → single iteration
(b) begin:end with begin < end (zero-sized empty loop)
And we may want to warn for step == 0.
* * *
In any case, it would be good to have a testcase that checks for
empty loops, i.e. begin > end (with step < 0) and begin > end (with step < 0)
And, I guess, it wouldn't harm to have a check for step < 0 with begin > end
(finite loop size but with negative values).
* * *
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
I think we want to have somewhere a comment regarding how iterators are
passed. I think the main one should be in libgomp/target.c but possibly
some hint about the format makes sense. For instance, as part of the
comment to (and/or) lower_omp_map_iterator_{expr,size}.
* * *
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x)));
+ avar = build_fold_addr_expr (avar);
+ avar = lower_omp_map_iterator_expr (avar, c, stmt);
+ gimplify_assign (x, avar, &ilist);
Twice: Line starts with space + tab.
* * *
case GIMPLE_OMP_TARGET:
if (!is_gimple_omp_offloaded (stmt))
{
*handled_ops_p = false;
return NULL_TREE;
}
/* FALLTHRU */
case GIMPLE_OMP_PARALLEL:
case GIMPLE_OMP_TASK:
do_parallel:
{
+ if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+ walk_body (convert_tramp_reference_stmt, convert_tramp_reference_op,
+ info, gimple_omp_target_iterator_loops_ptr (stmt));
I think it would be cleaner to move the walk_body before FALLTHRU.
* * *
--- a/gcc/tree.cc
+++ b/gcc/tree.cc
@@ -323,7 +323,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_EXCLUSIVE */
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
- 2, /* OMP_CLAUSE_MAP */
+ 3, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
3, /* OMP_CLAUSE__MAPPER_BINDING_ */
@@ -11767,6 +11767,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
{
int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)];
+ /* Do not walk the iterator operand of OpenMP MAP clauses. */
+ if (OMP_CLAUSE_HAS_ITERATORS (t))
+ len--;
This looks fragile, i.e. it assumes that we never add a new operator to
OMP_CLAUSE_MAP; I wonder whether should add somewhere a note about this?
I don't want to clutter omp_clause_num_ops but it seems to be the safest place?
* * *
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1003,6 +1003,105 @@ gomp_map_val (struct target_mem_desc *tgt, void
**hostaddrs, size_t i)
}
}
+static const char *
+kind_to_name (unsigned short kind)
+{
+ if (GOMP_MAP_IMPLICIT_P (kind))
+ kind &= ~GOMP_MAP_IMPLICIT;
+
+ switch (kind & 0xff)
+ {
+ case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+ case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+ case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+ case GOMP_MAP_TO: return "GOMP_MAP_TO";
+ case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET";
+ case GOMP_MAP_FROM: return "GOMP_MAP_FROM";
+ case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM";
+ case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER";
+ case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+ case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+ default: return "unknown";
+ }
+}
I think that's a quite useful function; I also have locally a similar function
to debug mapping issues. However, I think you should really cover all kinds,
e.g. I fail to see why 'map(iterator(), always, to:' should print 'unknown'.
If I grep correctly, I see the following GOMP_MAP items being used in target;
I have not checked whether all make sense, though - at least implicit can be
removed for the current code.
Probably, we should then also add the short flag to the function. In the
caller we can pass the fixed value ('true'), but that permits to later reuse
the function simply (locally or also in a committed version) for other map
kinds.
GOMP_MAP_ALLOC
GOMP_MAP_ALWAYS_FROM
GOMP_MAP_ALWAYS_POINTER
GOMP_MAP_ALWAYS_PRESENT_FROM
GOMP_MAP_ALWAYS_PRESENT_TO
GOMP_MAP_ALWAYS_PRESENT_TOFROM
GOMP_MAP_ALWAYS_TO
GOMP_MAP_ALWAYS_TOFROM
GOMP_MAP_ATTACH
GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
GOMP_MAP_DELETE
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
GOMP_MAP_DETACH
GOMP_MAP_FIRSTPRIVATE
GOMP_MAP_FIRSTPRIVATE_INT
GOMP_MAP_FORCE_ALLOC
GOMP_MAP_FORCE_DEVICEPTR
GOMP_MAP_FORCE_FROM
GOMP_MAP_FORCE_PRESENT
GOMP_MAP_FORCE_TO
GOMP_MAP_FORCE_TOFROM
GOMP_MAP_FROM
GOMP_MAP_IF_PRESENT
GOMP_MAP_IMPLICIT
GOMP_MAP_POINTER
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
GOMP_MAP_RELEASE
GOMP_MAP_STRUCT
GOMP_MAP_STRUCT_UNORD
GOMP_MAP_TO
GOMP_MAP_TOFROM
GOMP_MAP_TO_PSET
GOMP_MAP_USE_DEVICE_PTR
GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT
GOMP_MAP_VARS_DATA
GOMP_MAP_VARS_ENTER_DATA
GOMP_MAP_VARS_OPENACC
GOMP_MAP_VARS_TARGET
GOMP_MAP_ZERO_LEN_ARRAY_SECTION
+/* Map entries containing expanded iterators will be flattened and merged into
+ HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are
+ any iterators found. ITERATOR_COUNT holds the iteration count of the
+ iterator that generates each map (0 if not generated from an iterator).
(To me it reads better with an 'and': '… each map (and 0 if …)'.)
+ HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any
+ merging occurs. */
I think we should somewhere in target.c document the format of iterators,
either here
or for the call, i.e. SIZE == MAX_SIZE to indicate the iterator, kind contains
the
size and hostpointer points to an array of triplets (addr,kind,size) for each
iterated item.
That's hidden in the code, but it helps to have it spelled out to make debugging
of general mapping issues easier.
* * *
@@ -1019,6 +1118,11 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
+ bool iterators_p = false;
+ size_t *iterator_count = NULL;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds, &iterator_count);
I think it makes sense to append ' /* OpenMP */' to the 'if
(short_mapkind)' line.
@@ -1896,14 +2000,17 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
if (pragma_kind & GOMP_MAP_VARS_TARGET)
{
+ size_t map_num = 0;
for (i = 0; i < mapnum; i++)
- {
- cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
- gomp_copy_host2dev (devicep, aq,
- (void *) (tgt->tgt_start + i * sizeof (void *)),
- (void *) &cur_node.tgt_offset, sizeof (void *),
- true, cbufp);
- }
+ if (!iterator_count || iterator_count[i] <= 1)
+ {
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (tgt->tgt_start + map_num * sizeof
(void *)),
+ (void *) &cur_node.tgt_offset, sizeof (void *),
+ true, cbufp);
+ map_num++;
+ }
}
First, I think it is more readable with '== 1' instead of '<= 1' as == 0 cannot
occur. Second, I think that requires a comment why > 1 is ignored here.
* * *
Thanks for the patch,
Tobias
