From: Paul-Antoine Arras <[email protected]>
An offloaded simd loop with collapse(n > 1) never took the SIMT path:
scan_omp_1_stmt calls scan_omp_simd only for collapse(1), so the loop got no
_simt_ clause and ran on a single lane. The one-lane cap for collapsed loops
in expand_omp_simd was unreachable for the same reason.
This patch lets rectangular collapsed loops take the SIMT path. The lanes are
spread over the innermost loop, which gives contiguous accesses for C loop
nests and Fortran DO nests alike. The outer iteration variables stay uniform
across lanes. Each lane starts a row at its own offset in the innermost loop,
steps by VF times the step, and leaves the row on its own. No division or
modulo is added per iteration.
- In the composite case (a simd loop combined into a worksharing or
distribute construct, which hands it a share of the flattened iteration
space), the lane offset is added where N2VAR is computed, to both the
innermost iteration variable and the linear counter. The row loop exits
when the share ends in the current row; that test is uniform across lanes.
- In the non-composite case, the lane offset is added to the innermost
iteration variable at the start of each row, and the alternate IV is
stepped by VF.
The final values of a collapsed loop's iteration variables do not depend on
the lane, so lastprivate on them does not need the SIMT last-lane exchange.
That covers every Fortran DO variable and every C iteration variable declared
outside the loop. expand_omp_simd stores the final values in every lane
after the loop, and lower_lastprivate_clauses copies them out from the loop
index.
simt_collapse_p keeps these loops on one lane, each with a -fopt-info-omp
note:
- non-rectangular loops and pointer iterators;
- non-composite loops whose unsigned innermost iteration variable counts
down, where the step scaled by VF can wrap;
- lastprivate on variables other than the iteration variables, and linear;
- inscan reductions, and reductions on non-composite loops;
- loop bodies with integer division or modulo, address-taken locals, or
calls to non-internal functions. These give wrong results under amdgcn
SIMT for collapse(1) loops too.
On amdgcn, the change takes effect only with
--param=offload-gpu-mode=simt-directed. In the default (simd) mode,
omp_max_simt_vf returns 0 and the generated code is unchanged. On nvptx,
collapsed simd loops in offloaded code now get SIMT regions.
Assisted-by: Claude Opus 5.5 (testcases + a few bits and pieces in omp-expand)
gcc/ChangeLog:
* omp-expand.cc (expand_omp_simd): Adjust iteration variables, bounds
and steps for collapsed loops.
* omp-low.cc (struct omp_context): Add lastprivate_iv_map.
(delete_omp_context): Delete it.
(omp_find_iv_index): New function.
(scan_omp_for): Fill lastprivate_iv_map for collapsed simd loops.
(simt_collapse_reject, simt_collapse_body, simt_collapse_p): New
functions.
(scan_omp_1_stmt): Also call scan_omp_simd for collapsed loops that
simt_collapse_p accepts.
(lower_lastprivate_clauses): Under SIMT, copy out lastprivate
iteration variables from the loop index instead of using
GOMP_SIMT_LAST_LANE and GOMP_SIMT_XCHG_IDX.
libgomp/ChangeLog:
* testsuite/libgomp.c-c++-common/simt-collapse-1.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-2.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-3.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-4.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c:
New test.
* testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c: New test.
* testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c: New test.
* testsuite/libgomp.fortran/simt-collapse-1.f90: New test.
* testsuite/libgomp.fortran/simt-collapse-2.f90: New test.
* testsuite/libgomp.fortran/simt-collapse-simd-1.f90: New test.
* testsuite/libgomp.fortran/simt-collapse-simd-2.f90: New test.
---
Dependencies: the patch applies cleanly on master. However, to test on amdgcn,
it needs @ams's "SIMT for AMD GCN" and "SIMT reductions" series.
Testing on nvptx does not depend on anything else.
Known limitation: a non-composite inner loop counting up near the limit
of its type can still wrap, as collapse(1) loops can today.
Testing, amdgcn (gfx90a), SIMT series with and without the patch:
- simd mode, libgomp: no regressions. Device assembly of the 413 libgomp
tests with offload code is identical.
- SIMT-directed mode, libgomp: no regressions.
Testing, nvptx (sm_61), master with and without the patch: libgomp has
no regressions, and the new tests pass on the final patch. Five
reduction forms in the new tests are left out because they fail
to link ("unresolved symbol __atomic_compare_exchange_16") with and
without the patch.
OK for mainline once the series is in?
This is a forge pull request published on the gcc-patches mailing list mailing
list
as requested by Paul-Antoine Arras via Sourceware Forge
<[email protected]>.
Forge discussion: https://forge.sourceware.org/gcc/gcc/pulls/278
Get it locally using:
```
git fetch forge-upstream "+refs/versioned_pull/278/*:refs/versioned_pull/278/*"
git switch -c "pr-278-v1" "refs/versioned_pull/278/1/head"
```
Or, download the patch at: https://forge.sourceware.org/gcc/gcc/pulls/278.diff
Created on: 2026-11-13 18:03:51+00:00
Latest update: 2026-11-13 18:10:34+00:00
Changes: 25 changed files, 2562 additions, 38 deletions
Head revision: parras/gcc ref simt-collapse commit
54601f3e74d4630fcea2755ea4d3643bb1ff4b0e
Base revision: gcc/gcc ref trunk commit
785a4b3e6071366e99655a8001aef04da3e1064e r17-6233-g785a4b3e607136
Merge base: 785a4b3e6071366e99655a8001aef04da3e1064e
Requested Reviewers: burnus
Changed files:
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-2.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-3.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-4.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c
- A: libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c
- A: libgomp/testsuite/libgomp.fortran/simt-collapse-1.f90
- A: libgomp/testsuite/libgomp.fortran/simt-collapse-2.f90
- A: libgomp/testsuite/libgomp.fortran/simt-collapse-simd-1.f90
- A: libgomp/testsuite/libgomp.fortran/simt-collapse-simd-2.f90
- M: gcc/omp-expand.cc
- M: gcc/omp-low.cc
gcc/omp-expand.cc | 162 +++++++--
gcc/omp-low.cc | 223 +++++++++++-
.../libgomp.c-c++-common/simt-collapse-1.c | 344 ++++++++++++++++++
.../libgomp.c-c++-common/simt-collapse-2.c | 13 +
.../libgomp.c-c++-common/simt-collapse-3.c | 66 ++++
.../libgomp.c-c++-common/simt-collapse-4.c | 188 ++++++++++
.../simt-collapse-fallback-1.c | 281 ++++++++++++++
.../simt-collapse-fallback-2.c | 8 +
.../simt-collapse-fallback-3.c | 32 ++
.../simt-collapse-lastprivate-1.c | 153 ++++++++
.../simt-collapse-lastprivate-2.c | 12 +
.../simt-collapse-lastprivate-3.c | 152 ++++++++
.../simt-collapse-lastprivate-simd-1.c | 12 +
.../simt-collapse-lastprivate-simd-2.c | 11 +
.../simt-collapse-lastprivate-simd-3.c | 12 +
.../simt-collapse-reduction-1.c | 324 +++++++++++++++++
.../simt-collapse-reduction-simd-1.c | 12 +
.../simt-collapse-simd-1.c | 11 +
.../simt-collapse-simd-2.c | 11 +
.../simt-collapse-simd-3.c | 11 +
.../simt-collapse-simd-4.c | 11 +
.../libgomp.fortran/simt-collapse-1.f90 | 265 ++++++++++++++
.../libgomp.fortran/simt-collapse-2.f90 | 260 +++++++++++++
.../libgomp.fortran/simt-collapse-simd-1.f90 | 15 +
.../libgomp.fortran/simt-collapse-simd-2.f90 | 11 +
25 files changed, 2562 insertions(+), 38 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/simt-collapse-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/simt-collapse-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/simt-collapse-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/simt-collapse-4.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c
create mode 100644 libgomp/testsuite/libgomp.fortran/simt-collapse-1.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/simt-collapse-2.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/simt-collapse-simd-1.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/simt-collapse-simd-2.f90
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index 9304b16f4681a..55798a908231f 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -6802,28 +6802,31 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
is_simt = safelen_int > 1;
}
tree simt_lane = NULL_TREE, simt_maxlane = NULL_TREE;
+ tree vf = NULL_TREE;
if (is_simt)
{
simt_lane = create_tmp_var (unsigned_type_node);
gimple *g = gimple_build_call_internal (IFN_GOMP_SIMT_LANE, 0);
gimple_call_set_lhs (g, simt_lane);
gsi_insert_before (&gsi, g, GSI_SAME_STMT);
- tree offset = fold_build2 (MULT_EXPR, TREE_TYPE (step), step,
- fold_convert (TREE_TYPE (step), simt_lane));
- n1 = fold_convert (type, n1);
- if (POINTER_TYPE_P (type))
- n1 = fold_build_pointer_plus (n1, offset);
- else
- n1 = fold_build2 (PLUS_EXPR, type, n1, fold_convert (type, offset));
+ if (fd->collapse == 1)
+ {
+ /* For non-collapsed SIMT loops, add the lane offset to the low
+ bound. */
+ tree offset
+ = fold_build2 (MULT_EXPR, TREE_TYPE (step), step,
+ fold_convert (TREE_TYPE (step), simt_lane));
+ n1 = fold_convert (type, n1);
+ if (POINTER_TYPE_P (type))
+ n1 = fold_build_pointer_plus (n1, offset);
+ else
+ n1 = fold_build2 (PLUS_EXPR, type, n1, fold_convert (type, offset));
+ }
- /* Collapsed loops not handled for SIMT yet: limit to one lane only. */
- if (fd->collapse > 1)
- simt_maxlane = build_one_cst (unsigned_type_node);
- else if (safelen_int < omp_max_simt_vf ())
+ if (safelen_int < omp_max_simt_vf ())
simt_maxlane = build_int_cst (unsigned_type_node, safelen_int);
- tree vf
- = build_call_expr_internal_loc (UNKNOWN_LOCATION, IFN_GOMP_SIMT_VF,
- unsigned_type_node, 0);
+ vf = build_call_expr_internal_loc (UNKNOWN_LOCATION, IFN_GOMP_SIMT_VF,
+ unsigned_type_node, 0);
if (simt_maxlane)
vf = fold_build2 (MIN_EXPR, unsigned_type_node, vf, simt_maxlane);
vf = fold_convert (TREE_TYPE (step), vf);
@@ -7026,20 +7029,29 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
expand_omp_build_assign (&gsi, fd->loop.v, build_one_cst (type));
if (altv)
{
- t = fold_build2 (PLUS_EXPR, TREE_TYPE (altv), altv,
- build_one_cst (TREE_TYPE (altv)));
+ tree step;
+ if (is_simt && fd->collapse > 1)
+ step = fold_convert (TREE_TYPE (altv), vf);
+ else
+ step = build_one_cst (TREE_TYPE (altv));
+ t = fold_build2 (PLUS_EXPR, TREE_TYPE (altv), altv, step);
expand_omp_build_assign (&gsi, altv, t);
}
if (fd->collapse > 1)
{
i = fd->collapse - 1;
+ tree step = fd->loops[i].step;
+ if (is_simt)
+ /* For collapsed SIMT loops, scale the step by the vectorisation
+ factor. */
+ step = fold_build2 (MULT_EXPR, TREE_TYPE (fd->loops[i].v), step,
+ fold_convert (TREE_TYPE (fd->loops[i].v), vf));
if (POINTER_TYPE_P (TREE_TYPE (fd->loops[i].v)))
- t = fold_build_pointer_plus (fd->loops[i].v, fd->loops[i].step);
+ t = fold_build_pointer_plus (fd->loops[i].v, step);
else
{
- t = fold_convert (TREE_TYPE (fd->loops[i].v),
- fd->loops[i].step);
+ t = fold_convert (TREE_TYPE (fd->loops[i].v), step);
t = fold_build2 (PLUS_EXPR, TREE_TYPE (fd->loops[i].v),
fd->loops[i].v, t);
}
@@ -7112,7 +7124,8 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
}
/* Add 'V -= STEP * (SIMT_VF - 1)' after the loop. */
- if (is_simt)
+ if (is_simt
+ && (fd->collapse == 1 || gimple_omp_for_combined_into_p (fd->for_stmt)))
{
gsi = gsi_start_bb (l2_bb);
step = fold_build2 (MINUS_EXPR, TREE_TYPE (step), orig_step, step);
@@ -7123,6 +7136,43 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
expand_omp_build_assign (&gsi, fd->loop.v, t);
}
+ if (is_simt && fd->collapse > 1 && !broken_loop
+ && omp_find_clause (gimple_omp_for_clauses (fd->for_stmt),
+ OMP_CLAUSE_LASTPRIVATE))
+ {
+ /* After the last iteration, ensure all lanes get the same final IV
+ value. */
+ gsi = gsi_start_bb (l2_bb);
+ for (int k = 0; k < fd->collapse; k++)
+ {
+ struct omp_for_data_loop *l = &fd->loops[k];
+ tree itype = TREE_TYPE (l->v);
+ gcc_assert (!POINTER_TYPE_P (itype) && !l->m1 && !l->m2);
+
+ tree t = build_int_cst (itype, (l->cond_code == LT_EXPR ? -1 : 1));
+ t = fold_build2 (PLUS_EXPR, itype, fold_convert (itype, l->step), t);
+ t = fold_build2 (PLUS_EXPR, itype, t, fold_convert (itype, l->n2));
+ t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, l->n1));
+ if (TYPE_UNSIGNED (itype) && l->cond_code == GT_EXPR)
+ t = fold_build2 (TRUNC_DIV_EXPR, itype,
+ fold_build1 (NEGATE_EXPR, itype, t),
+ fold_build1 (NEGATE_EXPR, itype,
+ fold_convert (itype, l->step)));
+ else
+ t = fold_build2 (TRUNC_DIV_EXPR, itype, t,
+ fold_convert (itype, l->step));
+
+ tree niter = t;
+
+ t = fold_build2 (MULT_EXPR, itype, niter,
+ fold_convert (itype, l->step));
+ t = fold_build2 (PLUS_EXPR, itype, fold_convert (itype, l->n1), t);
+
+ tree final_iv = t;
+ expand_omp_build_assign (&gsi, l->v, final_iv);
+ }
+ }
+
/* Remove GIMPLE_OMP_RETURN. */
gsi = gsi_last_nondebug_bb (exit_bb);
gsi_remove (&gsi, true);
@@ -7194,7 +7244,34 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
gsi = gsi_after_labels (bb);
expand_omp_build_assign (&gsi, fd->loops[i].v, t);
+ basic_block early_exit = NULL;
+ if (is_simt && i + 2 == fd->collapse && n2var)
+ early_exit = bb;
+
bb = split_block (bb, last_nondebug_stmt (bb))->dest;
+
+ /* Composite case: leave when the share ends in this row. Test N2VAR
+ rather than V so that each lane leaves with its own V, which the
+ lastprivate copy-out needs. */
+ if (early_exit)
+ {
+ gsi = gsi_last_bb (early_exit);
+ gcond *c = gimple_build_cond (LT_EXPR, n2var, n2, NULL_TREE,
+ NULL_TREE);
+ gsi_insert_after (&gsi, c, GSI_NEW_STMT);
+ edge et = single_succ_edge (early_exit);
+ et->flags = EDGE_TRUE_VALUE;
+ et->probability
+ = profile_probability::guessed_always ().apply_scale (7, 8);
+ edge ef = make_edge (early_exit, l2_bb, EDGE_FALSE_VALUE);
+ ef->probability = et->probability.invert ();
+ set_immediate_dominator (
+ CDI_DOMINATORS, l2_bb,
+ nearest_common_dominator (
+ CDI_DOMINATORS,
+ get_immediate_dominator (CDI_DOMINATORS, l2_bb), early_exit));
+ }
+
gsi = gsi_start_bb (bb);
tree itype = TREE_TYPE (fd->loops[i].v);
if (fd->loops[i].m2)
@@ -7230,6 +7307,9 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
ne->probability = e->probability.invert ();
gsi = gsi_after_labels (init_bb);
+ if (is_simt && i + 2 == fd->collapse && n2var)
+ /* Composite case: restart V from the end of the previous row. */
+ expand_omp_build_assign (&gsi, fd->loop.v, n2var);
if (fd->loops[i + 1].m1)
{
tree t2 = fold_convert (TREE_TYPE (fd->loops[i + 1].v),
@@ -7250,6 +7330,24 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
t = fold_convert (TREE_TYPE (fd->loops[i + 1].v),
fd->loops[i + 1].n1);
expand_omp_build_assign (&gsi, fd->loops[i + 1].v, t);
+
+ if (is_simt && i + 2 == fd->collapse && !n2var)
+ {
+ /* Non-composite case: add the step scaled by the SIMT lane to the
+ inner IVs. */
+ itype = TREE_TYPE (fd->loops[i + 1].v);
+ if (POINTER_TYPE_P (itype))
+ itype = signed_type_for (itype);
+ tree ilane = fold_convert (itype, simt_lane);
+ t = fold_build2 (MULT_EXPR, itype, ilane,
+ fold_convert (itype, fd->loops[i + 1].step));
+ if (POINTER_TYPE_P (TREE_TYPE (fd->loops[i + 1].v)))
+ t = fold_build_pointer_plus (fd->loops[i + 1].v, t);
+ else
+ t = fold_build2 (PLUS_EXPR, itype, fd->loops[i + 1].v, t);
+ expand_omp_build_assign (&gsi, fd->loops[i + 1].v, t);
+ }
+
if (fd->loops[i + 1].m2)
{
if (i + 2 == fd->collapse && (n2var || altv))
@@ -7323,13 +7421,33 @@ expand_omp_simd (struct omp_region *region, struct
omp_for_data *fd)
e = split_block (init_bb, last_nondebug_stmt (init_bb));
gsi = gsi_after_labels (e->dest);
init_bb = e->dest;
- remove_edge (FALLTHRU_EDGE (entry_bb));
- make_edge (entry_bb, init_bb, EDGE_FALLTHRU);
+ redirect_edge_succ (find_edge (entry_bb, l1_bb), init_bb);
set_immediate_dominator (CDI_DOMINATORS, init_bb, entry_bb);
set_immediate_dominator (CDI_DOMINATORS, l1_bb, init_bb);
t = fold_build2 (MIN_EXPR, type, min_arg1, min_arg2);
t = fold_build2 (PLUS_EXPR, type, fd->loop.v, t);
expand_omp_build_assign (&gsi, n2var, t);
+ if (is_simt)
+ {
+ /* Join block: add lane offset to both the linear and the
+ inner IVs. */
+ tree itype = TREE_TYPE (fd->loops[i + 1].v);
+ if (POINTER_TYPE_P (itype))
+ itype = signed_type_for (itype);
+ tree ilane = fold_convert (itype, simt_lane);
+ t = fold_build2 (MULT_EXPR, itype, ilane,
+ fold_convert (itype, fd->loops[i + 1].step));
+ if (POINTER_TYPE_P (TREE_TYPE (fd->loops[i + 1].v)))
+ t = fold_build_pointer_plus (fd->loops[i + 1].v, t);
+ else
+ t = fold_build2 (PLUS_EXPR, itype, fd->loops[i + 1].v, t);
+ expand_omp_build_assign (&gsi, fd->loops[i + 1].v, t);
+ tree lane = fold_convert (type, simt_lane);
+ t = fold_build2 (MULT_EXPR, type, lane,
+ fold_convert (type, orig_step));
+ t = fold_build2 (PLUS_EXPR, type, fd->loop.v, t);
+ expand_omp_build_assign (&gsi, fd->loop.v, t);
+ }
}
if (i + 2 == fd->collapse && altv)
{
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 2046e01127a27..24327fc70c021 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -127,6 +127,10 @@ struct omp_context
corresponding tracking loop iteration variables. */
hash_map<tree, tree> *lastprivate_conditional_map;
+ /* And a hash map from the lastprivate iteration variables to their
+ corresponding loop index. */
+ hash_map<tree, int> *lastprivate_iv_map;
+
/* And a hash map from the allocate variables to their corresponding
allocators. */
hash_map<tree, tree> *allocate_map;
@@ -1104,6 +1108,7 @@ delete_omp_context (splay_tree_value value)
}
delete ctx->lastprivate_conditional_map;
+ delete ctx->lastprivate_iv_map;
delete ctx->allocate_map;
XDELETE (ctx);
@@ -2856,6 +2861,28 @@ check_oacc_kernel_gwv (gomp_for *stmt, omp_context *ctx)
return outer_mask | this_mask;
}
+/* If DECL is the iteration variable of loop K in the collapsed loop nest
+ STMT, meaning that the gimplifier put "DECL = V" at the start of the body,
+ return K. Otherwise return -1. */
+
+static int
+omp_find_iv_index (const gomp_for *stmt, tree decl)
+{
+ gimple_seq body = gimple_omp_body (stmt);
+ for (gimple_stmt_iterator gsi = gsi_start (body); !gsi_end_p (gsi);
+ gsi_next (&gsi))
+ {
+ gimple *s = gsi_stmt (gsi);
+ if (!gimple_assign_single_p (s))
+ break;
+ if (gimple_assign_lhs (s) == decl)
+ for (size_t k = 0; k < gimple_omp_for_collapse (stmt); k++)
+ if (gimple_assign_rhs1 (s) == gimple_omp_for_index (stmt, k))
+ return k;
+ }
+ return -1;
+}
+
/* Scan a GIMPLE_OMP_FOR. */
static omp_context *
@@ -3020,6 +3047,22 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
scan_sharing_clauses (clauses, ctx);
+ /* Record which lastprivate clauses reference iteration variables. */
+ if (gimple_omp_for_collapse (stmt) > 1
+ && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_SIMD)
+ for (tree clause = clauses; clause; clause = OMP_CLAUSE_CHAIN (clause))
+ if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_LASTPRIVATE)
+ {
+ tree decl = OMP_CLAUSE_DECL (clause);
+ int k = omp_find_iv_index (stmt, decl);
+ if (k != -1)
+ {
+ if (ctx->lastprivate_iv_map == NULL)
+ ctx->lastprivate_iv_map = new hash_map<tree, int>;
+ ctx->lastprivate_iv_map->put (decl, k);
+ }
+ }
+
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
for (i = 0; i < gimple_omp_for_collapse (stmt); i++)
{
@@ -4146,6 +4189,139 @@ setjmp_or_longjmp_p (const_tree fndecl)
return !strcmp (name, "setjmp") || !strcmp (name, "longjmp");
}
+/* Helper function for simt_collapse_p. */
+
+static bool
+simt_collapse_reject (const gimple *stmt, const char *reason)
+{
+ if (dump_enabled_p ())
+ dump_printf_loc (MSG_MISSED_OPTIMIZATION, stmt,
+ "collapsed loop cannot take the SIMT path: %s\n", reason);
+ return false;
+}
+
+/* Callback for walk_gimple_seq. Return integer_zero_node if the loop body
+ prevents the SIMT path, NULL_TREE otherwise. */
+
+static tree
+simt_collapse_body (gimple_stmt_iterator *gsi_p, bool *handled_ops_p,
+ struct walk_stmt_info *wi)
+{
+ gimple *stmt = gsi_stmt (*gsi_p);
+ *handled_ops_p = false;
+
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_ASSIGN:
+ if (gimple_assign_rhs_code (stmt) == TRUNC_DIV_EXPR
+ || gimple_assign_rhs_code (stmt) == TRUNC_MOD_EXPR)
+ {
+ *handled_ops_p = true;
+ simt_collapse_reject (stmt, "integer division or modulo");
+ return integer_zero_node;
+ }
+ break;
+ case GIMPLE_BIND:
+ for (tree var = gimple_bind_vars (as_a <gbind *> (stmt)); var;
+ var = DECL_CHAIN (var))
+ if (TREE_ADDRESSABLE (var))
+ {
+ *handled_ops_p = true;
+ simt_collapse_reject (stmt, "address-taken local variable");
+ return integer_zero_node;
+ }
+ break;
+ case GIMPLE_CALL:
+ if (!gimple_call_internal_p (stmt))
+ {
+ *handled_ops_p = true;
+ simt_collapse_reject (stmt, "external function call");
+ return integer_zero_node;
+ }
+ break;
+ default:
+ break;
+ }
+
+ return NULL_TREE;
+}
+
+/* Return TRUE if the collapsed loop in STMT can take the SIMT path. */
+
+static bool
+simt_collapse_p (const gimple *stmt)
+{
+ int collapse = gimple_omp_for_collapse (stmt);
+ gcc_assert (collapse > 1);
+
+ /* Non-rectangular loops cannot use the SIMT path. */
+ for (int i = 1; i < collapse; i++)
+ {
+ tree n1 = gimple_omp_for_initial (stmt, i);
+ tree n2 = gimple_omp_for_final (stmt, i);
+ if (TREE_CODE (n1) == TREE_VEC || TREE_CODE (n2) == TREE_VEC)
+ return simt_collapse_reject (stmt, "non-rectangular loop");
+ }
+
+ /* Pointer iterators cannot use the SIMT path. */
+ for (int i = 0; i < collapse; i++)
+ if (POINTER_TYPE_P (TREE_TYPE (gimple_omp_for_index (stmt, i))))
+ return simt_collapse_reject (stmt, "pointer iterator");
+
+ /* Non-composite loops whose innermost IV is unsigned and counts down cannot
+ use the SIMT path. */
+ if (!gimple_omp_for_combined_into_p (stmt))
+ {
+ tree innermost_iv = gimple_omp_for_index (stmt, collapse - 1);
+ enum tree_code cond = gimple_omp_for_cond (stmt, collapse - 1);
+ if (cond == NE_EXPR)
+ {
+ tree step = omp_get_for_step_from_incr (
+ gimple_location (stmt), gimple_omp_for_incr (stmt, collapse - 1));
+ gcc_checking_assert (integer_onep (step)
+ || integer_minus_onep (step));
+ cond = integer_onep (step) ? LT_EXPR : GT_EXPR;
+ }
+ if (TYPE_UNSIGNED (TREE_TYPE (innermost_iv))
+ && (cond == GT_EXPR || cond == GE_EXPR))
+ return simt_collapse_reject (
+ stmt, "non-composite loop with unsigned innermost IV counting down");
+ }
+
+ /* Check for OMP clauses that prevent the SIMT path. */
+ tree clauses = gimple_omp_for_clauses (stmt);
+ for (; clauses; clauses = OMP_CLAUSE_CHAIN (clauses))
+ {
+ switch (OMP_CLAUSE_CODE (clauses))
+ {
+ case OMP_CLAUSE_LASTPRIVATE:
+ if (omp_find_iv_index (as_a <const gomp_for *> (stmt),
+ OMP_CLAUSE_DECL (clauses))
+ == -1)
+ return simt_collapse_reject (
+ stmt, "lastprivate clause on a non-iteration variable");
+ break;
+ case OMP_CLAUSE_LINEAR:
+ return simt_collapse_reject (stmt, "linear clause");
+ case OMP_CLAUSE_REDUCTION:
+ if (OMP_CLAUSE_REDUCTION_INSCAN (clauses))
+ return simt_collapse_reject (stmt, "inscan reduction clause");
+ if (!gimple_omp_for_combined_into_p (stmt))
+ return simt_collapse_reject (
+ stmt, "reduction clause on non-composite loop");
+ break;
+ default:
+ break;
+ }
+ }
+
+ /* Walk loop body. */
+ struct walk_stmt_info wi;
+ memset (&wi, 0, sizeof (wi));
+ walk_gimple_seq (gimple_omp_body (stmt), simt_collapse_body, NULL, &wi);
+ return wi.callback_result == NULL_TREE;
+}
+
/* Helper function for scan_omp.
Callback for walk_gimple_stmt used to scan for OMP directives in
@@ -4286,7 +4462,7 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool
*handled_ops_p,
== GF_OMP_FOR_KIND_SIMD)
&& omp_maybe_offloaded_ctx (ctx)
&& omp_max_simt_vf ()
- && gimple_omp_for_collapse (stmt) == 1)
+ && (gimple_omp_for_collapse (stmt) == 1 || simt_collapse_p (stmt)))
scan_omp_simd (gsi, as_a <gomp_for *> (stmt), ctx);
else
scan_omp_for (as_a <gomp_for *> (stmt), ctx);
@@ -7433,23 +7609,38 @@ lower_lastprivate_clauses (tree clauses, tree
predicate, gimple_seq *body_p,
}
else if (maybe_simt)
{
- tree val = (DECL_HAS_VALUE_EXPR_P (new_var)
- ? DECL_VALUE_EXPR (new_var)
- : new_var);
- if (simtlast == NULL)
+ int *k = ctx->lastprivate_iv_map
+ ? ctx->lastprivate_iv_map->get (var)
+ : NULL;
+ if (k == NULL)
+ {
+ tree val = (DECL_HAS_VALUE_EXPR_P (new_var)
+ ? DECL_VALUE_EXPR (new_var)
+ : new_var);
+ if (simtlast == NULL)
+ {
+ simtlast = create_tmp_var (unsigned_type_node);
+ gcall *g
+ = gimple_build_call_internal (IFN_GOMP_SIMT_LAST_LANE,
+ 1, simtcond);
+ gimple_call_set_lhs (g, simtlast);
+ gimple_seq_add_stmt (this_stmt_list, g);
+ }
+ x = build_call_expr_internal_loc (UNKNOWN_LOCATION,
+ IFN_GOMP_SIMT_XCHG_IDX,
+ TREE_TYPE (val), 2, val,
+ simtlast);
+ new_var = unshare_expr (new_var);
+ gimplify_assign (new_var, x, this_stmt_list);
+ new_var = unshare_expr (new_var);
+ }
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE)
{
- simtlast = create_tmp_var (unsigned_type_node);
- gcall *g = gimple_build_call_internal
- (IFN_GOMP_SIMT_LAST_LANE, 1, simtcond);
- gimple_call_set_lhs (g, simtlast);
- gimple_seq_add_stmt (this_stmt_list, g);
+ OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c) = NULL;
+ gimplify_assign (new_var,
+ gimple_omp_for_index (ctx->stmt, *k),
+ &OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c));
}
- x = build_call_expr_internal_loc
- (UNKNOWN_LOCATION, IFN_GOMP_SIMT_XCHG_IDX,
- TREE_TYPE (val), 2, val, simtlast);
- new_var = unshare_expr (new_var);
- gimplify_assign (new_var, x, this_stmt_list);
- new_var = unshare_expr (new_var);
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_LASTPRIVATE
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-1.c
new file mode 100644
index 0000000000000..4b4958011b97f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-1.c
@@ -0,0 +1,344 @@
+/* Check that offloaded composite simd collapse(n) loops get a SIMT region
+ and compute the right result: partial last rows, inner trip counts below
+ and above the vector length, thread shares that start in the middle of a
+ row, shares that span many rows, collapse(3), a non-unit and a negative
+ inner step, schedule, dist_schedule, safelen and simdlen clauses, a
+ target parallel construct, and iteration variables of several types.
+
+ Each loop declares its iteration variables in the for-init, so that they
+ are private, not implicitly lastprivate. Every result array has guard
+ elements on both sides, which catch writes outside the iteration space. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+
+#define PAD 256
+#define MAXN 8192
+
+/* simt-collapse-2.c runs the same loops with num_teams and num_threads. */
+#ifndef SIMT_CLAUSES
+#define SIMT_CLAUSES(TEAMS, THREADS)
+#endif
+
+int buf[MAXN + 2 * PAD];
+
+static void
+clear (void)
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ buf[x] = -1;
+}
+
+/* BUF[PAD + I] must be EXPECT (I) for the I in [0, N), and -1 elsewhere. */
+
+static void
+check (int n, int (*expect) (int))
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ {
+ int i = x - PAD;
+ int want = (i >= 0 && i < n) ? expect (i) : -1;
+ if (buf[x] != want)
+ abort ();
+ }
+}
+
+static int
+ident (int i)
+{
+ return i * 7 + 3;
+}
+
+/* Inner trip count M, not a multiple of the vector length, and N * M not
+ divisible by the number of threads, so shares start in the middle of a
+ row. The bounds are runtime values, so one loop covers all shapes.
+ TEAMS and THREADS are used only by simt-collapse-2.c. Shapes with more
+ iterations than threads, whose shares span rows, are in
+ simt-collapse-3.c. */
+
+static void
+c2_gen (int n, int m, int teams, int threads)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ SIMT_CLAUSES (teams, threads) map(tofrom: buf) firstprivate(n, m)
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < m; j++)
+ buf[PAD + i * m + j] = (i * m + j) * 7 + 3;
+ check (n * m, ident);
+}
+
+static void
+c2_default (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (int i = 0; i < 61; i++)
+ for (int j = 0; j < 67; j++)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* Inner step 3: the lanes must be offset by 3 * lane in the inner variable,
+ but by 1 * lane in the linear counter. The index has no division, which
+ would take the one-lane fallback, so the rows have holes. */
+
+static int
+step3_expect (int i)
+{
+ return i % 3 == 0 ? ident (i) : -1;
+}
+
+static void
+c2_step3 (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ SIMT_CLAUSES (3, 5) map(tofrom: buf)
+ for (int i = 0; i < 19; i++)
+ for (int j = 0; j < 100; j += 3)
+ buf[PAD + i * 102 + j] = (i * 102 + j) * 7 + 3;
+ check (19 * 102, step3_expect);
+}
+
+/* Outer loop with a non-unit start and step. As above, the index has no
+ division, so the rows have holes. */
+
+static int
+outer_step_expect (int i)
+{
+ return (i / 17) % 5 == 0 ? ident (i) : -1;
+}
+
+static void
+c2_outer_step (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ SIMT_CLAUSES (3, 5) map(tofrom: buf)
+ for (int i = 5; i < 50; i += 5)
+ for (int j = 0; j < 17; j++)
+ buf[PAD + (i - 5) * 17 + j] = ((i - 5) * 17 + j) * 7 + 3;
+ check (41 * 17, outer_step_expect);
+}
+
+/* Decreasing inner loop. */
+
+static void
+c2_down (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ SIMT_CLAUSES (3, 5) map(tofrom: buf)
+ for (int i = 0; i < 23; i++)
+ for (int j = 40; j > 0; j--)
+ buf[PAD + i * 40 + (40 - j)] = (i * 40 + (40 - j)) * 7 + 3;
+ check (23 * 40, ident);
+}
+
+static void
+c3 (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(3) \
+ SIMT_CLAUSES (3, 5) map(tofrom: buf)
+ for (int i = 0; i < 7; i++)
+ for (int j = 0; j < 5; j++)
+ for (int k = 0; k < 19; k++)
+ buf[PAD + (i * 5 + j) * 19 + k] = ((i * 5 + j) * 19 + k) * 7 + 3;
+ check (7 * 5 * 19, ident);
+}
+
+/* Empty and one-iteration nests. */
+
+static void
+c2_empty (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (int i = 0; i < 0; i++)
+ for (int j = 0; j < 10; j++)
+ buf[PAD + i * 10 + j] = 0;
+ check (0, ident);
+}
+
+static void
+c2_single (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (int i = 0; i < 1; i++)
+ for (int j = 0; j < 1; j++)
+ buf[PAD + i * 1 + j] = 3;
+ check (1, ident);
+}
+
+/* Shares that span many rows, so the share end is reached before the end of
+ the outer loops. */
+
+static void
+c2_many_rows (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (int i = 0; i < 1000; i++)
+ for (int j = 0; j < 7; j++)
+ buf[PAD + i * 7 + j] = (i * 7 + j) * 7 + 3;
+ check (1000 * 7, ident);
+}
+
+static void
+c3_many_rows (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(3) map(tofrom:
buf)
+ for (int i = 0; i < 50; i++)
+ for (int j = 0; j < 7; j++)
+ for (int k = 0; k < 5; k++)
+ buf[PAD + (i * 7 + j) * 5 + k] = ((i * 7 + j) * 5 + k) * 7 + 3;
+ check (50 * 7 * 5, ident);
+}
+
+/* A 61x67 nest with the clauses in the second argument: shares handed out
+ by other schedules, and safelen or simdlen, which limit the number of
+ active lanes. */
+
+#define DO_PRAGMA(X) _Pragma (#X)
+#define C2_CLAUSES(NAME, ...) \
+static void \
+NAME (void) \
+{ \
+ DO_PRAGMA (omp target teams distribute parallel for simd collapse(2) \
+ map(tofrom: buf) __VA_ARGS__) \
+ for (int i = 0; i < 61; i++) \
+ for (int j = 0; j < 67; j++) \
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3; \
+ check (61 * 67, ident); \
+}
+
+C2_CLAUSES (c2_sched_static_1, schedule(static, 1))
+C2_CLAUSES (c2_dist_static_7, dist_schedule(static, 7))
+C2_CLAUSES (c2_dist_static_500, dist_schedule(static, 500))
+C2_CLAUSES (c2_safelen_2, safelen(2))
+C2_CLAUSES (c2_safelen_8, safelen(8))
+C2_CLAUSES (c2_safelen_33, safelen(33))
+C2_CLAUSES (c2_simdlen_4, simdlen(4))
+
+static void
+c3_safelen_8 (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(3) \
+ safelen(8) map(tofrom: buf)
+ for (int i = 0; i < 9; i++)
+ for (int j = 0; j < 13; j++)
+ for (int k = 0; k < 19; k++)
+ buf[PAD + (i * 13 + j) * 19 + k] = ((i * 13 + j) * 19 + k) * 7 + 3;
+ check (9 * 13 * 19, ident);
+}
+
+/* Without teams. */
+
+static void
+c2_parallel (void)
+{
+#pragma omp target parallel for simd collapse(2) map(tofrom: buf)
+ for (int i = 0; i < 61; i++)
+ for (int j = 0; j < 67; j++)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* Iteration variables of other types. For the decreasing loops, the
+ unsigned division path computes the inner trip count. */
+
+static void
+c2_unsigned (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (unsigned i = 0; i < 61; i++)
+ for (unsigned j = 0; j < 67; j++)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+static void
+c2_unsigned_down (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (unsigned i = 0; i < 61; i++)
+ for (unsigned j = 67; j > 0; j--)
+ buf[PAD + i * 67 + (67 - j)] = (i * 67 + (67 - j)) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* J is 133, 131, ..., 1. The index uses a shift, because a division would
+ take the one-lane fallback. */
+
+static void
+c2_unsigned_down_step2 (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (unsigned i = 0; i < 61; i++)
+ for (unsigned j = 133; j > 0; j -= 2)
+ buf[PAD + i * 67 + ((133 - j) >> 1)]
+ = (i * 67 + ((133 - j) >> 1)) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+static void
+c2_uchar_int (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (unsigned char i = 0; i < 61; i++)
+ for (int j = -67; j < 0; j++)
+ buf[PAD + i * 67 + (j + 67)] = (i * 67 + (j + 67)) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* Inner values outside the range of int. */
+
+static void
+c2_long_long_inner (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
buf)
+ for (int i = 0; i < 61; i++)
+ for (long long j = 0x100000000LL; j < 0x100000000LL + 67; j++)
+ buf[PAD + i * 67 + (int) (j - 0x100000000LL)]
+ = (i * 67 + (int) (j - 0x100000000LL)) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+#define RUN(F) do { clear (); F; } while (0)
+
+int
+main (void)
+{
+ RUN (c2_gen (37, 23, 3, 5));
+ RUN (c2_gen (41, 3, 3, 5));
+ RUN (c2_gen (53, 1, 2, 7));
+ RUN (c2_gen (11, 150, 3, 5));
+ RUN (c2_gen (9, 128, 2, 4));
+ RUN (c2_gen (13, 29, 1, 1));
+ RUN (c2_gen (7, 31, 8, 16));
+ RUN (c2_gen (3, 200, 3, 5));
+ RUN (c2_default ());
+ RUN (c2_step3 ());
+ RUN (c2_outer_step ());
+ RUN (c2_down ());
+ RUN (c3 ());
+ RUN (c2_empty ());
+ RUN (c2_single ());
+ RUN (c2_many_rows ());
+ RUN (c3_many_rows ());
+ RUN (c2_sched_static_1 ());
+ RUN (c2_dist_static_7 ());
+ RUN (c2_dist_static_500 ());
+ RUN (c2_safelen_2 ());
+ RUN (c2_safelen_8 ());
+ RUN (c2_safelen_33 ());
+ RUN (c2_simdlen_4 ());
+ RUN (c3_safelen_8 ());
+ RUN (c2_parallel ());
+ RUN (c2_unsigned ());
+ RUN (c2_unsigned_down ());
+ RUN (c2_unsigned_down_step2 ());
+ RUN (c2_uchar_int ());
+ RUN (c2_long_long_inner ());
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump "GOMP_SIMT_ENTER" "optimized" { target
offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-2.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-2.c
new file mode 100644
index 0000000000000..224c0a8d64c5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-2.c
@@ -0,0 +1,13 @@
+/* The loops of simt-collapse-1.c, with num_teams and num_threads clauses.
+ The c2_gen shapes, whose bounds are only known at run time, fault on the
+ GPU with the SIMT series: a bound that the host code copies into a vector
+ register while EXEC has a single lane set is used inside the SIMT region,
+ where the other lanes read stale values. The other shapes pass. */
+
+/* { dg-do run } */
+/* { dg-xfail-run-if "SIMT loops with runtime bounds and num_teams or
num_threads fault" { offload_device_gcn } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+
+#define SIMT_CLAUSES(TEAMS, THREADS) num_teams(TEAMS) num_threads(THREADS)
+#include "simt-collapse-1.c"
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-3.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-3.c
new file mode 100644
index 0000000000000..acbcd087712f3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-3.c
@@ -0,0 +1,66 @@
+/* Offloaded composite simd collapse(2) loop with runtime bounds and more
+ iterations than threads, so that thread shares span rows. With the SIMT
+ series, a runtime bound can be kept in a vector register that was written
+ before the SIMT region entry, where only one lane is active, and then be
+ read inside the region, collapsed or not.
+
+ The loop declares its iteration variables in the for-init, so that they
+ are private, not implicitly lastprivate. The result array has guard
+ elements on both sides, which catch writes outside the iteration space. */
+
+/* { dg-do run } */
+/* { dg-xfail-run-if "SIMT region reads vector registers written before its
entry" { offload_device_gcn } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+
+#define PAD 256
+#define MAXN 8192
+
+int buf[MAXN + 2 * PAD];
+
+static void
+clear (void)
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ buf[x] = -1;
+}
+
+/* BUF[PAD + I] must be I * 7 + 3 for the I in [0, N), and -1 elsewhere. */
+
+static void
+check (int n)
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ {
+ int i = x - PAD;
+ int want = (i >= 0 && i < n) ? i * 7 + 3 : -1;
+ if (buf[x] != want)
+ abort ();
+ }
+}
+
+static void
+c2_gen (int n, int m)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ map(tofrom: buf) firstprivate(n, m)
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < m; j++)
+ buf[PAD + i * m + j] = (i * m + j) * 7 + 3;
+ check (n * m);
+}
+
+#define RUN(F) do { clear (); F; } while (0)
+
+int
+main (void)
+{
+ RUN (c2_gen (31, 71));
+ RUN (c2_gen (64, 64));
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump "GOMP_SIMT_ENTER" "optimized" { target
offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-4.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-4.c
new file mode 100644
index 0000000000000..1284c2cf6e673
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-4.c
@@ -0,0 +1,188 @@
+/* Check that offloaded simd collapse(n) loops that are not composite get a
+ SIMT region and compute the right result. The whole nest runs in every
+ lane, and each row is split over the lanes. Covered: inner trip counts
+ below, equal to and above the vector length, runtime bounds, collapse(3),
+ a non-unit outer and inner step, a negative inner step, unsigned
+ iteration variables, and safelen and simdlen clauses.
+
+ Each loop declares its iteration variables in the for-init, so that they
+ are private, not implicitly lastprivate. The result array has guard
+ elements on both sides, which catch writes outside the iteration space. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+
+#define PAD 256
+#define MAXN 8192
+
+int buf[MAXN + 2 * PAD];
+
+static void
+clear (void)
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ buf[x] = -1;
+}
+
+/* BUF[PAD + I] must be EXPECT (I) for the I in [0, N), and -1 elsewhere. */
+
+static void
+check (int n, int (*expect) (int))
+{
+ for (int x = 0; x < MAXN + 2 * PAD; x++)
+ {
+ int i = x - PAD;
+ int want = (i >= 0 && i < n) ? expect (i) : -1;
+ if (buf[x] != want)
+ abort ();
+ }
+}
+
+static int
+ident (int i)
+{
+ return i * 7 + 3;
+}
+
+static void
+t2_default (void)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf)
+ for (int i = 0; i < 61; i++)
+ for (int j = 0; j < 67; j++)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* Runtime bounds, so one loop covers all shapes. */
+
+static void
+t2_gen (int n, int m)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf) firstprivate(n, m)
+ for (int i = 0; i < n; i++)
+ for (int j = 0; j < m; j++)
+ buf[PAD + i * m + j] = (i * m + j) * 7 + 3;
+ check (n * m, ident);
+}
+
+static void
+t3_default (void)
+{
+#pragma omp target simd collapse(3) map(tofrom: buf)
+ for (int i = 0; i < 9; i++)
+ for (int j = 0; j < 13; j++)
+ for (int k = 0; k < 19; k++)
+ buf[PAD + (i * 13 + j) * 19 + k] = ((i * 13 + j) * 19 + k) * 7 + 3;
+ check (9 * 13 * 19, ident);
+}
+
+/* Inner step 3: the lanes must be offset by 3 * lane. The index has no
+ division, which would take the one-lane fallback, so the rows have
+ holes. */
+
+static int
+step3_expect (int i)
+{
+ return i % 3 == 0 ? ident (i) : -1;
+}
+
+static void
+t2_step3 (void)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf)
+ for (int i = 0; i < 19; i++)
+ for (int j = 0; j < 200; j += 3)
+ buf[PAD + i * 201 + j] = (i * 201 + j) * 7 + 3;
+ check (19 * 201, step3_expect);
+}
+
+/* Outer loop with a non-unit start and step. As above, the index has no
+ division, so the rows have holes. */
+
+static int
+outer_step_expect (int i)
+{
+ return (i / 70) % 5 == 0 ? ident (i) : -1;
+}
+
+static void
+t2_outer_step (void)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf)
+ for (int i = 5; i < 100; i += 5)
+ for (int j = 0; j < 70; j++)
+ buf[PAD + (i - 5) * 70 + j] = ((i - 5) * 70 + j) * 7 + 3;
+ check (91 * 70, outer_step_expect);
+}
+
+static void
+t2_down (void)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf)
+ for (int i = 0; i < 61; i++)
+ for (int j = 66; j >= 0; j--)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+static void
+t2_unsigned (void)
+{
+#pragma omp target simd collapse(2) map(tofrom: buf)
+ for (unsigned i = 0; i < 61; i++)
+ for (unsigned j = 0; j < 67; j++)
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3;
+ check (61 * 67, ident);
+}
+
+/* safelen and simdlen limit the number of active lanes. */
+
+#define DO_PRAGMA(X) _Pragma (#X)
+#define T2_CLAUSES(NAME, ...) \
+static void \
+NAME (void) \
+{ \
+ DO_PRAGMA (omp target simd collapse(2) map(tofrom: buf) __VA_ARGS__) \
+ for (int i = 0; i < 61; i++) \
+ for (int j = 0; j < 67; j++) \
+ buf[PAD + i * 67 + j] = (i * 67 + j) * 7 + 3; \
+ check (61 * 67, ident); \
+}
+
+T2_CLAUSES (t2_safelen_2, safelen(2))
+T2_CLAUSES (t2_safelen_8, safelen(8))
+T2_CLAUSES (t2_safelen_33, safelen(33))
+T2_CLAUSES (t2_simdlen_4, simdlen(4))
+
+#define RUN(F) do { clear (); F; } while (0)
+
+int
+main (void)
+{
+ RUN (t2_default ());
+ RUN (t2_gen (1, 1));
+ RUN (t2_gen (5, 3));
+ RUN (t2_gen (3, 64));
+ RUN (t2_gen (7, 65));
+ RUN (t2_gen (2, 200));
+ RUN (t2_gen (100, 1));
+ RUN (t2_gen (0, 10));
+ RUN (t2_gen (10, 0));
+ RUN (t3_default ());
+ RUN (t2_step3 ());
+ RUN (t2_outer_step ());
+ RUN (t2_down ());
+ RUN (t2_unsigned ());
+ RUN (t2_safelen_2 ());
+ RUN (t2_safelen_8 ());
+ RUN (t2_safelen_33 ());
+ RUN (t2_simdlen_4 ());
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump "GOMP_SIMT_ENTER" "optimized" { target
offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c
new file mode 100644
index 0000000000000..bdac5790d816e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-1.c
@@ -0,0 +1,281 @@
+/* Check that offloaded simd collapse(2) loops of the forms the SIMT path
+ does not handle take the one-lane fallback, and still compute the right
+ result. Each loop declares its iteration variables in the for-init, so
+ that they are private, not implicitly lastprivate, and only the one
+ construct under test can trigger the fallback. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2 -fdump-tree-omplower-details" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+
+#define N 37
+#define M 23
+
+int a[N * M];
+
+#pragma omp begin declare target
+__attribute__((noinline)) static int
+f (int x)
+{
+ return 3 * x + 1;
+}
+#pragma omp end declare target
+
+static void
+clear (void)
+{
+ for (int x = 0; x < N * M; x++)
+ a[x] = -1;
+}
+
+/* A non-rectangular composite distribute/for simd collapse(2) ICEs in the
+ amdgcn offload compiler (expand_fn_using_insn), in mainline too, so use
+ a plain simd in a target region. */
+
+static void
+non_rect (void)
+{
+#pragma omp target map(tofrom: a)
+#pragma omp simd collapse(2)
+ for (int i = 0; i < M; i++)
+ for (int j = i; j < M; j++)
+ a[i * M + j] = i + j;
+}
+
+/* Start from a pointer parameter: initializing P from the mapped global
+ array A itself ICEs in scan_omp_1_op, in mainline too. */
+
+static void
+pointer_iv (int *b)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ map(tofrom: b[:N * M])
+ for (int *p = b; p < b + N * M; p += M)
+ for (int j = 0; j < M; j++)
+ p[j] = j;
+}
+
+static int
+lastprivate_clause (void)
+{
+ int x = -1;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ map(tofrom: a) lastprivate(x)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ x = i * M + j;
+ a[x] = x;
+ }
+ return x;
+}
+
+/* A non-iterator linear clause is not allowed on a construct combined
+ with distribute, so use a plain simd in a target region. K is stepped
+ in the body: without that, GCC leaves it at its initial value in a
+ collapsed loop, in mainline too. */
+
+static void
+linear_clause (void)
+{
+ int k = 0;
+#pragma omp target map(tofrom: a)
+#pragma omp simd collapse(2) linear(k: 1)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ a[i * M + j] = k;
+ k++;
+ }
+}
+
+/* A reduction on a non-composite loop takes the fallback. The simd loop
+ is nested in a loop construct, the shape on which a SIMT reduction hangs
+ on GCN. Reductions on composite loops take the SIMT path, see
+ simt-collapse-reduction-1.c. */
+
+static int
+reduction_non_composite (void)
+{
+ int s = 0;
+#pragma omp target parallel for reduction(+: s)
+ for (int o = 0; o < 4; o++)
+ {
+ int t = 0;
+#pragma omp simd collapse(2) reduction(+: t)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ t += i * M + j + o;
+ s += t;
+ }
+ return s;
+}
+
+static void
+call (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
a)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ a[i * M + j] = f (i * M + j);
+}
+
+static void
+addressable_local (void)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
a)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ int t = i * M + j;
+ int *q = &t;
+ a[i * M + j] = *q + 1;
+ }
+}
+
+static void
+int_div (int d)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
a)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ a[i * M + j] = (i * M + j) / d;
+}
+
+static void
+int_mod (int d)
+{
+#pragma omp target teams distribute parallel for simd collapse(2) map(tofrom:
a)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ a[i * M + j] = (i * M + j) % d;
+}
+
+/* In a non-composite loop, each lane steps the innermost variable by the
+ step times the SIMT width and tests it against its bound, so an unsigned
+ variable counting down can wrap past 0 and keep running. Check each
+ spelling of the condition: > (the front ends also turn != 0 into it),
+ >= and != on a non-zero bound. */
+
+static void
+unsigned_down_gt (void)
+{
+#pragma omp target map(tofrom: a)
+#pragma omp simd collapse(2)
+ for (int i = 0; i < N; i++)
+ for (unsigned j = M; j > 0; j--)
+ a[i * M + j - 1] = i * M + j - 1;
+}
+
+static void
+unsigned_down_ge (void)
+{
+#pragma omp target map(tofrom: a)
+#pragma omp simd collapse(2)
+ for (int i = 0; i < N; i++)
+ for (unsigned j = M; j >= 1; j--)
+ a[i * M + j - 1] = i * M + j - 1;
+}
+
+static void
+unsigned_down_ne (void)
+{
+#pragma omp target map(tofrom: a)
+#pragma omp simd collapse(2)
+ for (int i = 0; i < N; i++)
+ for (unsigned j = M + 5; j != 5; j--)
+ a[i * M + j - 6] = i * M + j - 6;
+}
+
+int
+main (void)
+{
+ clear ();
+ non_rect ();
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ if (a[i * M + j] != (i < M && j >= i ? i + j : -1))
+ abort ();
+
+ clear ();
+ pointer_iv (a);
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x % M)
+ abort ();
+
+ clear ();
+ if (lastprivate_clause () != N * M - 1)
+ abort ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x)
+ abort ();
+
+ clear ();
+ linear_clause ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x)
+ abort ();
+
+ if (reduction_non_composite ()
+ != 4 * (N * M * (N * M - 1) / 2) + N * M * 6)
+ abort ();
+
+
+ clear ();
+ call ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != 3 * x + 1)
+ abort ();
+
+ clear ();
+ addressable_local ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x + 1)
+ abort ();
+
+ clear ();
+ int_div (3);
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x / 3)
+ abort ();
+
+ clear ();
+ int_mod (3);
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x % 3)
+ abort ();
+
+ clear ();
+ unsigned_down_gt ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x)
+ abort ();
+
+ clear ();
+ unsigned_down_ge ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x)
+ abort ();
+
+ clear ();
+ unsigned_down_ne ();
+ for (int x = 0; x < N * M; x++)
+ if (a[x] != x)
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: non-rectangular loop" 1 "omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: pointer iterator" 1 "omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: lastprivate clause" 1 "omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: linear clause" 1 "omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: reduction clause on non-composite loop" 1 "omplower" { target
offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: external function call" 1 "omplower" { target offload_target_amdgcn } } }
*/
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: address-taken local variable" 1 "omplower" { target offload_target_amdgcn
} } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: integer division or modulo" 2 "omplower" { target offload_target_amdgcn }
} } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: non-composite loop with unsigned innermost IV counting down" 3 "omplower"
{ target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path" 12 "omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-tree-dump-not "_simt_" "omplower" { target
offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c
new file mode 100644
index 0000000000000..f5825c6d325f4
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-2.c
@@ -0,0 +1,8 @@
+/* The loops of simt-collapse-fallback-1.c, compiled without any dump or
+ -fopt-info option: taking the fallback must not depend on dumping. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-fallback-1.c"
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c
new file mode 100644
index 0000000000000..994679ea78de3
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-fallback-3.c
@@ -0,0 +1,32 @@
+/* An inscan reduction on a collapsed simd loop takes the one-lane fallback.
+ Compile only: the loop must not ICE in scan_omp_1_op, and simt_collapse_p
+ must reject it. (The inscan result of such a loop is wrong on the host
+ too, so there is nothing to run.) The scan splits the loop in two, so
+ there are two notes. */
+
+/* { dg-do compile } */
+/* { dg-additional-options "-O2 -fdump-tree-omplower-details" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+
+#define N 37
+#define M 23
+
+int a[N * M];
+
+void
+reduction_inscan (void)
+{
+ int s = 0;
+#pragma omp target map(tofrom: a, s)
+#pragma omp parallel for simd collapse(2) reduction(inscan, +: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ s += i * M + j;
+#pragma omp scan inclusive(s)
+ a[i * M + j] = s;
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path: inscan reduction clause" 2 "omplower" { target offload_target_amdgcn } }
} */
+/* { dg-final { scan-tree-dump-times "collapsed loop cannot take the SIMT
path" 2 "omplower" { target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c
new file mode 100644
index 0000000000000..3a07bb39c704f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-1.c
@@ -0,0 +1,153 @@
+/* Check that offloaded simd collapse(n) loops whose iteration variables are
+ lastprivate get a SIMT region, and leave the iteration variables with
+ their values after sequential execution of the nest. The loops here are
+ not composite; simt-collapse-lastprivate-2.c includes this file with
+ COMPOSITE defined for the composite form.
+
+ The iteration variables are declared outside the nest, so they are
+ implicitly lastprivate, except in t_int_x, which has an explicit
+ lastprivate clause. They are declare target variables, not mapped ones,
+ which simt-collapse-lastprivate-3.c covers. Each nest runs once on the
+ host into REF and once offloaded into BUF, and the elements visited and
+ the final values must match. For a nest with no iterations, only that
+ nothing is visited is checked. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+#include <string.h>
+
+#define SZ 20000
+
+#pragma omp begin declare target
+int buf[SZ];
+int n, m;
+int i, j, k;
+unsigned ui, uj;
+#pragma omp end declare target
+
+int ref[SZ];
+
+#define DO_PRAGMA(X) _Pragma (#X)
+#ifdef COMPOSITE
+#define DIR(...) DO_PRAGMA (omp teams distribute parallel for simd __VA_ARGS__)
+#else
+#define DIR(...) DO_PRAGMA (omp simd __VA_ARGS__)
+#endif
+
+static void
+check (long vi, long vj, long ri, long rj, int visits)
+{
+ if (memcmp (buf, ref, sizeof buf) != 0)
+ abort ();
+ if (visits && (vi != ri || vj != rj))
+ abort ();
+}
+
+/* Run the collapse(2) nest IL, JL with iteration variables I and J of type
+ T, marking element IDX, on the host and offloaded with CLAUSES. */
+
+#define TEST2(NAME, T, I, J, IL, JL, IDX, ...) \
+static void \
+NAME (void) \
+{ \
+ int visits = 0; \
+ memset (ref, 0, sizeof ref); \
+ I = 77, J = 77; \
+ for (IL) \
+ for (JL) \
+ ref[IDX] = 1, visits++; \
+ T ri = I, rj = J; \
+ memset (buf, 0, sizeof buf); \
+ I = 77, J = 77; \
+ DO_PRAGMA (omp target update to(buf, n, m, I, J)) \
+ DO_PRAGMA (omp target) \
+ DIR (collapse(2) __VA_ARGS__)
\
+ for (IL) \
+ for (JL) \
+ buf[IDX] = 1; \
+ DO_PRAGMA (omp target update from(buf, I, J))
\
+ check (I, J, ri, rj, visits);
\
+}
+
+TEST2 (t_int, int, i, j, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j)
+TEST2 (t_int_x, int, i, j, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j,
+ lastprivate(i, j))
+TEST2 (t_safelen2, int, i, j, i = 0; i < n; i++, j = 0; j < m; j++,
+ i * m + j, safelen(2))
+TEST2 (t_safelen8, int, i, j, i = 0; i < n; i++, j = 0; j < m; j++,
+ i * m + j, safelen(8))
+TEST2 (t_step3, int, i, j, i = 0; i < n; i++, j = 0; j < m; j += 3,
+ i * (m + 1) + j)
+TEST2 (t_outer5, int, i, j, i = 5; i < n * 5 + 5; i += 5, j = 0; j < m; j++,
+ (i - 5) * m + j)
+TEST2 (t_down, int, i, j, i = 0; i < n; i++, j = m - 1; j >= 0; j--,
+ i * m + j)
+TEST2 (t_unsigned, unsigned, ui, uj, ui = 0; ui < n; ui++, uj = 0; uj < m;
+ uj++, ui * m + uj)
+#if !defined COMPOSITE && !defined SIMD_MODE
+/* Not composite, an unsigned innermost variable counting down takes the
+ one-lane fallback. The composite form is in
+ simt-collapse-lastprivate-3.c. The simd-mode test defines SIMD_MODE:
+ in simd mode, this loop faults on GCN, with or without lastprivate,
+ independently of SIMT. */
+TEST2 (t_udown3, unsigned, ui, uj, ui = 0; ui < n; ui++, uj = m; uj > 2;
+ uj -= 3, ui * (m + 1) + uj)
+#endif
+
+static void
+t3 (void)
+{
+ int visits = 0;
+ memset (ref, 0, sizeof ref);
+ i = j = k = 77;
+ for (i = 0; i < n; i++)
+ for (j = 0; j < m; j++)
+ for (k = 0; k < 19; k++)
+ ref[(i * m + j) * 19 + k] = 1, visits++;
+ int ri = i, rj = j, rk = k;
+ memset (buf, 0, sizeof buf);
+ i = j = k = 77;
+#pragma omp target update to(buf, n, m, i, j, k)
+#pragma omp target
+ DIR (collapse(3))
+ for (i = 0; i < n; i++)
+ for (j = 0; j < m; j++)
+ for (k = 0; k < 19; k++)
+ buf[(i * m + j) * 19 + k] = 1;
+#pragma omp target update from(buf, i, j, k)
+ check (i, j, ri, rj, visits);
+ if (visits && k != rk)
+ abort ();
+}
+
+int
+main (void)
+{
+ static const int nm[][2] = {
+ { 5, 70 }, { 5, 3 }, { 3, 64 }, { 7, 65 }, { 2, 200 }, { 100, 1 },
+ { 1, 1 }, { 2, 3 }, { 0, 10 }, { 10, 0 }
+ };
+ for (unsigned x = 0; x < sizeof nm / sizeof nm[0]; x++)
+ {
+ n = nm[x][0], m = nm[x][1];
+ t_int ();
+ t_int_x ();
+ t_safelen2 ();
+ t_safelen8 ();
+ t_step3 ();
+ t_outer5 ();
+ t_down ();
+ t_unsigned ();
+#if !defined COMPOSITE && !defined SIMD_MODE
+ t_udown3 ();
+#endif
+ t3 ();
+ }
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump-times "GOMP_SIMT_ENTER" 9 "optimized" {
target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c
new file mode 100644
index 0000000000000..676fe398d0ba2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-2.c
@@ -0,0 +1,12 @@
+/* The loops of simt-collapse-lastprivate-1.c in composite form,
+ target teams distribute parallel for simd, without t_udown3. */
+
+/* { dg-do run } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#define COMPOSITE
+#include "simt-collapse-lastprivate-1.c"
+
+/* { dg-final { scan-offload-tree-dump-times "GOMP_SIMT_ENTER" 9 "optimized" {
target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c
new file mode 100644
index 0000000000000..6764135d0430a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-3.c
@@ -0,0 +1,152 @@
+/* The loops of simt-collapse-lastprivate-1.c with mapped iteration
+ variables, in both forms, and the composite form of t_udown3. With the
+ SIMT series, values that the region reads, such as .omp_data_i, through
+ which the mapped variables are copied out, can be kept in vector
+ registers that were written before the SIMT region entry, where only one
+ lane is active.
+
+ Each nest runs once on the host into REF and once offloaded into BUF,
+ and the elements visited and the final values must match. For a nest
+ with no iterations, only that nothing is visited is checked. */
+
+/* { dg-do run } */
+/* { dg-xfail-run-if "SIMT region reads vector registers written before its
entry" { offload_device_gcn } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+#include <string.h>
+
+#define SZ 20000
+
+int buf[SZ], ref[SZ];
+
+#define DO_PRAGMA(X) _Pragma (#X)
+
+static void
+check (long vi, long vj, long ri, long rj, int visits)
+{
+ if (memcmp (buf, ref, sizeof buf) != 0)
+ abort ();
+ if (visits && (vi != ri || vj != rj))
+ abort ();
+}
+
+/* Run the collapse(2) nest IL, JL with iteration variables i and j of type
+ T, marking element IDX, on the host and offloaded with CLAUSES, as
+ NAME_nc (target simd) and NAME_c (target teams distribute parallel for
+ simd). */
+
+#define TEST2_FORM(NAME, DIR, T, IL, JL, IDX, ...) \
+static void \
+NAME (int n, int m) \
+{ \
+ int visits = 0; \
+ T i = 77, j = 77; \
+ memset (ref, 0, sizeof ref); \
+ for (IL) \
+ for (JL) \
+ ref[IDX] = 1, visits++; \
+ T ri = i, rj = j; \
+ memset (buf, 0, sizeof buf); \
+ i = 77, j = 77; \
+ DO_PRAGMA (omp target map(tofrom: buf, i, j) firstprivate(n, m)) \
+ DO_PRAGMA (omp DIR collapse(2) __VA_ARGS__) \
+ for (IL) \
+ for (JL) \
+ buf[IDX] = 1; \
+ check (i, j, ri, rj, visits);
\
+}
+
+#define TEST2(NAME, T, IL, JL, IDX, ...) \
+ TEST2_FORM (NAME##_nc, simd, T, IL, JL, IDX, __VA_ARGS__) \
+ TEST2_FORM (NAME##_c, teams distribute parallel for simd, T, IL, JL, \
+ IDX, __VA_ARGS__)
+
+TEST2 (t_int, int, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j)
+TEST2 (t_int_x, int, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j,
+ lastprivate(i, j))
+TEST2 (t_safelen2, int, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j,
+ safelen(2))
+TEST2 (t_safelen8, int, i = 0; i < n; i++, j = 0; j < m; j++, i * m + j,
+ safelen(8))
+TEST2 (t_step3, int, i = 0; i < n; i++, j = 0; j < m; j += 3,
+ i * (m + 1) + j)
+TEST2 (t_outer5, int, i = 5; i < n * 5 + 5; i += 5, j = 0; j < m; j++,
+ (i - 5) * m + j)
+TEST2 (t_down, int, i = 0; i < n; i++, j = m - 1; j >= 0; j--, i * m + j)
+TEST2 (t_unsigned, unsigned, i = 0; i < n; i++, j = 0; j < m; j++,
+ i * m + j)
+/* Not composite, an unsigned innermost variable counting down takes the
+ one-lane fallback. The simd-mode test defines SIMD_MODE: in simd mode,
+ the non-composite loop faults on GCN, with or without lastprivate,
+ independently of SIMT. */
+TEST2 (t_udown3, unsigned, i = 0; i < n; i++, j = m; j > 2; j -= 3,
+ i * (m + 1) + j)
+
+#define TEST3(NAME, DIR) \
+static void \
+NAME (int n, int m) \
+{ \
+ int visits = 0; \
+ int i = 77, j = 77, k = 77; \
+ memset (ref, 0, sizeof ref); \
+ for (i = 0; i < n; i++) \
+ for (j = 0; j < m; j++) \
+ for (k = 0; k < 19; k++) \
+ ref[(i * m + j) * 19 + k] = 1, visits++; \
+ int ri = i, rj = j, rk = k; \
+ memset (buf, 0, sizeof buf); \
+ i = 77, j = 77, k = 77; \
+ DO_PRAGMA (omp target map(tofrom: buf, i, j, k) firstprivate(n, m)) \
+ DO_PRAGMA (omp DIR collapse(3)) \
+ for (i = 0; i < n; i++) \
+ for (j = 0; j < m; j++) \
+ for (k = 0; k < 19; k++) \
+ buf[(i * m + j) * 19 + k] = 1; \
+ check (i, j, ri, rj, visits);
\
+ if (visits && k != rk) \
+ abort (); \
+}
+
+TEST3 (t3_nc, simd)
+TEST3 (t3_c, teams distribute parallel for simd)
+
+int
+main (void)
+{
+ static const int nm[][2] = {
+ { 5, 70 }, { 5, 3 }, { 3, 64 }, { 7, 65 }, { 2, 200 }, { 100, 1 },
+ { 1, 1 }, { 2, 3 }, { 0, 10 }, { 10, 0 }
+ };
+ for (unsigned x = 0; x < sizeof nm / sizeof nm[0]; x++)
+ {
+ int n = nm[x][0], m = nm[x][1];
+ t_int_nc (n, m);
+ t_int_c (n, m);
+ t_int_x_nc (n, m);
+ t_int_x_c (n, m);
+ t_safelen2_nc (n, m);
+ t_safelen2_c (n, m);
+ t_safelen8_nc (n, m);
+ t_safelen8_c (n, m);
+ t_step3_nc (n, m);
+ t_step3_c (n, m);
+ t_outer5_nc (n, m);
+ t_outer5_c (n, m);
+ t_down_nc (n, m);
+ t_down_c (n, m);
+ t_unsigned_nc (n, m);
+ t_unsigned_c (n, m);
+#ifndef SIMD_MODE
+ t_udown3_nc (n, m);
+#endif
+ t_udown3_c (n, m);
+ t3_nc (n, m);
+ t3_c (n, m);
+ }
+ return 0;
+}
+
+/* { dg-final { scan-offload-tree-dump-times "GOMP_SIMT_ENTER" 19 "optimized"
{ target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c
new file mode 100644
index 0000000000000..fd06dabbeb849
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-1.c
@@ -0,0 +1,12 @@
+/* The loops of simt-collapse-lastprivate-1.c in the default offload mode, simd
+ on GCN: they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#define SIMD_MODE
+#include "simt-collapse-lastprivate-1.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c
new file mode 100644
index 0000000000000..1138817fb762e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-2.c
@@ -0,0 +1,11 @@
+/* The loops of simt-collapse-lastprivate-2.c in the default offload mode, simd
+ on GCN: they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-lastprivate-2.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c
new file mode 100644
index 0000000000000..d0f3f3b0e9325
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-lastprivate-simd-3.c
@@ -0,0 +1,12 @@
+/* The loops of simt-collapse-lastprivate-3.c in the default offload mode, simd
+ on GCN: they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#define SIMD_MODE
+#include "simt-collapse-lastprivate-3.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c
new file mode 100644
index 0000000000000..9656b9602c42f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-1.c
@@ -0,0 +1,324 @@
+/* Check reductions on offloaded composite simd collapse(n) loops: every
+ reduction operator on int, unsigned, long long, float and double, two
+ reductions in one loop, collapse(3), a non-unit inner step, safelen,
+ num_teams and num_threads, iteration variables that are lastprivate, and
+ complex addition. Inputs are chosen so that every result is exact in any
+ evaluation order, and each result is compared with a serial host loop.
+
+ Forms that the lowering does not give a SIMT region (user-defined and
+ array section reductions, complex multiplication, && and || on floating
+ point) take the plain simd path. They must still compute the right result,
+ and must not be rejected by simt_collapse_p.
+
+ The bounds are constants. A loop whose bounds are firstprivate scalars
+ needs the live-in scalars to be correct in every lane, which the SIMT
+ series does not yet guarantee on GCN, with or without a reduction. The
+ reductions that simt_collapse_p rejects (on a non-composite loop, or
+ inscan) are in simt-collapse-fallback-1.c. */
+
+/* { dg-do run } */
+/* The - reduction operator is deprecated in OpenMP 5.2, but still has to
work. */
+/* { dg-additional-options "-O2 -fdump-tree-omplower-details
-Wno-deprecated-openmp" } */
+/* { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } } */
+/* Two reductions of different types in one loop, complex, user-defined and
+ array section reductions do not link for nvptx ("unresolved symbol
+ __atomic_compare_exchange_16"), with or without the SIMT lowering. */
+/* { dg-additional-options "-DNO_COMBINED_REDUCTIONS" { target
offload_target_nvptx } } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include <stdlib.h>
+
+#define N 37
+#define M 101
+
+static int nfail;
+
+/* Input values, as functions of the flat index K. No division or modulo:
+ those keep a collapsed loop off the SIMT path. */
+#define v_add(k) (((k) & 7) - 3)
+#define v_mul(k) (((k) & 511) == 0 ? 2 : (((k) & 1023) == 513 ? -1 : 1))
+#define v_mm(k) ((((k) * 37) & 1023) - 512 + ((k) == 2345 ? 900 : 0) \
+ - ((k) == 77 ? 900 : 0))
+#define v_and(k) (~(1 << ((k) & 15)) | (((k) & 15) == 5 ? (1 << 5) : 0))
+#define v_or(k) (((k) & 15) == 12 ? 0 : 1 << ((k) & 15))
+#define v_xor(k) (((k) * 2654435761u) >> 7)
+#define v_land(k) ((k) != 1234)
+#define v_lor(k) ((k) == 2222)
+
+#define DO_PRAGMA(X) _Pragma (#X)
+
+#define U_ADD(s, x) s += x
+#define U_MUL(s, x) s *= x
+#define U_MAX(s, x) s = s > x ? s : x
+#define U_MIN(s, x) s = s < x ? s : x
+#define U_AND(s, x) s &= x
+#define U_OR(s, x) s |= x
+#define U_XOR(s, x) s ^= x
+#define U_LAND(s, x) s = s && x
+#define U_LOR(s, x) s = s || x
+
+#define FAIL(NAME) \
+ do { __builtin_printf ("FAIL %s\n", NAME); nfail++; } while (0)
+
+/* A reduction with operator OP of type T over the nest, starting from INIT,
+ with input VAL and update UPD. */
+
+#define RED(T, NAME, OP, INIT, VAL, UPD) \
+static void \
+test_##NAME##_##T (void) \
+{ \
+ T s = INIT, h = INIT;
\
+ DO_PRAGMA (omp target teams distribute parallel for simd collapse(2) \
+ reduction(OP: s)) \
+ for (int i = 0; i < N; i++) \
+ for (int j = 0; j < M; j++)
\
+ {
\
+ T x = (T) VAL (i * M + j); \
+ UPD (s, x); \
+ }
\
+ for (int k = 0; k < N * M; k++) \
+ { \
+ T x = (T) VAL (k); \
+ UPD (h, x); \
+ } \
+ if (s != h) \
+ FAIL (#NAME " " #T); \
+}
+
+typedef unsigned int uint;
+typedef long long llong;
+
+/* && and || on float and double are not supported by the SIMT reduction,
+ and take the plain simd path. */
+#define ARITH(T) \
+ RED (T, add, +, 0, v_add, U_ADD) \
+ RED (T, sub, -, 0, v_add, U_ADD) \
+ RED (T, mul, *, 1, v_mul, U_MUL) \
+ RED (T, max, max, (T) -1000, v_mm, U_MAX) \
+ RED (T, min, min, (T) 1000, v_mm, U_MIN) \
+ RED (T, land, &&, 1, v_land, U_LAND) \
+ RED (T, lor, ||, 0, v_lor, U_LOR)
+#define BITS(T)
\
+ RED (T, and, &, (T) ~0, v_and, U_AND)
\
+ RED (T, or, |, 0, v_or, U_OR)
\
+ RED (T, xor, ^, 0, v_xor, U_XOR)
+
+ARITH (int) BITS (int)
+ARITH (uint) BITS (uint)
+ARITH (llong) BITS (llong)
+ARITH (float)
+ARITH (double)
+
+#define CALLA(T) \
+ test_add_##T (); test_sub_##T (); test_mul_##T (); test_max_##T (); \
+ test_min_##T (); test_land_##T (); test_lor_##T ();
+#define CALLB(T) test_and_##T (); test_or_##T (); test_xor_##T ();
+
+#ifndef NO_COMBINED_REDUCTIONS
+/* Two reductions of different types and operators in one loop. */
+
+static void
+test_two (void)
+{
+ int s = 0, hs = 0;
+ double m = -1000.0, hm = -1000.0;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: s) reduction(max: m)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ s += v_add (i * M + j);
+ m = m > v_mm (i * M + j) ? m : v_mm (i * M + j);
+ }
+ for (int k = 0; k < N * M; k++)
+ {
+ hs += v_add (k);
+ hm = hm > v_mm (k) ? hm : v_mm (k);
+ }
+ if (s != hs || m != hm)
+ FAIL ("two reductions");
+}
+#endif
+
+static void
+test_collapse3 (void)
+{
+ int s = 0, h = 0;
+#pragma omp target teams distribute parallel for simd collapse(3) \
+ reduction(+: s)
+ for (int i = 0; i < 5; i++)
+ for (int j = 0; j < 7; j++)
+ for (int k = 0; k < 67; k++)
+ s += v_add ((i * 7 + j) * 67 + k);
+ for (int k = 0; k < 5 * 7 * 67; k++)
+ h += v_add (k);
+ if (s != h)
+ FAIL ("collapse(3)");
+}
+
+static void
+test_inner_step3 (void)
+{
+ int s = 0, h = 0;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < 200; j += 3)
+ s += v_add (i * 200 + j);
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < 200; j += 3)
+ h += v_add (i * 200 + j);
+ if (s != h)
+ FAIL ("inner step 3");
+}
+
+static void
+test_safelen (void)
+{
+ int s = 0, h = 0;
+#pragma omp target teams distribute parallel for simd collapse(2) safelen(8) \
+ reduction(+: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ s += v_add (i * M + j);
+ for (int k = 0; k < N * M; k++)
+ h += v_add (k);
+ if (s != h)
+ FAIL ("safelen");
+}
+
+static void
+test_num_teams_threads (void)
+{
+ int s = 0, h = 0;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: s) num_teams(3) num_threads(5)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ s += v_add (i * M + j);
+ for (int k = 0; k < N * M; k++)
+ h += v_add (k);
+ if (s != h)
+ FAIL ("num_teams and num_threads");
+}
+
+/* Iteration variables declared outside the loop are lastprivate. */
+
+static void
+test_lastprivate_iv (void)
+{
+ int s = 0, h = 0, i, j;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: s) map(tofrom: i, j)
+ for (i = 0; i < N; i++)
+ for (j = 0; j < M; j++)
+ s += v_add (i * M + j);
+ for (int k = 0; k < N * M; k++)
+ h += v_add (k);
+ if (s != h || i != N || j != M)
+ FAIL ("lastprivate iteration variables");
+}
+
+#ifndef NO_COMBINED_REDUCTIONS
+static void
+test_complex_add (void)
+{
+ _Complex double s = 0, h = 0;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ s += v_add (i * M + j) + 1.0i * v_add (i * M + j + 1);
+ for (int k = 0; k < N * M; k++)
+ h += v_add (k) + 1.0i * v_add (k + 1);
+ if (s != h)
+ FAIL ("complex add");
+}
+
+/* The next three get no SIMT region from the lowering. */
+
+static void
+test_complex_mul (void)
+{
+ _Complex double s = 1, h = 1;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(*: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ s *= ((i * M + j) == 100 || (i * M + j) == 900) ? 1.0i : 1.0;
+ for (int k = 0; k < N * M; k++)
+ h *= (k == 100 || k == 900) ? 1.0i : 1.0;
+ if (s != h)
+ FAIL ("complex mul");
+}
+
+#pragma omp declare reduction (mymax: int : \
+ omp_out = omp_out > omp_in ? omp_out : omp_in) \
+ initializer (omp_priv = -1000000)
+
+static void
+test_user_defined (void)
+{
+ int s = -1000000, h = -1000000;
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(mymax: s)
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ {
+ int x = v_mm (i * M + j);
+ s = s > x ? s : x;
+ }
+ for (int k = 0; k < N * M; k++)
+ h = h > v_mm (k) ? h : v_mm (k);
+ if (s != h)
+ FAIL ("user-defined reduction");
+}
+
+static void
+test_array_section (void)
+{
+ int a[4] = { 0, 0, 0, 0 }, h[4] = { 0, 0, 0, 0 };
+#pragma omp target teams distribute parallel for simd collapse(2) \
+ reduction(+: a[:4])
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ a[(i + j) & 3] += v_add (i * M + j);
+ for (int i = 0; i < N; i++)
+ for (int j = 0; j < M; j++)
+ h[(i + j) & 3] += v_add (i * M + j);
+ for (int x = 0; x < 4; x++)
+ if (a[x] != h[x])
+ FAIL ("array section");
+}
+#endif
+
+int
+main (void)
+{
+ CALLA (int) CALLB (int)
+ CALLA (uint) CALLB (uint)
+ CALLA (llong) CALLB (llong)
+ CALLA (float)
+ CALLA (double)
+#ifndef NO_COMBINED_REDUCTIONS
+ test_two ();
+#endif
+ test_collapse3 ();
+ test_inner_step3 ();
+ test_safelen ();
+ test_num_teams_threads ();
+ test_lastprivate_iv ();
+#ifndef NO_COMBINED_REDUCTIONS
+ test_complex_add ();
+ test_complex_mul ();
+ test_user_defined ();
+ test_array_section ();
+#endif
+ if (nfail)
+ abort ();
+ return 0;
+}
+
+/* No composite loop of this file is rejected by simt_collapse_p. */
+/* { dg-final { scan-tree-dump-not "collapsed loop cannot take the SIMT path"
"omplower" { target offload_target_amdgcn } } } */
+/* { dg-final { scan-offload-tree-dump "GOMP_SIMT_ENTER" "optimized" { target
offload_target_amdgcn } } } */
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c
new file mode 100644
index 0000000000000..e7960c0d63049
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-reduction-simd-1.c
@@ -0,0 +1,12 @@
+/* The loops of simt-collapse-reduction-1.c in the default offload mode, simd
+ on GCN: they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2 -Wno-deprecated-openmp" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+/* { dg-additional-options "-DNO_COMBINED_REDUCTIONS" { target
offload_target_nvptx } } */
+
+#include "simt-collapse-reduction-1.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c
new file mode 100644
index 0000000000000..4dcfa29d99537
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-1.c
@@ -0,0 +1,11 @@
+/* The loops of simt-collapse-1.c in the default offload mode, simd on GCN:
+ they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-1.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c
new file mode 100644
index 0000000000000..69b5bd470731b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-2.c
@@ -0,0 +1,11 @@
+/* The loops of simt-collapse-2.c in the default offload mode, simd on GCN:
+ they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-2.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c
new file mode 100644
index 0000000000000..965c53ec05a38
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-3.c
@@ -0,0 +1,11 @@
+/* The loops of simt-collapse-3.c in the default offload mode, simd on GCN:
+ they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-3.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c
b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c
new file mode 100644
index 0000000000000..b3884ece4ffeb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/simt-collapse-simd-4.c
@@ -0,0 +1,11 @@
+/* The loops of simt-collapse-4.c in the default offload mode, simd on GCN:
+ they must get no SIMT region, and compute the same results. */
+
+/* { dg-do run } */
+/* { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } } */
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } } */
+
+#include "simt-collapse-4.c"
+
+/* { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } } */
diff --git a/libgomp/testsuite/libgomp.fortran/simt-collapse-1.f90
b/libgomp/testsuite/libgomp.fortran/simt-collapse-1.f90
new file mode 100644
index 0000000000000..fb546bfa53687
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/simt-collapse-1.f90
@@ -0,0 +1,265 @@
+! Check that offloaded simd collapse(n) DO loops that are not composite get
+! a SIMT region, and leave the DO variables with their values after
+! sequential execution of the nest. The DO variables are implicitly
+! lastprivate, except in t_int_x, which has an explicit lastprivate clause.
+! They are declare target variables; simt-collapse-2.f90 has the composite
+! form and mapped DO variables. Each nest runs once on the host into REF
+! and once offloaded into BUF, and the elements visited and the final
+! values must match. For a nest with no iterations, only that nothing is
+! visited is checked.
+
+! { dg-do run }
+! { dg-additional-options "-O2" }
+! { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } }
+! { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } }
+
+module simt_collapse_1_m
+ implicit none
+ integer, parameter :: sz = 20000
+ integer :: buf(0:sz-1), n, m, i, j, k
+ integer(2) :: j2
+ integer(8) :: i8
+ !$omp declare target(buf, n, m, i, j, k, j2, i8)
+ integer :: ref(0:sz-1), visits
+contains
+ subroutine check (v1, v2, v3, r1, r2, r3)
+ integer(8) :: v1, v2, v3, r1, r2, r3
+ if (any (buf /= ref)) stop 1
+ if (visits /= 0 .and. (v1 /= r1 .or. v2 /= r2 .or. v3 /= r3)) stop 2
+ end subroutine
+
+ subroutine t_int
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = 0, m - 1
+ ref(j*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2)
+ do j = 0, n - 1
+ do i = 0, m - 1
+ buf(j*m+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_int_x
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = 0, m - 1
+ ref(j*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2) lastprivate(i, j)
+ do j = 0, n - 1
+ do i = 0, m - 1
+ buf(j*m+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_safelen8
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = 0, m - 1
+ ref(j*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2) safelen(8)
+ do j = 0, n - 1
+ do i = 0, m - 1
+ buf(j*m+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_one
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 1, n
+ do i = 1, m
+ ref((j-1)*m+i-1) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2)
+ do j = 1, n
+ do i = 1, m
+ buf((j-1)*m+i-1) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_step3
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = 0, m - 1, 3
+ ref(j*(m+1)+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2)
+ do j = 0, n - 1
+ do i = 0, m - 1, 3
+ buf(j*(m+1)+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_outer5
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 5, n * 5, 5
+ do i = 0, m - 1
+ ref((j-5)*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2)
+ do j = 5, n * 5, 5
+ do i = 0, m - 1
+ buf((j-5)*m+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_down
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = m - 1, 0, -1
+ ref(j*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target
+ !$omp simd collapse(2)
+ do j = 0, n - 1
+ do i = m - 1, 0, -1
+ buf(j*m+i) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ ! The DO variables' kinds differ from each other and from the count's.
+ subroutine t_kind
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i8 = 77; j2 = 77
+ do j2 = 0, n - 1
+ do i8 = 0, m - 1
+ ref(j2*m+i8) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i8; rj = j2
+ buf = 0; i8 = 77; j2 = 77
+ !$omp target update to(buf, i8, j2)
+ !$omp target
+ !$omp simd collapse(2)
+ do j2 = 0, n - 1
+ do i8 = 0, m - 1
+ buf(j2*m+i8) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i8, j2)
+ call check (i8, int (j2, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine t_c3
+ integer(8) :: ri, rj, rk
+ ref = 0; visits = 0; i = 77; j = 77; k = 77
+ do k = 0, n - 1
+ do j = 0, 2
+ do i = 0, m - 1
+ ref((k*3+j)*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ end do
+ ri = i; rj = j; rk = k
+ buf = 0; i = 77; j = 77; k = 77
+ !$omp target update to(buf, i, j, k)
+ !$omp target
+ !$omp simd collapse(3)
+ do k = 0, n - 1
+ do j = 0, 2
+ do i = 0, m - 1
+ buf((k*3+j)*m+i) = 1
+ end do
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf, i, j, k)
+ call check (int (i, 8), int (j, 8), int (k, 8), ri, rj, rk)
+ end subroutine
+end module
+
+program main
+ use simt_collapse_1_m
+ implicit none
+ integer :: s
+ integer, parameter :: ns(10) = [5, 5, 3, 7, 2, 100, 1, 2, 0, 10]
+ integer, parameter :: ms(10) = [70, 3, 64, 65, 200, 1, 1, 3, 10, 0]
+ do s = 1, 10
+ n = ns(s); m = ms(s)
+ !$omp target update to(n, m)
+ call t_int
+ call t_int_x
+ call t_safelen8
+ call t_one
+ call t_step3
+ call t_outer5
+ call t_down
+ call t_kind
+ call t_c3
+ end do
+end program
+
+! { dg-final { scan-offload-tree-dump-times "GOMP_SIMT_ENTER" 9 "optimized" {
target offload_target_amdgcn } } }
diff --git a/libgomp/testsuite/libgomp.fortran/simt-collapse-2.f90
b/libgomp/testsuite/libgomp.fortran/simt-collapse-2.f90
new file mode 100644
index 0000000000000..4426bf6e07236
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/simt-collapse-2.f90
@@ -0,0 +1,260 @@
+! Offloaded simd collapse(n) DO loops that get a SIMT region: the composite
+! form with declare target DO variables, mapped DO variables in both forms,
+! and stencils on mapped arrays shaped like SPEC ACCEL swim (collapse(2))
+! and seismic (collapse(3)). With the SIMT series, values that the region
+! reads, such as .omp_data_i and the loop bounds passed through it, can be
+! kept in vector registers that were written before the SIMT region entry,
+! where only one lane is active. With the current series only the seismic
+! routine still fails this way; the other routines pass.
+!
+! Each nest runs once on the host and once offloaded, and the elements
+! visited, the results and the final DO variables must match. For a nest
+! with no iterations, only that nothing is visited is checked.
+
+! { dg-do run }
+! { dg-xfail-run-if "SIMT region reads vector registers written before its
entry" { offload_device_gcn } }
+! { dg-additional-options "-O2" }
+! { dg-additional-options "--param=offload-gpu-mode=simt-directed" { target
offload_target_amdgcn } }
+! { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } }
+
+module simt_collapse_2_m
+ implicit none
+ integer, parameter :: sz = 20000
+ integer :: buf(0:sz-1), n, m, i, j, k
+ integer(2) :: j2
+ integer(8) :: i8
+ !$omp declare target(buf, n, m, i, j, k, j2, i8)
+ integer :: ref(0:sz-1), visits
+contains
+ subroutine check (v1, v2, v3, r1, r2, r3)
+ integer(8) :: v1, v2, v3, r1, r2, r3
+ if (any (buf /= ref)) stop 1
+ if (visits /= 0 .and. (v1 /= r1 .or. v2 /= r2 .or. v3 /= r3)) stop 2
+ end subroutine
+
+ ! Composite, declare target DO variables.
+
+ subroutine c_int
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i = 77; j = 77
+ do j = 0, n - 1
+ do i = 0, m - 1
+ ref(j*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i; rj = j
+ buf = 0; i = 77; j = 77
+ !$omp target update to(buf, i, j)
+ !$omp target teams distribute parallel do simd collapse(2)
+ do j = 0, n - 1
+ do i = 0, m - 1
+ buf(j*m+i) = 1
+ end do
+ end do
+ !$omp target update from(buf, i, j)
+ call check (int (i, 8), int (j, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine c_kind
+ integer(8) :: ri, rj
+ ref = 0; visits = 0; i8 = 77; j2 = 77
+ do j2 = 0, n - 1
+ do i8 = 0, m - 1
+ ref(j2*m+i8) = 1; visits = visits + 1
+ end do
+ end do
+ ri = i8; rj = j2
+ buf = 0; i8 = 77; j2 = 77
+ !$omp target update to(buf, i8, j2)
+ !$omp target teams distribute parallel do simd collapse(2)
+ do j2 = 0, n - 1
+ do i8 = 0, m - 1
+ buf(j2*m+i8) = 1
+ end do
+ end do
+ !$omp target update from(buf, i8, j2)
+ call check (i8, int (j2, 8), 0_8, ri, rj, 0_8)
+ end subroutine
+
+ subroutine c_c3
+ integer(8) :: ri, rj, rk
+ ref = 0; visits = 0; i = 77; j = 77; k = 77
+ do k = 0, n - 1
+ do j = 0, 2
+ do i = 0, m - 1
+ ref((k*3+j)*m+i) = 1; visits = visits + 1
+ end do
+ end do
+ end do
+ ri = i; rj = j; rk = k
+ buf = 0; i = 77; j = 77; k = 77
+ !$omp target update to(buf, i, j, k)
+ !$omp target teams distribute parallel do simd collapse(3)
+ do k = 0, n - 1
+ do j = 0, 2
+ do i = 0, m - 1
+ buf((k*3+j)*m+i) = 1
+ end do
+ end do
+ end do
+ !$omp target update from(buf, i, j, k)
+ call check (int (i, 8), int (j, 8), int (k, 8), ri, rj, rk)
+ end subroutine
+
+ ! Mapped DO variables. The composite construct is split, because an
+ ! iteration variable cannot be mapped on the combined one.
+
+ subroutine nc_mapped
+ integer :: p, q
+ integer(8) :: rp, rq
+ ref = 0; visits = 0; p = 77; q = 77
+ do q = 0, n - 1
+ do p = 0, m - 1
+ ref(q*m+p) = 1; visits = visits + 1
+ end do
+ end do
+ rp = p; rq = q
+ buf = 0; p = 77; q = 77
+ !$omp target update to(buf)
+ !$omp target map(tofrom: p, q)
+ !$omp simd collapse(2)
+ do q = 0, n - 1
+ do p = 0, m - 1
+ buf(q*m+p) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf)
+ call check (int (p, 8), int (q, 8), 0_8, rp, rq, 0_8)
+ end subroutine
+
+ subroutine c_mapped
+ integer :: p, q
+ integer(8) :: rp, rq
+ ref = 0; visits = 0; p = 77; q = 77
+ do q = 0, n - 1
+ do p = 0, m - 1
+ ref(q*m+p) = 1; visits = visits + 1
+ end do
+ end do
+ rp = p; rq = q
+ buf = 0; p = 77; q = 77
+ !$omp target update to(buf)
+ !$omp target map(tofrom: p, q)
+ !$omp teams distribute parallel do simd collapse(2)
+ do q = 0, n - 1
+ do p = 0, m - 1
+ buf(q*m+p) = 1
+ end do
+ end do
+ !$omp end target
+ !$omp target update from(buf)
+ call check (int (p, 8), int (q, 8), 0_8, rp, rq, 0_8)
+ end subroutine
+
+ ! swim's CALC1 shape: lanes on the first index, local DO variables.
+
+ subroutine swim (mm, nn)
+ integer :: mm, nn
+ integer, allocatable :: p(:,:), u(:,:), v(:,:), cu(:,:), z(:,:), &
+ cuh(:,:), zh(:,:)
+ integer :: ii, jj
+ allocate (p(mm+1,nn+1), u(mm+1,nn+1), v(mm+1,nn+1), cu(mm+1,nn+1), &
+ z(mm+1,nn+1), cuh(mm+1,nn+1), zh(mm+1,nn+1))
+ do jj = 1, nn + 1
+ do ii = 1, mm + 1
+ p(ii,jj) = mod (ii * 7 + jj * 13, 101)
+ u(ii,jj) = mod (ii * 3 + jj * 5, 37) - 18
+ v(ii,jj) = mod (ii * 11 + jj * 2, 29) - 14
+ end do
+ end do
+ cu = -1; z = -1; cuh = -1; zh = -1
+ do jj = 1, nn
+ do ii = 1, mm
+ cuh(ii+1,jj) = (p(ii+1,jj) + p(ii,jj)) * u(ii+1,jj)
+ zh(ii+1,jj+1) = (v(ii+1,jj+1) - v(ii,jj+1) - u(ii+1,jj+1) &
+ + u(ii+1,jj)) * (p(ii,jj) + p(ii+1,jj) &
+ + p(ii+1,jj+1) + p(ii,jj+1))
+ end do
+ end do
+ !$omp target teams distribute parallel do simd collapse(2) &
+ !$omp map(to: p, u, v) map(tofrom: cu, z)
+ do jj = 1, nn
+ do ii = 1, mm
+ cu(ii+1,jj) = (p(ii+1,jj) + p(ii,jj)) * u(ii+1,jj)
+ z(ii+1,jj+1) = (v(ii+1,jj+1) - v(ii,jj+1) - u(ii+1,jj+1) &
+ + u(ii+1,jj)) * (p(ii,jj) + p(ii+1,jj) &
+ + p(ii+1,jj+1) + p(ii,jj+1))
+ end do
+ end do
+ if (any (cu /= cuh) .or. any (z /= zh)) stop 3
+ end subroutine
+
+ ! seismic's velocity-update shape, collapse(3).
+
+ subroutine seismic (nx, ny, nz)
+ integer :: nx, ny, nz
+ integer, allocatable :: sxx(:,:,:), sxy(:,:,:), vx(:,:,:), vxh(:,:,:)
+ integer :: ii, jj, kk
+ allocate (sxx(nx+1,ny+1,nz+1), sxy(nx+1,ny+1,nz+1), vx(nx,ny,nz), &
+ vxh(nx,ny,nz))
+ do kk = 1, nz + 1
+ do jj = 1, ny + 1
+ do ii = 1, nx + 1
+ sxx(ii,jj,kk) = mod (ii * 5 + jj * 3 + kk * 7, 61) - 30
+ sxy(ii,jj,kk) = mod (ii * 2 + jj * 9 + kk, 43) - 21
+ end do
+ end do
+ end do
+ vx = 3; vxh = 3
+ do kk = 2, nz
+ do jj = 2, ny
+ do ii = 2, nx
+ vxh(ii,jj,kk) = vxh(ii,jj,kk) &
+ + 2 * (sxx(ii,jj,kk) - sxx(ii-1,jj,kk)) &
+ + 3 * (sxy(ii,jj,kk) - sxy(ii,jj-1,kk))
+ end do
+ end do
+ end do
+ !$omp target teams distribute parallel do simd collapse(3) &
+ !$omp map(to: sxx, sxy) map(tofrom: vx)
+ do kk = 2, nz
+ do jj = 2, ny
+ do ii = 2, nx
+ vx(ii,jj,kk) = vx(ii,jj,kk) &
+ + 2 * (sxx(ii,jj,kk) - sxx(ii-1,jj,kk)) &
+ + 3 * (sxy(ii,jj,kk) - sxy(ii,jj-1,kk))
+ end do
+ end do
+ end do
+ if (any (vx /= vxh)) stop 4
+ end subroutine
+end module
+
+program main
+ use simt_collapse_2_m
+ implicit none
+ integer :: s
+ integer, parameter :: ns(10) = [5, 5, 3, 7, 2, 100, 1, 2, 0, 10]
+ integer, parameter :: ms(10) = [70, 3, 64, 65, 200, 1, 1, 3, 10, 0]
+ integer, parameter :: sm(5) = [1, 5, 64, 65, 513]
+ integer, parameter :: sn(5) = [1, 70, 64, 129, 257]
+ integer, parameter :: xs(5) = [2, 5, 66, 65, 130]
+ integer, parameter :: ys(5) = [2, 4, 3, 65, 41]
+ integer, parameter :: zs(5) = [2, 70, 5, 9, 17]
+ do s = 1, 10
+ n = ns(s); m = ms(s)
+ !$omp target update to(n, m)
+ call c_int
+ call c_kind
+ call c_c3
+ call nc_mapped
+ call c_mapped
+ end do
+ do s = 1, 5
+ call swim (sm(s), sn(s))
+ call seismic (xs(s), ys(s), zs(s))
+ end do
+end program
+
+! { dg-final { scan-offload-tree-dump-times "GOMP_SIMT_ENTER" 7 "optimized" {
target offload_target_amdgcn } } }
diff --git a/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-1.f90
b/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-1.f90
new file mode 100644
index 0000000000000..4c3f569c75d95
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-1.f90
@@ -0,0 +1,15 @@
+! The loops of simt-collapse-1.f90 in the default offload mode, simd on GCN:
+! they must get no SIMT region, and compute the same results.
+
+! { dg-do run }
+! { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } }
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } }
+
+! In simd mode, the offload compiler warns on t_one's loop, whose accesses
+! stay within BUF.
+! { dg-prune-output "iteration \[0-9\]+ invokes undefined behavior" }
+
+include "simt-collapse-1.f90"
+
+! { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } }
diff --git a/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-2.f90
b/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-2.f90
new file mode 100644
index 0000000000000..2c65e485b5dac
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/simt-collapse-simd-2.f90
@@ -0,0 +1,11 @@
+! The loops of simt-collapse-2.f90 in the default offload mode, simd on GCN:
+! they must get no SIMT region, and compute the same results.
+
+! { dg-do run }
+! { dg-skip-if "tests the default offload mode" { *-*-* } {
"--param=offload-gpu-mode=simt*" } { "" } }
+! { dg-additional-options "-O2" }
+! { dg-additional-options "-foffload-options=-fdump-tree-optimized" { target
offload_target_amdgcn } }
+
+include "simt-collapse-2.f90"
+
+! { dg-final { scan-offload-tree-dump-not "GOMP_SIMT_ENTER" "optimized" {
target offload_target_amdgcn } } }
--
2.56.0
