* [PATCH 0/3] openmp: Add support for iterators in OpenMP mapping clauses (C/C++)
@ 2024-05-24 19:57 Kwok Cheung Yeung
2024-05-24 19:59 ` [PATCH 1/3] openmp: Refactor handling of iterators Kwok Cheung Yeung
` (2 more replies)
0 siblings, 3 replies; 4+ messages in thread
From: Kwok Cheung Yeung @ 2024-05-24 19:57 UTC (permalink / raw)
To: Jakub Jelinek, Tobias Burnus, gcc-patches
This series of patches adds support for OpenMP iterators in the 'map'
clause of the 'target' construct (and it's derivatives such as 'target
enter data'), and the 'to' and 'from' constructs of the 'target update'
construct, currently for C and C++ only.
The approach in this patch differs from Tobias' WFC patch
(https://gcc.gnu.org/pipermail/gcc-patches/2021-December/586237.html) in
that it does not rely on generating a callback function - instead,
during Gimplification it generates loop(s) to evaluate every iteration
of the iterator expression, and the results (i.e. addresses, as the
expression should be an lvalue) are placed into a new array. This array
is then used as the 'hostaddrs' entry for that particular map. Libgomp
detects this (the corresponding size entry is set to SIZE_MAX, which
shouldn't normally occur) and inserts the contents of the array into the
map information before continuing on as normal.
Caveats:
- In section 2.21.7.1 of the OpenMP 5.1 standard, it states that 'If an
expression that is used to form a list item in a map clause contains an
iterator identifier, the list item instances that would result from
different values of the iterator must not have the same containing array
and must not have base pointers that share original storage' - this is
currently not enforced (it would prohibit something like map
(iterator(i=0:10), to: x[i]) while x is an int[]). As the expression in
the iterator is more-or-less unbound, it would be very difficult to
determine this at compile time. At runtime in libgomp, I suppose we
could check every iterator-derived mapping to ensure that they all
access unique entries in mem_map?
- The clause finishing currently generates spurious firstprivate maps -
the patch currently just ignores them when in iterator clauses, but is
there a better way of doing this?
- Clause reordering does not work too well with iterators. I believe the
current approach to reordering on trunk is a bit buggy in the first
place, so I just added enough to get the clauses through the pass
without ICEing.
The GCC gomp tests and all the libgomp tests have been run without
regressions on an x86-64 host with NVPTX offloading. Testing on AMD GCN
to follow.
Kwok
^ permalink raw reply [flat|nested] 4+ messages in thread
* [PATCH 1/3] openmp: Refactor handling of iterators
2024-05-24 19:57 [PATCH 0/3] openmp: Add support for iterators in OpenMP mapping clauses (C/C++) Kwok Cheung Yeung
@ 2024-05-24 19:59 ` Kwok Cheung Yeung
2024-05-24 20:01 ` [PATCH 2/3] openmp: Add support for iterators in map clauses (C/C++) Kwok Cheung Yeung
2024-05-24 20:02 ` [PATCH 3/3] openmp: Add support for iterators in to/from " Kwok Cheung Yeung
2 siblings, 0 replies; 4+ messages in thread
From: Kwok Cheung Yeung @ 2024-05-24 19:59 UTC (permalink / raw)
To: Jakub Jelinek, Tobias Burnus, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 361 bytes --]
This patch factors out the code to calculate the number of iterations
required and to generate the iteration loop into separate functions from
gimplify_omp_depend for reuse later.
I have also replaced the 'TREE_CODE (*tp) == TREE_LIST && ...' checks
used for detecting an iterator clause with a macro OMP_ITERATOR_DECL_P,
as it needs to be done frequently.
[-- Attachment #2: 0001-openmp-Refactor-handling-of-iterators.patch --]
[-- Type: text/plain, Size: 21693 bytes --]
From 0439fce03c2b5fb2802eaf65831e28f548ca074b Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcyeung@baylibre.com>
Date: Tue, 12 Mar 2024 20:51:38 +0000
Subject: [PATCH 1/3] openmp: Refactor handling of iterators
Move code to calculate the iteration size and to generate the iterator
expansion loop into separate functions.
Use OMP_ITERATOR_DECL_P to check for iterators in clause declarations.
2024-05-24 Kwok Cheung Yeung <kcyeung@baylibre.com>
gcc/c-family/
* c-omp.cc (c_finish_omp_depobj): Use OMP_ITERATOR_DECL_P.
gcc/c/
* c-typeck.cc (handle_omp_array_sections): Use OMP_ITERATOR_DECL_P.
(c_finish_omp_clauses): Likewise.
gcc/cp/
* pt.cc (tsubst_omp_clause_decl): Use OMP_ITERATOR_DECL_P.
* semantics.cc (handle_omp_array_sections): Likewise.
(finish_omp_clauses): Likewise.
gcc/
* gimplify.cc (gimplify_omp_affinity): Use OMP_ITERATOR_DECL_P.
(compute_iterator_count): New.
(build_iterator_loop): New.
(gimplify_omp_depend): Use OMP_ITERATOR_DECL_P, compute_iterator_count
and build_iterator_loop.
* tree-inline.cc (copy_tree_body_r): Use OMP_ITERATOR_DECL_P.
* tree-pretty-print.cc (dump_omp_clause): Likewise.
* tree.h (OMP_ITERATOR_DECL_P): New macro.
---
gcc/c-family/c-omp.cc | 4 +-
gcc/c/c-typeck.cc | 13 +-
gcc/cp/pt.cc | 4 +-
gcc/cp/semantics.cc | 8 +-
gcc/gimplify.cc | 326 +++++++++++++++++++--------------------
gcc/tree-inline.cc | 5 +-
gcc/tree-pretty-print.cc | 8 +-
gcc/tree.h | 6 +
8 files changed, 175 insertions(+), 199 deletions(-)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index c0e02aa422f..b56e49da62c 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -744,9 +744,7 @@ c_finish_omp_depobj (location_t loc, tree depobj,
kind = OMP_CLAUSE_DEPEND_KIND (clause);
t = OMP_CLAUSE_DECL (clause);
gcc_assert (t);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
error_at (OMP_CLAUSE_LOCATION (clause),
"%<iterator%> modifier may not be specified on "
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index 7ecca9f58c6..b0fe80cf224 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -14218,9 +14218,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
tree *tp = &OMP_CLAUSE_DECL (c);
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY)
- && TREE_CODE (*tp) == TREE_LIST
- && TREE_PURPOSE (*tp)
- && TREE_CODE (TREE_PURPOSE (*tp)) == TREE_VEC)
+ && OMP_ITERATOR_DECL_P (*tp))
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
@@ -15409,9 +15407,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_AFFINITY:
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_PURPOSE (t) != last_iterators)
last_iterators_remove
@@ -15511,10 +15507,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
}
}
- if (TREE_CODE (OMP_CLAUSE_DECL (c)) == TREE_LIST
- && TREE_PURPOSE (OMP_CLAUSE_DECL (c))
- && (TREE_CODE (TREE_PURPOSE (OMP_CLAUSE_DECL (c)))
- == TREE_VEC))
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
TREE_VALUE (OMP_CLAUSE_DECL (c)) = t;
else
OMP_CLAUSE_DECL (c) = t;
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index e77c48e463e..26db4f6e0cf 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17520,9 +17520,7 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain,
return decl;
/* Handle OpenMP iterators. */
- if (TREE_CODE (decl) == TREE_LIST
- && TREE_PURPOSE (decl)
- && TREE_CODE (TREE_PURPOSE (decl)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (decl))
{
tree ret;
if (iterator_cache[0] == TREE_PURPOSE (decl))
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index f90c304a65b..a48b3d2fcc5 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -5884,9 +5884,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort)
tree *tp = &OMP_CLAUSE_DECL (c);
if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY)
- && TREE_CODE (*tp) == TREE_LIST
- && TREE_PURPOSE (*tp)
- && TREE_CODE (TREE_PURPOSE (*tp)) == TREE_VEC)
+ && OMP_ITERATOR_DECL_P (*tp))
tp = &TREE_VALUE (*tp);
tree first = handle_omp_array_sections_1 (c, *tp, types,
maybe_zero_len, first_non_one,
@@ -8191,9 +8189,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_AFFINITY:
t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_PURPOSE (t) != last_iterators)
last_iterators_remove
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index b0ed58ed0f9..cb7358640f0 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -8640,9 +8640,7 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY)
{
tree t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_VALUE (t) == null_pointer_node)
continue;
@@ -8747,6 +8745,159 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p)
return;
}
+/* Returns a tree expression containing the total iteration count of the
+ iterator clause decl T. */
+
+static tree
+compute_iterator_count (tree t, gimple_seq *pre_p)
+{
+ tree tcnt = size_one_node;
+ for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+ {
+ if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR
+ || gimplify_expr (&TREE_VEC_ELT (it, 2), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR
+ || gimplify_expr (&TREE_VEC_ELT (it, 3), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR
+ || (gimplify_expr (&TREE_VEC_ELT (it, 4), pre_p, NULL,
+ is_gimple_val, fb_rvalue) == GS_ERROR))
+ return NULL_TREE;
+ tree var = TREE_VEC_ELT (it, 0);
+ tree begin = TREE_VEC_ELT (it, 1);
+ tree end = TREE_VEC_ELT (it, 2);
+ tree step = TREE_VEC_ELT (it, 3);
+ tree orig_step = TREE_VEC_ELT (it, 4);
+ tree type = TREE_TYPE (var);
+ tree stype = TREE_TYPE (step);
+ location_t loc = DECL_SOURCE_LOCATION (var);
+ tree endmbegin;
+ /* Compute count for this iterator as
+ orig_step > 0
+ ? (begin < end ? (end - begin + (step - 1)) / step : 0)
+ : (begin > end ? (end - begin + (step + 1)) / step : 0)
+ and compute product of those for the entire clause. */
+ if (POINTER_TYPE_P (type))
+ endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, stype, end, begin);
+ else
+ endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, end, begin);
+ tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step,
+ build_int_cst (stype, 1));
+ tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step,
+ build_int_cst (stype, 1));
+ tree pos = fold_build2_loc (loc, PLUS_EXPR, stype,
+ unshare_expr (endmbegin), stepm1);
+ pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, pos, step);
+ tree neg = fold_build2_loc (loc, PLUS_EXPR, stype, endmbegin, stepp1);
+ if (TYPE_UNSIGNED (stype))
+ {
+ neg = fold_build1_loc (loc, NEGATE_EXPR, stype, neg);
+ step = fold_build1_loc (loc, NEGATE_EXPR, stype, step);
+ }
+ neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, neg, step);
+ step = NULL_TREE;
+ tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, begin, end);
+ pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos,
+ build_int_cst (stype, 0));
+ cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, end, begin);
+ neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg,
+ build_int_cst (stype, 0));
+ tree osteptype = TREE_TYPE (orig_step);
+ cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, orig_step,
+ build_int_cst (osteptype, 0));
+ tree cnt = fold_build3_loc (loc, COND_EXPR, stype, cond, pos, neg);
+ cnt = fold_convert_loc (loc, sizetype, cnt);
+ if (gimplify_expr (&cnt, pre_p, NULL, is_gimple_val,
+ fb_rvalue) == GS_ERROR)
+ return NULL_TREE;
+ tcnt = size_binop_loc (loc, MULT_EXPR, tcnt, cnt);
+ }
+ if (gimplify_expr (&tcnt, pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
+ return NULL_TREE;
+
+ return tcnt;
+}
+
+/* Build loops iterating over the space defined by the iterators in clause C.
+ Returns a pointer to the BIND_EXPR_BODY in the innermost loop body.
+ LAST_BIND is set to point to the BIND_EXPR containing the whole loop. */
+
+static tree *
+build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind)
+{
+ tree t = OMP_CLAUSE_DECL (c);
+ gcc_assert (OMP_ITERATOR_DECL_P (t));
+
+ if (*last_bind)
+ gimplify_and_add (*last_bind, pre_p);
+ tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
+ *last_bind = build3 (BIND_EXPR, void_type_node,
+ BLOCK_VARS (block), NULL, block);
+ TREE_SIDE_EFFECTS (*last_bind) = 1;
+ SET_EXPR_LOCATION (*last_bind, OMP_CLAUSE_LOCATION (c));
+ tree *p = &BIND_EXPR_BODY (*last_bind);
+ for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ tree begin = TREE_VEC_ELT (it, 1);
+ tree end = TREE_VEC_ELT (it, 2);
+ tree step = TREE_VEC_ELT (it, 3);
+ tree orig_step = TREE_VEC_ELT (it, 4);
+ tree type = TREE_TYPE (var);
+ location_t loc = DECL_SOURCE_LOCATION (var);
+ /* Emit:
+ var = begin;
+ goto cond_label;
+ beg_label:
+ ...
+ var = var + step;
+ cond_label:
+ if (orig_step > 0) {
+ if (var < end) goto beg_label;
+ } else {
+ if (var > end) goto beg_label;
+ }
+ for each iterator, with inner iterators added to
+ the ... above. */
+ tree beg_label = create_artificial_label (loc);
+ tree cond_label = NULL_TREE;
+ tree tem = build2_loc (loc, MODIFY_EXPR, void_type_node, var, begin);
+ append_to_statement_list_force (tem, p);
+ tem = build_and_jump (&cond_label);
+ append_to_statement_list_force (tem, p);
+ tem = build1 (LABEL_EXPR, void_type_node, beg_label);
+ append_to_statement_list (tem, p);
+ tree bind = build3 (BIND_EXPR, void_type_node, NULL_TREE,
+ NULL_TREE, NULL_TREE);
+ TREE_SIDE_EFFECTS (bind) = 1;
+ SET_EXPR_LOCATION (bind, loc);
+ append_to_statement_list_force (bind, p);
+ if (POINTER_TYPE_P (type))
+ tem = build2_loc (loc, POINTER_PLUS_EXPR, type,
+ var, fold_convert_loc (loc, sizetype, step));
+ else
+ tem = build2_loc (loc, PLUS_EXPR, type, var, step);
+ tem = build2_loc (loc, MODIFY_EXPR, void_type_node, var, tem);
+ append_to_statement_list_force (tem, p);
+ tem = build1 (LABEL_EXPR, void_type_node, cond_label);
+ append_to_statement_list (tem, p);
+ tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, var, end);
+ tree pos = fold_build3_loc (loc, COND_EXPR, void_type_node, cond,
+ build_and_jump (&beg_label), void_node);
+ cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, var, end);
+ tree neg = fold_build3_loc (loc, COND_EXPR, void_type_node, cond,
+ build_and_jump (&beg_label), void_node);
+ tree osteptype = TREE_TYPE (orig_step);
+ cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, orig_step,
+ build_int_cst (osteptype, 0));
+ tem = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, pos, neg);
+ append_to_statement_list_force (tem, p);
+ p = &BIND_EXPR_BODY (bind);
+ }
+
+ return p;
+}
+
/* If *LIST_P contains any OpenMP depend clauses with iterators,
lower all the depend clauses by populating corresponding depend
array. Returns 0 if there are no such depend clauses, or
@@ -8791,89 +8942,12 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
tree t = OMP_CLAUSE_DECL (c);
if (first_loc == UNKNOWN_LOCATION)
first_loc = OMP_CLAUSE_LOCATION (c);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_PURPOSE (t) != last_iter)
{
- tree tcnt = size_one_node;
- for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
- {
- if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL,
- is_gimple_val, fb_rvalue) == GS_ERROR
- || gimplify_expr (&TREE_VEC_ELT (it, 2), pre_p, NULL,
- is_gimple_val, fb_rvalue) == GS_ERROR
- || gimplify_expr (&TREE_VEC_ELT (it, 3), pre_p, NULL,
- is_gimple_val, fb_rvalue) == GS_ERROR
- || (gimplify_expr (&TREE_VEC_ELT (it, 4), pre_p, NULL,
- is_gimple_val, fb_rvalue)
- == GS_ERROR))
- return 2;
- tree var = TREE_VEC_ELT (it, 0);
- tree begin = TREE_VEC_ELT (it, 1);
- tree end = TREE_VEC_ELT (it, 2);
- tree step = TREE_VEC_ELT (it, 3);
- tree orig_step = TREE_VEC_ELT (it, 4);
- tree type = TREE_TYPE (var);
- tree stype = TREE_TYPE (step);
- location_t loc = DECL_SOURCE_LOCATION (var);
- tree endmbegin;
- /* Compute count for this iterator as
- orig_step > 0
- ? (begin < end ? (end - begin + (step - 1)) / step : 0)
- : (begin > end ? (end - begin + (step + 1)) / step : 0)
- and compute product of those for the entire depend
- clause. */
- if (POINTER_TYPE_P (type))
- endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR,
- stype, end, begin);
- else
- endmbegin = fold_build2_loc (loc, MINUS_EXPR, type,
- end, begin);
- tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype,
- step,
- build_int_cst (stype, 1));
- tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step,
- build_int_cst (stype, 1));
- tree pos = fold_build2_loc (loc, PLUS_EXPR, stype,
- unshare_expr (endmbegin),
- stepm1);
- pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype,
- pos, step);
- tree neg = fold_build2_loc (loc, PLUS_EXPR, stype,
- endmbegin, stepp1);
- if (TYPE_UNSIGNED (stype))
- {
- neg = fold_build1_loc (loc, NEGATE_EXPR, stype, neg);
- step = fold_build1_loc (loc, NEGATE_EXPR, stype, step);
- }
- neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype,
- neg, step);
- step = NULL_TREE;
- tree cond = fold_build2_loc (loc, LT_EXPR,
- boolean_type_node,
- begin, end);
- pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos,
- build_int_cst (stype, 0));
- cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node,
- end, begin);
- neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg,
- build_int_cst (stype, 0));
- tree osteptype = TREE_TYPE (orig_step);
- cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node,
- orig_step,
- build_int_cst (osteptype, 0));
- tree cnt = fold_build3_loc (loc, COND_EXPR, stype,
- cond, pos, neg);
- cnt = fold_convert_loc (loc, sizetype, cnt);
- if (gimplify_expr (&cnt, pre_p, NULL, is_gimple_val,
- fb_rvalue) == GS_ERROR)
- return 2;
- tcnt = size_binop_loc (loc, MULT_EXPR, tcnt, cnt);
- }
- if (gimplify_expr (&tcnt, pre_p, NULL, is_gimple_val,
- fb_rvalue) == GS_ERROR)
+ tree tcnt = compute_iterator_count (t, pre_p);
+ if (!tcnt)
return 2;
last_iter = TREE_PURPOSE (t);
last_count = tcnt;
@@ -9027,92 +9101,10 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
gcc_unreachable ();
}
tree t = OMP_CLAUSE_DECL (c);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_PURPOSE (t) != last_iter)
- {
- if (last_bind)
- gimplify_and_add (last_bind, pre_p);
- tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
- last_bind = build3 (BIND_EXPR, void_type_node,
- BLOCK_VARS (block), NULL, block);
- TREE_SIDE_EFFECTS (last_bind) = 1;
- SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
- tree *p = &BIND_EXPR_BODY (last_bind);
- for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
- {
- tree var = TREE_VEC_ELT (it, 0);
- tree begin = TREE_VEC_ELT (it, 1);
- tree end = TREE_VEC_ELT (it, 2);
- tree step = TREE_VEC_ELT (it, 3);
- tree orig_step = TREE_VEC_ELT (it, 4);
- tree type = TREE_TYPE (var);
- location_t loc = DECL_SOURCE_LOCATION (var);
- /* Emit:
- var = begin;
- goto cond_label;
- beg_label:
- ...
- var = var + step;
- cond_label:
- if (orig_step > 0) {
- if (var < end) goto beg_label;
- } else {
- if (var > end) goto beg_label;
- }
- for each iterator, with inner iterators added to
- the ... above. */
- tree beg_label = create_artificial_label (loc);
- tree cond_label = NULL_TREE;
- tem = build2_loc (loc, MODIFY_EXPR, void_type_node,
- var, begin);
- append_to_statement_list_force (tem, p);
- tem = build_and_jump (&cond_label);
- append_to_statement_list_force (tem, p);
- tem = build1 (LABEL_EXPR, void_type_node, beg_label);
- append_to_statement_list (tem, p);
- tree bind = build3 (BIND_EXPR, void_type_node, NULL_TREE,
- NULL_TREE, NULL_TREE);
- TREE_SIDE_EFFECTS (bind) = 1;
- SET_EXPR_LOCATION (bind, loc);
- append_to_statement_list_force (bind, p);
- if (POINTER_TYPE_P (type))
- tem = build2_loc (loc, POINTER_PLUS_EXPR, type,
- var, fold_convert_loc (loc, sizetype,
- step));
- else
- tem = build2_loc (loc, PLUS_EXPR, type, var, step);
- tem = build2_loc (loc, MODIFY_EXPR, void_type_node,
- var, tem);
- append_to_statement_list_force (tem, p);
- tem = build1 (LABEL_EXPR, void_type_node, cond_label);
- append_to_statement_list (tem, p);
- tree cond = fold_build2_loc (loc, LT_EXPR,
- boolean_type_node,
- var, end);
- tree pos
- = fold_build3_loc (loc, COND_EXPR, void_type_node,
- cond, build_and_jump (&beg_label),
- void_node);
- cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node,
- var, end);
- tree neg
- = fold_build3_loc (loc, COND_EXPR, void_type_node,
- cond, build_and_jump (&beg_label),
- void_node);
- tree osteptype = TREE_TYPE (orig_step);
- cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node,
- orig_step,
- build_int_cst (osteptype, 0));
- tem = fold_build3_loc (loc, COND_EXPR, void_type_node,
- cond, pos, neg);
- append_to_statement_list_force (tem, p);
- p = &BIND_EXPR_BODY (bind);
- }
- last_body = p;
- }
+ last_body = build_iterator_loop (c, pre_p, &last_bind);
last_iter = TREE_PURPOSE (t);
if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR)
{
diff --git a/gcc/tree-inline.cc b/gcc/tree-inline.cc
index f31a34ac410..05dea9473a0 100644
--- a/gcc/tree-inline.cc
+++ b/gcc/tree-inline.cc
@@ -1453,10 +1453,7 @@ copy_tree_body_r (tree *tp, int *walk_subtrees, void *data)
|| OMP_CLAUSE_CODE (*tp) == OMP_CLAUSE_DEPEND))
{
tree t = OMP_CLAUSE_DECL (*tp);
- if (t
- && TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (t && OMP_ITERATOR_DECL_P (t))
{
*walk_subtrees = 0;
OMP_CLAUSE_DECL (*tp) = copy_node (t);
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index f9ad8562078..011f44bfd3d 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -805,9 +805,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
pp_string (pp, "affinity(");
{
tree t = OMP_CLAUSE_DECL (clause);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
dump_omp_iterators (pp, TREE_PURPOSE (t), spc, flags);
pp_colon (pp);
@@ -847,9 +845,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
}
{
tree t = OMP_CLAUSE_DECL (clause);
- if (TREE_CODE (t) == TREE_LIST
- && TREE_PURPOSE (t)
- && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC)
+ if (OMP_ITERATOR_DECL_P (t))
{
dump_omp_iterators (pp, TREE_PURPOSE (t), spc, flags);
pp_colon (pp);
diff --git a/gcc/tree.h b/gcc/tree.h
index ee2aae332a4..e8568a69f95 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -2147,6 +2147,12 @@ class auto_suppress_location_wrappers
#define OMP_CLAUSE_OPERAND(NODE, I) \
OMP_CLAUSE_ELT_CHECK (NODE, I)
+/* True if the clause decl NODE contains an iterator. */
+#define OMP_ITERATOR_DECL_P(NODE) \
+ (TREE_CODE (NODE) == TREE_LIST \
+ && TREE_PURPOSE (NODE) \
+ && TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC)
+
/* In a BLOCK (scope) node:
Variables declared in the scope NODE. */
#define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars)
--
2.34.1
^ permalink raw reply [flat|nested] 4+ messages in thread
* [PATCH 2/3] openmp: Add support for iterators in map clauses (C/C++)
2024-05-24 19:57 [PATCH 0/3] openmp: Add support for iterators in OpenMP mapping clauses (C/C++) Kwok Cheung Yeung
2024-05-24 19:59 ` [PATCH 1/3] openmp: Refactor handling of iterators Kwok Cheung Yeung
@ 2024-05-24 20:01 ` Kwok Cheung Yeung
2024-05-24 20:02 ` [PATCH 3/3] openmp: Add support for iterators in to/from " Kwok Cheung Yeung
2 siblings, 0 replies; 4+ messages in thread
From: Kwok Cheung Yeung @ 2024-05-24 20:01 UTC (permalink / raw)
To: Jakub Jelinek, Tobias Burnus, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1889 bytes --]
This patch modifies the C and C++ parsers to accept an iterator as a map
type modifier, encoded in the same way as the depend and affinity
clauses. When finishing the clauses, clauses with iterators are treated
separately from ones without to avoid clashes (e.g. iterating over x[i]
will likely generate clauses to map x).
During gimplification, gimplify_omp_map_iterators is called during
scanning if a map clause encountered has any iterators. This scans all
the remaining clauses in one go, as iterators may be shared between
clauses. Later clauses with iterators are simply skipped over.
For each map clause with an iterator, gimplify_omp_map_iterators
generates a loop (or multiple loops, if the iterator is
multidimensional) to iterate over the iterator expression, storing the
result in a new array (constant-sized for now, we could dynamically
allocate the array for non-constant iteration bounds). The data array
stores the total number of iterations in the first element, then the
address generated by the iterator expression and the OMP_CLAUSE_SIZE
(since the iteration variables may occur within the size tree) for each
iteration. The clause is then rewritten to point to the new array. The
original clause decl is no longer directly relevant, but is kept around
for informational purposes and to help with clause sorting. The original
OMP_CLAUSE_SIZE is set to NULL_TREE.
When OMP lowering clauses with iterators, the data array holding the
expanded iterator info is allocated to a field in the omp_data, and the
size is set to SIZE_MAX to mark the entry as coming from an expanded
iterator.
Libgomp has a new function gomp_merge_iterator_maps which identifies
data coming from an iterator, and effectively creates new maps
on-the-fly from the iterator info array, inserting them into the list of
mappings at the point where iterator data occurred.
[-- Attachment #2: 0002-openmp-Add-support-for-iterators-in-map-clauses-C-C.patch --]
[-- Type: text/plain, Size: 43490 bytes --]
From b2e8ff46929d5a2781781486ec942b344056d78b Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcyeung@baylibre.com>
Date: Tue, 12 Mar 2024 22:51:06 +0000
Subject: [PATCH 2/3] 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.
2024-05-24 Kwok Cheung Yeung <kcyeung@baylibre.com>
gcc/c/
* c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier.
* c-typeck.cc (c_finish_omp_clauses): Call recursively on iterator
clauses.
gcc/cp/
* parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier.
* semantics.cc (finish_omp_clauses): Call recursively on iterator
clauses.
gcc/
* gimplify.cc (find_var_decl): New.
(check_iterator_var_usage): New.
(gimplify_omp_map_iterators): New.
(omp_group_iterator): New.
(omp_get_attachment): Replace OMP_CLAUSE_DECL with
OMP_ITERATOR_CLAUSE_DECL.
(omp_group_last): Keep decls with and without iterators in separate
groups.
(omp_index_mapping_groups_1): Replace OMP_CLAUSE_DECL with
OMP_ITERATOR_CLAUSE_DECL.
(omp_tsort_mapping_groups_1): Likewise.
(omp_resolve_clause_dependencies): Likewise. Prevent removal of
mapping if groups do not use the same iterators.
(omp_build_struct_sibling_lists): Replace OMP_CLAUSE_DECL with
OMP_ITERATOR_CLAUSE_DECL.
(gimplify_scan_omp_clauses): Call gimplify_omp_map_iterators once to
handle clauses with iterators, then skip subsequent iterator clauses.
* omp-low.cc (scan_sharing_clauses): Add field for iterator clauses.
(lower_omp_target): Add map entries for iterator clauses.
* tree-pretty-print.cc (dump_omp_map_iterators): New.
(dump_omp_clause): Call dump_omp_map_iterators for iterators in map
clauses.
* tree.h (OMP_ITERATOR_CLAUSE_DECL): New.
gcc/testsuite/
* c-c++-common/gomp/map-6.c (foo): Amend expected error message.
* c-c++-common/gomp/target-iterator-1.c: New.
* c-c++-common/gomp/target-iterator-2.c: New.
* c-c++-common/gomp/target-iterator-3.c: New.
libgomp/
* target.c (gomp_merge_iterator_maps): New.
(gomp_map_vars_internal): Call gomp_merge_iterator_maps. Free
allocated variables.
* testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New.
---
gcc/c/c-parser.cc | 60 ++++-
gcc/c/c-typeck.cc | 68 ++++++
gcc/cp/parser.cc | 64 ++++-
gcc/cp/semantics.cc | 65 ++++++
gcc/gimplify.cc | 220 +++++++++++++++++-
gcc/omp-low.cc | 52 ++++-
gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +-
.../c-c++-common/gomp/target-iterator-1.c | 23 ++
.../c-c++-common/gomp/target-iterator-2.c | 17 ++
.../c-c++-common/gomp/target-iterator-3.c | 20 ++
gcc/tree-pretty-print.cc | 24 +-
gcc/tree.h | 7 +
libgomp/target.c | 83 +++++++
.../target-map-iterators-1.c | 44 ++++
.../target-map-iterators-2.c | 42 ++++
.../target-map-iterators-3.c | 54 +++++
16 files changed, 823 insertions(+), 30 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 00f8bf4376e..2281148561c 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -18772,7 +18772,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list)
map ( [map-type-modifier[,] ...] map-kind: variable-list )
map-type-modifier:
- always | close */
+ always | close | present | iterator (iterators-definition) */
static tree
c_parser_omp_clause_map (c_parser *parser, tree list)
@@ -18787,15 +18787,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int pos = 1;
int map_kind_pos = 0;
- while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+ int iterator_length = 0;
+ for (;;)
{
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+ c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
+ if (tok->type != CPP_NAME)
+ break;
+
+ const char *p = IDENTIFIER_POINTER (tok->value);
+ c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN)
+ {
+ unsigned n = pos + 2;
+ if (c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN)
+ {
+ iterator_length = n - pos + 1;
+ pos = n;
+ next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
pos++;
}
@@ -18803,6 +18823,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int always_modifier = 0;
int close_modifier = 0;
int present_modifier = 0;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -18844,10 +18865,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
}
present_modifier++;
}
+ else if (strcmp ("iterator", p) == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ if (iterators)
+ {
+ c_parser_error (parser, "too many %<iterator%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ iterators = c_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ continue;
+ }
else
{
c_parser_error (parser, "%<map%> clause with map-type modifier other "
- "than %<always%>, %<close%> or %<present%>");
+ "than %<always%>, %<close%>, %<iterator%> "
+ "or %<present%>");
parens.skip_until_found_close (parser);
return list;
}
@@ -18896,8 +18931,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
true);
+ if (iterators)
+ {
+ tree block = pop_scope ();
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (iterators)
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+ }
parens.skip_until_found_close (parser);
return nl;
diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc
index b0fe80cf224..e29bf37a44c 100644
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -15526,6 +15526,74 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ tree iterators = TREE_PURPOSE (t);
+ if (c_omp_finish_iterators (iterators))
+ {
+ t = error_mark_node;
+ break;
+ }
+
+ /* Find the end of the group of clauses that use the same
+ iterator.*/
+ tree end_clause;
+ for (end_clause = c; end_clause;
+ end_clause = OMP_CLAUSE_CHAIN (end_clause))
+ {
+ tree nc = OMP_CLAUSE_CHAIN (end_clause);
+ /* Remove iterator temporarily. */
+ OMP_CLAUSE_DECL (end_clause) =
+ TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+ if (!nc
+ || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+ || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+ break;
+ }
+ tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+ /* Temporarily split off the group of clauses with the same
+ iterator. */
+ OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+ tree new_clauses = c_finish_omp_clauses (c, ort);
+
+ /* Replace the iterators and splice the new clauses in. */
+ tree *clause_p = &new_clauses;
+ while (*clause_p)
+ {
+ /* Skip unwanted clause types.
+ FIXME: Is this the right thing to do? */
+ bool skip = false;
+ if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ skip = false;
+ break;
+ default:
+ skip = true;
+ break;
+ }
+ if (skip)
+ *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+ else
+ {
+ OMP_CLAUSE_DECL (*clause_p)
+ = build_tree_list (iterators,
+ OMP_CLAUSE_DECL (*clause_p));
+
+ clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+ }
+ }
+ *clause_p = next_clause;
+ *pc = new_clauses;
+ pc = clause_p;
+ continue;
+ }
+
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
grp_start_p = pc;
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 476ddc0d63a..6dc67851f96 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41672,16 +41672,35 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
int pos = 1;
int map_kind_pos = 0;
- while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
- || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
+ int iterator_length = 0;
+ for (;;)
{
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON)
+ cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos);
+ if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE))
+ break;
+
+ cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1);
+ if (tok->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
+ && next_tok->type == CPP_OPEN_PAREN)
+ {
+ size_t n = cp_parser_skip_balanced_tokens (parser, 2);
+ if (cp_lexer_peek_nth_token (parser->lexer, n - 1)->type
+ == CPP_CLOSE_PAREN)
+ {
+ iterator_length = n - pos;
+ pos = n - 1;
+ next_tok = cp_lexer_peek_nth_token (parser->lexer, n);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
pos++;
}
@@ -41689,6 +41708,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
bool always_modifier = false;
bool close_modifier = false;
bool present_modifier = false;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -41738,10 +41758,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
}
present_modifier = true;
}
+ else if (strcmp ("iterator", p) == 0
+ && cp_lexer_peek_nth_token (parser->lexer, 2)->type
+ == CPP_OPEN_PAREN)
+ {
+ if (iterators)
+ {
+ cp_parser_error (parser, "too many %<iterator%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ begin_scope (sk_omp, NULL);
+ iterators = cp_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ continue;
+ }
else
{
cp_parser_error (parser, "%<map%> clause with map-type modifier other"
- " than %<always%>, %<close%> or %<present%>");
+ " than %<always%>, %<close%>, %<iterator%>"
+ " or %<present%>");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -41805,8 +41844,21 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
NULL, true);
finish_scope ();
+ if (iterators)
+ {
+ tree block = poplevel (1, 1, 0);
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (iterators)
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+ }
return nlist;
}
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index a48b3d2fcc5..aed0b7c024f 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8362,6 +8362,71 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ tree iterators = TREE_PURPOSE (t);
+ if (cp_omp_finish_iterators (iterators))
+ {
+ t = error_mark_node;
+ break;
+ }
+
+ /* Find the end of the group of clauses that use the same
+ iterator.*/
+ tree end_clause;
+ for (end_clause = c; end_clause;
+ end_clause = OMP_CLAUSE_CHAIN (end_clause))
+ {
+ tree nc = OMP_CLAUSE_CHAIN (end_clause);
+ /* Remove iterator temporarily. */
+ OMP_CLAUSE_DECL (end_clause) =
+ TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+ if (!nc
+ || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+ || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+ break;
+ }
+ tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+ /* Temporarily split off the group of clauses with the same
+ iterator. */
+ OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+ tree new_clauses = finish_omp_clauses (c, ort);
+
+ /* Replace the iterators and splice the new clauses in. */
+ tree *clause_p = &new_clauses;
+ while (*clause_p)
+ {
+ OMP_CLAUSE_DECL (*clause_p)
+ = build_tree_list (iterators,
+ OMP_CLAUSE_DECL (*clause_p));
+ /* Skip unwanted clause types.
+ FIXME: Is this the right thing to do? */
+ bool skip = false;
+ if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ skip = false;
+ break;
+ default:
+ skip = true;
+ break;
+ }
+ if (skip)
+ *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+ else
+ clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+ }
+ *clause_p = next_clause;
+ *pc = new_clauses;
+ pc = clause_p;
+ continue;
+ }
+
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
grp_start_p = pc;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index cb7358640f0..cadb2a3d96d 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9253,6 +9253,172 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
return 1;
}
+/* Callback for walk_tree to find a VAR_DECL for the given tree. */
+
+static tree
+find_var_decl (tree *tp, int *, void *data)
+{
+ tree t = *tp;
+
+ if (TREE_CODE (t) == VAR_DECL && t == (tree) data)
+ return t;
+
+ return NULL_TREE;
+}
+
+/* Check for clause decls in iterators that do not use all the iterator
+ variables. */
+
+static bool
+check_iterator_var_usage (tree c)
+{
+ tree decl = OMP_CLAUSE_DECL (c);
+ bool error = false;
+ gcc_assert (OMP_ITERATOR_DECL_P (decl));
+
+ for (tree it = TREE_PURPOSE (decl); it; it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ tree t = walk_tree (&TREE_VALUE (decl), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "iterator variable %qD not used in clause expression",
+ var);
+ error = true;
+ }
+ }
+ return !error;
+}
+
+static void
+gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
+{
+ tree last_iter = NULL_TREE;
+ tree last_bind = NULL_TREE;
+ tree last_count = NULL_TREE;
+ tree last_index = NULL_TREE;
+ tree *last_body = NULL;
+
+ while (tree c = *list_p)
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ {
+ list_p = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ tree t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ if (!check_iterator_var_usage (c))
+ {
+ *list_p = OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ if (TREE_PURPOSE (t) != last_iter)
+ {
+ tree tcnt = compute_iterator_count (t, pre_p);
+ if (!tcnt)
+ {
+ *list_p = OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+ last_iter = TREE_PURPOSE (t);
+ last_count = tcnt;
+
+ last_body = build_iterator_loop (c, pre_p, &last_bind);
+ last_index = create_tmp_var (sizetype);
+ SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
+
+ /* idx = -1; */
+ /* This should be initialized to before the individual elements,
+ as idx is pre-incremented in the loop body. */
+ gimple *g = gimple_build_assign (last_index, size_int (-1));
+ gimple_seq_add_stmt (pre_p, g);
+
+ /* IN LOOP BODY: */
+ /* idx += 2; */
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, last_index,
+ size_binop (PLUS_EXPR, last_index,
+ size_int (2)));
+ append_to_statement_list_force (tem, last_body);
+ }
+
+ /* Create array to hold expanded values. */
+ tree last_count_2 = size_binop (MULT_EXPR, last_count, size_int (2));
+ tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1));
+ tree elems = NULL_TREE;
+ if (TREE_CONSTANT (arr_length))
+ {
+ tree type = build_array_type (ptr_type_node,
+ build_index_type (arr_length));
+ elems = create_tmp_var_raw (type);
+ TREE_ADDRESSABLE (elems) = 1;
+ gimple_add_tmp_var (elems);
+ }
+ else
+ {
+ /* Handle dynamic sizes. */
+ sorry ("Dynamic iterator sizes not implemented yet.");
+ }
+
+ /* elems[0] = count; */
+ tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
+ NULL_TREE, NULL_TREE);
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, last_count);
+ gimplify_and_add (tem, pre_p);
+
+ /* IN LOOP BODY: */
+ /* elems[idx] = &<expr>; */
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems, last_index, NULL_TREE,
+ NULL_TREE);
+ tree rhs = build1 (ADDR_EXPR, ptr_type_node, TREE_VALUE (t));
+ tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, rhs);
+ append_to_statement_list_force (tem, last_body);
+
+ /* elems[idx+1] = OMP_CLAUSE_SIZE (c); */
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems,
+ size_binop (PLUS_EXPR, last_index, size_int (1)),
+ NULL_TREE, NULL_TREE);
+ tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, OMP_CLAUSE_SIZE (c));
+ append_to_statement_list_force (tem, last_body);
+
+ /* Replace iterator information. */
+ TREE_PURPOSE (t) = make_tree_vec (2);
+ TREE_VEC_ELT (TREE_PURPOSE (t), 0) = last_iter;
+ TREE_VEC_ELT (TREE_PURPOSE (t), 1) = elems;
+
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ }
+ else if (last_bind)
+ {
+ bool saved_into_ssa = gimplify_ctxp->into_ssa;
+ gimplify_ctxp->into_ssa = false;
+ gimplify_and_add (last_bind, pre_p);
+ gimplify_ctxp->into_ssa = saved_into_ssa;
+ last_bind = NULL_TREE;
+ }
+
+ list_p = &OMP_CLAUSE_CHAIN (c);
+ }
+
+ if (last_bind)
+ {
+ bool saved_into_ssa = gimplify_ctxp->into_ssa;
+ gimplify_ctxp->into_ssa = false;
+ gimplify_and_add (last_bind, pre_p);
+ gimplify_ctxp->into_ssa = saved_into_ssa;
+ }
+}
+
/* True if mapping node C maps, or unmaps, a (Fortran) array descriptor. */
static bool
@@ -9461,6 +9627,22 @@ omp_get_base_pointer (tree expr)
return NULL_TREE;
}
+/* Return the iterator for a mapping group, or NULL if there isn't one. */
+
+static tree
+omp_group_iterator (omp_mapping_group *grp)
+{
+ tree c = grp->grp_end;
+ if (!OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ return NULL_TREE;
+
+ tree iter = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+ if (TREE_VEC_LENGTH (iter) == 2)
+ iter = TREE_VEC_ELT (iter, 0);
+
+ return iter;
+}
+
/* An attach or detach operation depends directly on the address being
attached/detached. Return that address, or none if there are no
attachments/detachments. */
@@ -9515,7 +9697,7 @@ omp_get_attachment (omp_mapping_group *grp)
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
case GOMP_MAP_DETACH:
- return OMP_CLAUSE_DECL (node);
+ return OMP_ITERATOR_CLAUSE_DECL (node);
default:
internal_error ("unexpected mapping node");
@@ -9527,7 +9709,7 @@ omp_get_attachment (omp_mapping_group *grp)
node = OMP_CLAUSE_CHAIN (node);
if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
- return OMP_CLAUSE_DECL (node);
+ return OMP_ITERATOR_CLAUSE_DECL (node);
else
internal_error ("unexpected mapping node");
return error_mark_node;
@@ -9539,7 +9721,7 @@ omp_get_attachment (omp_mapping_group *grp)
return OMP_CLAUSE_DECL (*grp->grp_start);
if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
- return OMP_CLAUSE_DECL (*grp->grp_start);
+ return OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
else
internal_error ("unexpected mapping node");
return error_mark_node;
@@ -9593,7 +9775,9 @@ omp_group_last (tree *start_p)
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
- || omp_map_clause_descriptor_p (nc)))
+ || omp_map_clause_descriptor_p (nc))
+ && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+ == OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc)))
{
tree nc2 = OMP_CLAUSE_CHAIN (nc);
if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
@@ -9896,7 +10080,7 @@ omp_index_mapping_groups_1 (hash_map<tree_operand_hash_no_se,
node && j < chained;
node = OMP_CLAUSE_CHAIN (node), j++)
{
- tree decl = OMP_CLAUSE_DECL (node);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (node);
/* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF,
meaning node-hash lookups don't work. This is a workaround for
that, but ideally we should just create the INDIRECT_REF at
@@ -10083,7 +10267,7 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
}
}
- tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
while (decl)
{
@@ -10622,7 +10806,7 @@ omp_resolve_clause_dependencies (enum tree_code code,
FOR_EACH_VEC_ELT (*groups, i, grp)
{
tree grp_end = grp->grp_end;
- tree decl = OMP_CLAUSE_DECL (grp_end);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (grp_end);
gcc_assert (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP);
@@ -10809,7 +10993,9 @@ omp_resolve_clause_dependencies (enum tree_code code,
{
omp_mapping_group *struct_group;
if (omp_mapped_by_containing_struct (grpmap, decl, &struct_group)
- && *grp->grp_start == grp_end)
+ && *grp->grp_start == grp_end
+ && omp_group_iterator (grp)
+ == omp_group_iterator (struct_group))
{
omp_check_mapping_compatibility (OMP_CLAUSE_LOCATION (grp_end),
struct_group, grp);
@@ -11729,7 +11915,7 @@ omp_build_struct_sibling_lists (enum tree_code code,
FOR_EACH_VEC_ELT (*groups, i, grp)
{
tree c = grp->grp_end;
- tree decl = OMP_CLAUSE_DECL (c);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (c);
tree grp_end = grp->grp_end;
auto_vec<omp_addr_token *> addr_tokens;
tree sentinel = OMP_CLAUSE_CHAIN (grp_end);
@@ -12014,6 +12200,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree c;
tree *orig_list_p = list_p;
int handled_depend_iterators = -1;
+ bool handled_map_iterators = false;
int nowait = -1;
ctx = new_omp_context (region_type);
@@ -12395,6 +12582,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
decl = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ if (!handled_map_iterators)
+ {
+ gimplify_omp_map_iterators (list_p, pre_p);
+ handled_map_iterators = true;
+ continue;
+ }
+ /* Skip declarations with iterators. */
+ break;
+ }
+
if (error_operand_p (decl))
{
remove = true;
@@ -14076,6 +14275,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
if (remove)
break;
+ /* No further gimplfication required for clauses with iterators. */
+ if (OMP_ITERATOR_DECL_P (decl))
+ break;
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
{
/* Sanity check: attach/detach map kinds use the size as a bias,
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 4d003f42098..d0f0c5d884a 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1516,6 +1516,22 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_MAP:
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ /* FIXME: Is this the right way to handle these? */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ break;
+ 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) OMP_ITERATOR_CLAUSE_DECL (c),
+ (splay_tree_value) field);
+ break;
+ }
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
decl = OMP_CLAUSE_DECL (c);
@@ -12734,11 +12750,21 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
#endif
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ {
+ /* Skip firstprivate pointers/references.
+ FIXME: Is this the right thing to do? */
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ map_cnt++;
+ continue;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
oacc_firstprivate:
- var = OMP_CLAUSE_DECL (c);
+ var = OMP_ITERATOR_CLAUSE_DECL (c);
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
@@ -13019,12 +13045,31 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_FROM:
oacc_firstprivate_map:
nc = c;
- ovar = OMP_CLAUSE_DECL (c);
+ ovar = OMP_ITERATOR_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
break;
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ tree x = build_sender_ref (OMP_ITERATOR_CLAUSE_DECL (c), ctx);
+ tree iterator = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+ tree array = TREE_VEC_ELT (iterator, 1);
+ tree array_addr = build1 (ADDR_EXPR, ptr_type_node, array);
+ gimplify_assign (x, array_addr, &ilist);
+ purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, size_int (SIZE_MAX));
+
+ unsigned HOST_WIDE_INT tkind = OMP_CLAUSE_MAP_KIND (c);
+ gcc_checking_assert (tkind
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
+ gcc_checking_assert (tkind
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+ build_int_cstu (tkind_type, tkind));
+ break;
+ }
if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13904,6 +13949,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
+ /* FIXME: Handle firstprivate mappings for iterators properly. */
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ break;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
{
diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c
index 014ed35ab41..13e3b58cc92 100644
--- a/gcc/testsuite/c-c++-common/gomp/map-6.c
+++ b/gcc/testsuite/c-c++-common/gomp/map-6.c
@@ -13,19 +13,19 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c
new file mode 100644
index 00000000000..7d6c8dc6255
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, int **y)
+{
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* { dg-message "unsupported map expression" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i]) /* { dg-message "unsupported map expression" } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c
new file mode 100644
index 00000000000..39efda74fdc
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target map(iterator(i=0:10), to: x) /* { dg-error "iterator variable .i. not used in clause expression" }*/
+ ;
+
+ #pragma omp target map(iterator(i=0:10, j=0:20), to: x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/
+ ;
+
+ #pragma omp target map(iterator(i=0:10, j=0:20, k=0:30), to: x[i], y[j], z[k])
+ /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */
+ /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */
+ /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c
new file mode 100644
index 00000000000..22becdda559
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+ #pragma omp target map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+ map(iterator(i=0:DIM1), from: z[i][:DIM2])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:from:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:to:" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:attach:" 4 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index 011f44bfd3d..feb9c93a809 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -451,6 +451,19 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags)
pp_right_paren (pp);
}
+static void
+dump_omp_map_iterators (pretty_printer *pp, tree iter, int spc,
+ dump_flags_t flags)
+{
+ if (TREE_VEC_LENGTH (iter) == 6)
+ dump_omp_iterators (pp, iter, spc, flags);
+ else
+ {
+ dump_omp_iterators (pp, TREE_VEC_ELT (iter, 0), spc, flags);
+ pp_string (pp, ":iterator_array=");
+ dump_generic_node (pp, TREE_VEC_ELT (iter, 1), spc, flags, false);
+ }
+}
/* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN.
@@ -461,6 +474,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
{
const char *name;
const char *modifier = NULL;
+ tree decl = NULL_TREE;
switch (OMP_CLAUSE_CODE (clause))
{
case OMP_CLAUSE_PRIVATE:
@@ -911,6 +925,13 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
pp_string (pp, "map(");
if (OMP_CLAUSE_MAP_READONLY (clause))
pp_string (pp, "readonly,");
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_ALLOC:
@@ -1025,8 +1046,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
gcc_unreachable ();
}
pp_colon (pp);
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ dump_generic_node (pp, decl, spc, flags, false);
print_clause_size:
if (OMP_CLAUSE_SIZE (clause))
{
diff --git a/gcc/tree.h b/gcc/tree.h
index e8568a69f95..e01408c2e43 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -2153,6 +2153,13 @@ class auto_suppress_location_wrappers
&& TREE_PURPOSE (NODE) \
&& TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC)
+/* Return the iterator expression if NODE contains an iterator.
+ Return the clause decl if NODE does not. */
+#define OMP_ITERATOR_CLAUSE_DECL(NODE) \
+ (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (NODE)) \
+ ? TREE_VALUE (OMP_CLAUSE_DECL (NODE)) \
+ : OMP_CLAUSE_DECL (NODE))
+
/* In a BLOCK (scope) node:
Variables declared in the scope NODE. */
#define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars)
diff --git a/libgomp/target.c b/libgomp/target.c
index 5ec19ae489e..07ba840b495 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -972,6 +972,77 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
}
}
+/* 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. HOSTADDRS, SIZES and KINDS must be freed afterwards
+ if any merging occurs. */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+ void **kinds)
+{
+ bool iterator_p = false;
+ size_t map_count = 0;
+ unsigned short **skinds = (unsigned short **) kinds;
+
+ for (size_t i = 0; i < *mapnum; i++)
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ map_count += iterator_array[0];
+ iterator_p = true;
+ }
+ else
+ map_count++;
+
+ if (!iterator_p)
+ return false;
+
+ gomp_debug (1,
+ "Expanding iterator maps - number of map entries: %ld -> %ld\n",
+ *mapnum, map_count);
+ void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *));
+ size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+ unsigned short *new_kinds
+ = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short));
+ size_t new_idx = 0;
+
+ for (size_t i = 0; i < *mapnum; i++)
+ {
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ size_t count = iterator_array[0];
+ for (int j = 1; j < count * 2 + 1; j += 2)
+ {
+ new_hostaddrs[new_idx] = (void *) iterator_array[j];
+ new_sizes[new_idx] = iterator_array[j+1];
+ new_kinds[new_idx] = (*skinds)[i];
+ gomp_debug (1,
+ "Expanding map %ld: "
+ "hostaddrs[%ld] = %p, sizes[%ld] = %ld\n",
+ i, new_idx, new_hostaddrs[new_idx],
+ new_idx, new_sizes[new_idx]);
+ new_idx++;
+ }
+ }
+ else
+ {
+ new_hostaddrs[new_idx] = (*hostaddrs)[i];
+ new_sizes[new_idx] = (*sizes)[i];
+ new_kinds[new_idx] = (*skinds)[i];
+ new_idx++;
+ }
+ }
+
+ *mapnum = map_count;
+ *hostaddrs = new_hostaddrs;
+ *sizes = new_sizes;
+ *kinds = new_kinds;
+
+ return true;
+}
+
static inline __attribute__((always_inline)) struct target_mem_desc *
gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
@@ -988,6 +1059,10 @@ 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;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds);
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
@@ -1873,6 +1948,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ }
+
return tgt;
}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
new file mode 100644
index 00000000000..900a0ba2d64
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c
@@ -0,0 +1,44 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[])
+{
+ int expected = 0;
+
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ return expected;
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y;
+
+ int expected = mkarray (x);
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) map(from: y)
+ {
+ y = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+ }
+
+ return y - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
new file mode 100644
index 00000000000..bad0f7f17b8
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays from target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+ for (int i = 0; i < DIM1; i++)
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y, expected;
+
+ mkarray (x);
+
+ #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \
+ map(from: expected)
+ {
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = (i+1) * (j+1);
+ expected += x[i][j];
+ }
+ }
+
+ y = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+
+ return y - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
new file mode 100644
index 00000000000..e3da479e6cb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c
@@ -0,0 +1,54 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators, with multiple iterators and function calls in the iterator
+ expression. */
+
+#include <stdlib.h>
+
+#define DIM1 16
+#define DIM2 15
+
+int mkarrays (int *x[], int *y[])
+{
+ int expected = 0;
+
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ y[i] = (int *) malloc (sizeof (int));
+ *y[i] = rand ();
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j] * *y[i];
+ }
+ }
+
+ return expected;
+}
+
+int f (int i, int j)
+{
+ return i * 4 + j;
+}
+
+int main (void)
+{
+ int *x[DIM1], *y[DIM1];
+ int sum;
+
+ int expected = mkarrays (x, y);
+
+ #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \
+ map(iterator(i=0:DIM1), to: y[i][:1]) \
+ map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j] * y[i][0];
+ }
+
+ return sum - expected;
+}
--
2.34.1
^ permalink raw reply [flat|nested] 4+ messages in thread
* [PATCH 3/3] openmp: Add support for iterators in to/from clauses (C/C++)
2024-05-24 19:57 [PATCH 0/3] openmp: Add support for iterators in OpenMP mapping clauses (C/C++) Kwok Cheung Yeung
2024-05-24 19:59 ` [PATCH 1/3] openmp: Refactor handling of iterators Kwok Cheung Yeung
2024-05-24 20:01 ` [PATCH 2/3] openmp: Add support for iterators in map clauses (C/C++) Kwok Cheung Yeung
@ 2024-05-24 20:02 ` Kwok Cheung Yeung
2 siblings, 0 replies; 4+ messages in thread
From: Kwok Cheung Yeung @ 2024-05-24 20:02 UTC (permalink / raw)
To: Jakub Jelinek, Tobias Burnus, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 83 bytes --]
This patch extends the previous patch to cover to/from clauses in
'target update'.
[-- Attachment #2: 0003-openmp-Add-support-for-iterators-in-to-from-clauses-.patch --]
[-- Type: text/plain, Size: 24502 bytes --]
From 99addc124535307b50fbdeb66c4f90bb0cbeb041 Mon Sep 17 00:00:00 2001
From: Kwok Cheung Yeung <kcyeung@baylibre.com>
Date: Mon, 15 Apr 2024 13:50:22 +0100
Subject: [PATCH 3/3] openmp: Add support for iterators in to/from clauses
(C/C++)
This adds support for iterators in 'to' and 'from' clauses in the
'target update' OpenMP directive.
2024-05-24 Kwok Cheung Yeung <kcyeung@baylibre.com>
gcc/c/
* c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier.
gcc/cp/
* parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier.
gcc/
* gimplify.cc (gimplify_omp_map_iterators): Gimplify iterators in
to/from clauses.
(gimplify_scan_omp_clauses): Call gimplify_omp_map_iterators once to
handle clauses with iterators, then skip subsequent iterator clauses.
* omp-low.cc (scan_sharing_clauses): Skip firstprivate handling for
to/from clauses with iterators.
(lower_omp_target): Handle kinds for to/from clauses with iterators.
* tree-pretty-print.cc (dump_omp_clause): Call dump_omp_map_iterators
for to/from clauses with iterators.
gcc/testsuite/
* c-c++-common/gomp/target-update-iterator-1.c: New.
* c-c++-common/gomp/target-update-iterator-2.c: New.
* c-c++-common/gomp/target-update-iterator-3.c: New.
libgomp/
* target.c (gomp_update): Call gomp_merge_iterator_maps. Free
allocated variables.
* testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New.
* testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New.
---
gcc/c/c-parser.cc | 105 ++++++++++++++--
gcc/cp/parser.cc | 116 ++++++++++++++++--
gcc/gimplify.cc | 17 ++-
gcc/omp-low.cc | 24 +++-
.../gomp/target-update-iterator-1.c | 20 +++
.../gomp/target-update-iterator-2.c | 17 +++
.../gomp/target-update-iterator-3.c | 17 +++
gcc/tree-pretty-print.cc | 20 ++-
libgomp/target.c | 12 ++
.../target-update-iterators-1.c | 65 ++++++++++
.../target-update-iterators-2.c | 57 +++++++++
.../target-update-iterators-3.c | 66 ++++++++++
12 files changed, 509 insertions(+), 27 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc
index 2281148561c..6353b15d64f 100644
--- a/gcc/c/c-parser.cc
+++ b/gcc/c/c-parser.cc
@@ -19185,8 +19185,11 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list)
to ( variable-list )
OpenMP 5.1:
- from ( [present :] variable-list )
- to ( [present :] variable-list ) */
+ from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+ to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+ motion-modifier:
+ present | iterator (iterators-definition) */
static tree
c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
@@ -19197,15 +19200,88 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
if (!parens.require_open (parser))
return list;
+ int pos = 1, colon_pos = 0;
+ int iterator_length = 0;
+ while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+ {
+ if (c_parser_peek_nth_token_raw (parser, pos + 1)->type
+ == CPP_OPEN_PAREN)
+ {
+ unsigned int n = pos + 2;
+ if (c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && (c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN))
+ {
+ iterator_length = n - pos + 1;
+ pos = n;
+ }
+ }
+ if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+ pos += 2;
+ else
+ pos++;
+ if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON)
+ {
+ colon_pos = pos;
+ break;
+ }
+ }
+
bool present = false;
- c_token *token = c_parser_peek_token (parser);
+ tree iterators = NULL_TREE;
- if (token->type == CPP_NAME
- && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0
- && c_parser_peek_2nd_token (parser)->type == CPP_COLON)
+ for (pos = 1; pos < colon_pos; pos++)
{
- present = true;
- c_parser_consume_token (parser);
+ c_token *token = c_parser_peek_token (parser);
+
+ if (token->type == CPP_COMMA)
+ {
+ c_parser_consume_token (parser);
+ continue;
+ }
+ if (token->type == CPP_NAME)
+ {
+ const char *name = IDENTIFIER_POINTER (token->value);
+ if (strcmp (name, "present") == 0)
+ {
+ if (present)
+ {
+ c_parser_error (parser, "too many %<present%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ present = true;
+ c_parser_consume_token (parser);
+ }
+ else if (strcmp (name, "iterator") == 0)
+ {
+ if (iterators)
+ {
+ c_parser_error (parser, "too many %<iterator%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ iterators = c_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ }
+ else
+ {
+ if (kind == OMP_CLAUSE_TO)
+ c_parser_error (parser, "%<to%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ else
+ c_parser_error (parser, "%<from%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ }
+ }
+
+ if (colon_pos)
+ {
+ gcc_assert (pos == colon_pos);
+ gcc_assert (c_parser_next_token_is (parser, CPP_COLON));
c_parser_consume_token (parser);
}
@@ -19216,6 +19292,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind,
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+ if (iterators)
+ {
+ tree block = pop_scope ();
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
+ if (iterators)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
return nl;
}
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 6dc67851f96..f26da16ca13 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -41612,8 +41612,11 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc)
to ( variable-list )
OpenMP 5.1:
- from ( [present :] variable-list )
- to ( [present :] variable-list ) */
+ from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+ to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list )
+
+ motion-modifier:
+ present | iterator (iterators-definition) */
static tree
cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
@@ -41622,15 +41625,94 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
return list;
+ size_t pos = 1, colon_pos = 0;
+ int iterator_length = 0;
+ while (cp_lexer_nth_token_is (parser->lexer, pos, CPP_NAME))
+ {
+ if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN))
+ {
+ unsigned int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
+ {
+ iterator_length = n - pos;
+ pos = n - 1;
+ }
+ }
+ if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_COMMA))
+ pos += 2;
+ else
+ pos++;
+ if (cp_lexer_nth_token_is (parser->lexer, pos, CPP_COLON))
+ {
+ colon_pos = pos;
+ break;
+ }
+ }
+
bool present = false;
- cp_token *token = cp_lexer_peek_token (parser->lexer);
+ tree iterators = NULL_TREE;
+ for (pos = 1; pos < colon_pos; pos++)
+ {
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
- if (token->type == CPP_NAME
- && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0
- && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON))
+ if (token->type == CPP_COMMA)
+ {
+ cp_lexer_consume_token (parser->lexer);
+ continue;
+ }
+ if (token->type == CPP_NAME)
+ {
+ const char *name = IDENTIFIER_POINTER (token->u.value);
+ if (strcmp (name, "present") == 0)
+ {
+ if (present)
+ {
+ cp_parser_error (parser, "too many %<present%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ present = true;
+ cp_lexer_consume_token (parser->lexer);
+ }
+ else if (strcmp (name, "iterator") == 0)
+ {
+ if (iterators)
+ {
+ cp_parser_error (parser, "too many %<iterator%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ begin_scope (sk_omp, NULL);
+ iterators = cp_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ }
+ else
+ {
+ if (kind == OMP_CLAUSE_TO)
+ cp_parser_error (parser, "%<to%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ else
+ cp_parser_error (parser, "%<from%> clause with motion modifier "
+ "other than %<iterator%> or %<present%>");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ }
+ }
+
+ if (colon_pos)
{
- present = true;
- cp_lexer_consume_token (parser->lexer);
+ gcc_assert (pos == colon_pos);
+ gcc_assert (cp_lexer_next_token_is (parser->lexer, CPP_COLON));
cp_lexer_consume_token (parser->lexer);
}
@@ -41639,6 +41721,19 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind,
for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_MOTION_PRESENT (c) = 1;
+ if (iterators)
+ {
+ tree block = poplevel (1, 1, 0);
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
+ if (iterators)
+ for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+
return nl;
}
@@ -41684,9 +41779,8 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
&& strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
&& next_tok->type == CPP_OPEN_PAREN)
{
- size_t n = cp_parser_skip_balanced_tokens (parser, 2);
- if (cp_lexer_peek_nth_token (parser->lexer, n - 1)->type
- == CPP_CLOSE_PAREN)
+ int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
{
iterator_length = n - pos;
pos = n - 1;
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index cadb2a3d96d..3cf88cf82b5 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9304,7 +9304,9 @@ gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
while (tree c = *list_p)
{
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM)
{
list_p = &OMP_CLAUSE_CHAIN (c);
continue;
@@ -12937,6 +12939,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_FROM:
case OMP_CLAUSE__CACHE_:
decl = OMP_CLAUSE_DECL (c);
+
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ if (!handled_map_iterators)
+ {
+ gimplify_omp_map_iterators (list_p, pre_p);
+ handled_map_iterators = true;
+ continue;
+ }
+ /* Skip declarations with iterators. */
+ break;
+ }
+
if (error_operand_p (decl))
{
remove = true;
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index d0f0c5d884a..9e94ad329cd 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1519,8 +1519,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
{
/* FIXME: Is this the right way to handle these? */
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
break;
tree field
= build_decl (OMP_CLAUSE_LOCATION (c),
@@ -13061,7 +13063,23 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, size_int (SIZE_MAX));
- unsigned HOST_WIDE_INT tkind = OMP_CLAUSE_MAP_KIND (c);
+ unsigned HOST_WIDE_INT tkind = 0;
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ tkind = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_ALWAYS_PRESENT_TO : GOMP_MAP_TO);
+ break;
+ case OMP_CLAUSE_FROM:
+ tkind = (OMP_CLAUSE_MOTION_PRESENT (c)
+ ? GOMP_MAP_ALWAYS_PRESENT_FROM : GOMP_MAP_FROM);
+ break;
+ case OMP_CLAUSE_MAP:
+ tkind = OMP_CLAUSE_MAP_KIND (c);
+ break;
+ default:
+ gcc_unreachable ();
+ }
gcc_checking_assert (tkind
< (HOST_WIDE_INT_C (1U) << talign_shift));
gcc_checking_assert (tkind
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
new file mode 100644
index 00000000000..3a64f511da4
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, float **y)
+{
+ #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2])
+
+ #pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+ /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */
+
+ #pragma omp target update to (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error ".to. clause with motion modifier other than .iterator. or .present. before .something." } */
+ /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
new file mode 100644
index 00000000000..3789a559b6f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target update to(iterator(i=0:10): x) /* { dg-error "iterator variable .i. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update from(iterator(i=0:10, j=0:20): x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/
+ ;
+
+ #pragma omp target update to(iterator(i=0:10, j=0:20, k=0:30): x[i], y[j], z[k])
+ /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */
+ /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */
+ /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */
+ ;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
new file mode 100644
index 00000000000..d8672b3a242
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+ #pragma omp target update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3], y[i][j][:DIM3])
+ #pragma omp target update from (iterator(i=0:DIM1): z[i][:DIM2])
+}
+
+/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+" 1 "gimple" } } */
diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc
index feb9c93a809..a0fe270ab1b 100644
--- a/gcc/tree-pretty-print.cc
+++ b/gcc/tree-pretty-print.cc
@@ -1099,16 +1099,28 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
pp_string (pp, "from(");
if (OMP_CLAUSE_MOTION_PRESENT (clause))
pp_string (pp, "present:");
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
+ dump_generic_node (pp, decl, spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE_TO:
pp_string (pp, "to(");
if (OMP_CLAUSE_MOTION_PRESENT (clause))
pp_string (pp, "present:");
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
+ dump_generic_node (pp, decl, spc, flags, false);
goto print_clause_size;
case OMP_CLAUSE__CACHE_:
diff --git a/libgomp/target.c b/libgomp/target.c
index 07ba840b495..71a6428af49 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2221,6 +2221,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
size_t i;
struct splay_tree_key_s cur_node;
const int typemask = short_mapkind ? 0xff : 0x7;
+ bool iterators_p = false;
if (!devicep)
return;
@@ -2228,6 +2229,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
if (mapnum == 0)
return;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds);
+
gomp_mutex_lock (&devicep->lock);
if (devicep->state == GOMP_DEVICE_FINALIZED)
{
@@ -2321,6 +2326,13 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs,
}
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ }
}
static struct gomp_offload_icv_list *
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
new file mode 100644
index 00000000000..5a4cad5c219
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c
@@ -0,0 +1,65 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update to the target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[])
+{
+ int expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ return expected;
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int sum;
+ int expected = mkarray (x);
+
+ #pragma omp target enter data map(to: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ if (sum != expected)
+ return 1;
+
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] *= rand ();
+ expected += x[i][j];
+ }
+
+ #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2])
+
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ return sum != expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
new file mode 100644
index 00000000000..949cc266d84
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c
@@ -0,0 +1,57 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update from the target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] = 0;
+ }
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int sum, expected;
+
+ mkarray (x);
+
+ #pragma omp target enter data map(alloc: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ #pragma omp target map(from: expected)
+ {
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = (i + 1) * (j + 2);
+ expected += x[i][j];
+ }
+ }
+
+ /* Host copy of x should remain unchanged. */
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ if (sum != 0)
+ return 1;
+
+ #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2])
+
+ /* Host copy should now be updated. */
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ return sum - expected;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
new file mode 100644
index 00000000000..852635e50f6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+
+/* Test target enter data and target update to the target using map
+ iterators with a function. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ x[i][j] = rand ();
+ }
+}
+
+int f (int i)
+{
+ return i * 2;
+}
+
+int main (void)
+{
+ int *x[DIM1], x_new[DIM1][DIM2];
+ int sum, expected;
+
+ mkarray (x);
+
+ #pragma omp target enter data map(alloc: x[:DIM1])
+ #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2])
+
+ /* Update x on host. */
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x_new[i][j] = x[i][j];
+ x[i][j] = (i + 1) * (j + 2);
+ }
+
+ /* Update a subset of x on target. */
+ #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2])
+
+ #pragma omp target map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j];
+ }
+
+ /* Calculate expected value on host. */
+ for (int i = 0; i < DIM1/2; i++)
+ for (int j = 0; j < DIM2; j++)
+ x_new[f (i)][j] = x[f (i)][j];
+
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ expected += x_new[i][j];
+
+ return sum - expected;
+}
--
2.34.1
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2024-05-24 20:03 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2024-05-24 19:57 [PATCH 0/3] openmp: Add support for iterators in OpenMP mapping clauses (C/C++) Kwok Cheung Yeung
2024-05-24 19:59 ` [PATCH 1/3] openmp: Refactor handling of iterators Kwok Cheung Yeung
2024-05-24 20:01 ` [PATCH 2/3] openmp: Add support for iterators in map clauses (C/C++) Kwok Cheung Yeung
2024-05-24 20:02 ` [PATCH 3/3] openmp: Add support for iterators in to/from " Kwok Cheung Yeung
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).