On Tue, Sep 29, 2026 at 4:44 PM Andrew Stubbs <[email protected]> wrote: > > This updated patch series is my experimental SIMT support for AMD GCN. > > For now, I'd like the target independent parts reviewed (and approved!) > please. > > - [1/9] Add --param offload-gpu-mode > - [3/9] Add TARGET_SIMT_STACK_VARIABLE hooks > - [8/9] timevar.def: Add time tracking for SIMT passes > > The others are provided for context, not yet completely finalized, and > I can self-approve the GCN parts anyway. > > Compared to the first draft series, there's now C++ and Fortran support, > a lot of back-end bug fixes, and I've eliminated most of the SIMD-mode > test regressions. > > SIMT support already exists for NVPTX, and the middle-end code is > reused for GCN, but the backend is a somewhat different problem as the PTX > intermediary assembler language hides a lot of the complexity of > thread-divergence, stack handling, and register allocation, whereas GCN > targets the hardware directly. > > This patch series is enough to build newlib, libgcc, libstdc++, > libgfortran, libgomp, and the BabelStream benchmark so that it mostly > runs correctly (support for reductions remains on the to-do list, but > the placeholder at least does not crash). Other projects are not yet > tested. They might just work, or they might find new edge cases. I'll > get to those in the coming months. The patches expect that the "address > vector" patches have already been applied. > > A new option is added to enable the SIMT mode. It's deliberately a > "param" as I don't intend users to mess with it (once everything is > working properly) unless they're deep into advanced performance tuning. > > --param offload-gpu-mode=simd|simt|simt-directed|default > > What is implemented so far is the "simt-directed" mode in which OpenMP > code must use directives: > > #pragma omp target teams distribute parallel for > for (int i = 0; i < OUTERLIMIT; i++) > { > #pragma omp simd > for (int j = 0; j < INNERLIMIT; j++) > { > // SIMT code here > > This is the same level at which NVPTX enables SIMT, and is roughly the > same as what the SIMD would do (albeit that doesn't actually require any > directive). > > The "simt" mode (no "-directed") is planned to integrate SIMT > parallelization with OpenMP threads so that the "omp simd" directive is > not needed and SIMT is invisible to the end user.
Is it planned to have SIMT vs. non-SIMT switchable per loop (nest) or per function (given OMP will outline the nest anyway)? If you mark the --param as Optimization it would be in the per-function optimization node ... > In this mode, Libgomp > would report that the team has 1024 threads, rather than the current 16, > much like AMD's LLVM implementation does. This next step is mostly a > libgomp task (the backend will not generally need to be aware of it), > and may be made available NVPTX also (possibly). Eventually, I expect > that "simt" will be the default mode (the default remains "simd", for > now). > > Besides the option, I needed to add two new hooks: > TARGET_SIMT_STACK_VARIABLE_PADDING and > TARGET_SIMT_STACK_VARIABLE_OFFSET. These are used to duplicate the > local variables so there is one copy per SIMT thread. > > There's still some TODO and FIXME in the code. Most of these will be > fixed before the final commit. Some are merely nice-to-have and may end > up committed as-is. > > I've broken the new passes out into separate patches for the sake of > clarity, but they won't work in isolation; the whole stack is indended > to be present. > > Thanks in advance for the reviews. > > Andrew > > Andrew Stubbs (9): > Add --param offload-gpu-mode > amdgcn: Add -msimt option and multilib > Add TARGET_SIMT_STACK_VARIABLE hooks > amdgcn: backend SIMT supporting implementation > libgomp, amdgcn: rework team_malloc for SIMT mode > amdgcn: Add gcn-simt-vectorize.cc > amdgcn: Add gcn-simt-cbranch.cc > timevar.def: Add time tracking for SIMT passes > testsuite, amdgcn: testsuite changes > > gcc/cfgexpand.cc | 14 + > gcc/config.gcc | 2 +- > gcc/config/gcn/constraints.md | 3 + > gcc/config/gcn/gcn-builtins.def | 7 + > gcc/config/gcn/gcn-passes.def | 25 + > gcc/config/gcn/gcn-protos.h | 3 + > gcc/config/gcn/gcn-simt-cbranch.cc | 2026 +++++++++++++++ > gcc/config/gcn/gcn-simt-vectorize.cc | 2230 +++++++++++++++++ > gcc/config/gcn/gcn-valu.md | 235 +- > gcc/config/gcn/gcn.cc | 506 +++- > gcc/config/gcn/gcn.h | 10 + > gcc/config/gcn/gcn.md | 585 ++++- > gcc/config/gcn/gcn.opt | 4 + > gcc/config/gcn/mkoffload.cc | 13 + > gcc/config/gcn/predicates.md | 21 +- > gcc/config/gcn/t-gcn-hsa | 14 +- > gcc/doc/params.texi | 24 + > gcc/doc/tm.texi | 17 + > gcc/doc/tm.texi.in | 4 + > gcc/expr.cc | 13 + > gcc/flag-types.h | 9 + > gcc/function.cc | 12 + > gcc/omp-general.cc | 15 +- > gcc/params.opt | 19 + > gcc/target.def | 22 + > .../gcc.target/gcn/avgpr-spill-double.c | 1 + > .../gcc.target/gcn/avgpr-spill-int.c | 1 + > .../gcc.target/gcn/avgpr-spill-long.c | 1 + > .../gcc.target/gcn/avgpr-spill-short.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_1.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_1_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_2.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_2_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_3.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_3_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_4.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_4_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_5.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_5_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_6.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_6_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_7.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_7_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fmaxnm_8.c | 2 + > .../gcc.target/gcn/cond_fmaxnm_8_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_1.c | 2 + > .../gcc.target/gcn/cond_fminnm_1_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_2.c | 2 + > .../gcc.target/gcn/cond_fminnm_2_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_3.c | 2 + > .../gcc.target/gcn/cond_fminnm_3_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_4.c | 2 + > .../gcc.target/gcn/cond_fminnm_4_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_5.c | 2 + > .../gcc.target/gcn/cond_fminnm_5_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_6.c | 2 + > .../gcc.target/gcn/cond_fminnm_6_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_7.c | 2 + > .../gcc.target/gcn/cond_fminnm_7_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_fminnm_8.c | 2 + > .../gcc.target/gcn/cond_fminnm_8_run.c | 2 + > gcc/testsuite/gcc.target/gcn/cond_shift_3.c | 1 + > .../gcc.target/gcn/cond_shift_3_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_shift_4.c | 1 + > .../gcc.target/gcn/cond_shift_4_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_shift_8.c | 1 + > .../gcc.target/gcn/cond_shift_8_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_shift_9.c | 1 + > .../gcc.target/gcn/cond_shift_9_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_smax_1.c | 1 + > .../gcc.target/gcn/cond_smax_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_smin_1.c | 1 + > .../gcc.target/gcn/cond_smin_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_umax_1.c | 1 + > .../gcc.target/gcn/cond_umax_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/cond_umin_1.c | 1 + > .../gcc.target/gcn/cond_umin_1_run.c | 1 + > .../gcc.target/gcn/math-builtins-1.c | 2 + > .../gcc.target/gcn/sgpr-spill-double.c | 5 + > gcc/testsuite/gcc.target/gcn/sgpr-spill-int.c | 5 + > .../gcc.target/gcn/sgpr-spill-long.c | 5 + > .../gcc.target/gcn/sgpr-spill-short.c | 5 + > gcc/testsuite/gcc.target/gcn/simd-math-1.c | 2 + > gcc/testsuite/gcc.target/gcn/simd-math-2.c | 2 + > gcc/testsuite/gcc.target/gcn/simd-math-3-16.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-3-2.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-3-32.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-3-4.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-3-8.c | 1 + > .../gcc.target/gcn/simd-math-3-char-16.c | 1 + > .../gcc.target/gcn/simd-math-3-char-2.c | 1 + > .../gcc.target/gcn/simd-math-3-char-32.c | 1 + > .../gcc.target/gcn/simd-math-3-char-4.c | 1 + > .../gcc.target/gcn/simd-math-3-char-8.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run-16.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run-2.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run-32.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run-4.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run-8.c | 1 + > .../gcc.target/gcn/simd-math-3-char-run.c | 1 + > .../gcc.target/gcn/simd-math-3-char.c | 1 + > .../gcc.target/gcn/simd-math-3-long-16.c | 1 + > .../gcc.target/gcn/simd-math-3-long-2.c | 1 + > .../gcc.target/gcn/simd-math-3-long-32.c | 1 + > .../gcc.target/gcn/simd-math-3-long-4.c | 1 + > .../gcc.target/gcn/simd-math-3-long-8.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run-16.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run-2.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run-32.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run-4.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run-8.c | 1 + > .../gcc.target/gcn/simd-math-3-long-run.c | 1 + > .../gcc.target/gcn/simd-math-3-long.c | 1 + > .../gcc.target/gcn/simd-math-3-run-16.c | 1 + > .../gcc.target/gcn/simd-math-3-run-2.c | 1 + > .../gcc.target/gcn/simd-math-3-run-32.c | 1 + > .../gcc.target/gcn/simd-math-3-run-4.c | 1 + > .../gcc.target/gcn/simd-math-3-run-8.c | 1 + > .../gcc.target/gcn/simd-math-3-run.c | 1 + > .../gcc.target/gcn/simd-math-3-short-16.c | 1 + > .../gcc.target/gcn/simd-math-3-short-2.c | 1 + > .../gcc.target/gcn/simd-math-3-short-32.c | 1 + > .../gcc.target/gcn/simd-math-3-short-4.c | 1 + > .../gcc.target/gcn/simd-math-3-short-8.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run-16.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run-2.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run-32.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run-4.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run-8.c | 1 + > .../gcc.target/gcn/simd-math-3-short-run.c | 1 + > .../gcc.target/gcn/simd-math-3-short.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-3.c | 2 + > .../gcc.target/gcn/simd-math-4-char-run.c | 1 + > .../gcc.target/gcn/simd-math-4-char.c | 1 + > .../gcc.target/gcn/simd-math-4-long-run.c | 1 + > .../gcc.target/gcn/simd-math-4-long.c | 1 + > .../gcc.target/gcn/simd-math-4-run.c | 1 + > .../gcc.target/gcn/simd-math-4-short-run.c | 1 + > .../gcc.target/gcn/simd-math-4-short.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-4.c | 2 + > gcc/testsuite/gcc.target/gcn/simd-math-5-16.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-5-32.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-5-4.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-5-8.c | 1 + > .../gcc.target/gcn/simd-math-5-char-16.c | 1 + > .../gcc.target/gcn/simd-math-5-char-32.c | 1 + > .../gcc.target/gcn/simd-math-5-char-4.c | 1 + > .../gcc.target/gcn/simd-math-5-char-8.c | 1 + > .../gcc.target/gcn/simd-math-5-char-run-16.c | 1 + > .../gcc.target/gcn/simd-math-5-char-run-32.c | 1 + > .../gcc.target/gcn/simd-math-5-char-run-4.c | 1 + > .../gcc.target/gcn/simd-math-5-char-run-8.c | 1 + > .../gcc.target/gcn/simd-math-5-char-run.c | 1 + > .../gcc.target/gcn/simd-math-5-char.c | 1 + > .../gcc.target/gcn/simd-math-5-long-16.c | 1 + > .../gcc.target/gcn/simd-math-5-long-32.c | 1 + > .../gcc.target/gcn/simd-math-5-long-4.c | 1 + > .../gcc.target/gcn/simd-math-5-long-8.c | 1 + > .../gcc.target/gcn/simd-math-5-long-run-16.c | 1 + > .../gcc.target/gcn/simd-math-5-long-run-32.c | 1 + > .../gcc.target/gcn/simd-math-5-long-run-4.c | 1 + > .../gcc.target/gcn/simd-math-5-long-run-8.c | 1 + > .../gcc.target/gcn/simd-math-5-long-run.c | 1 + > .../gcc.target/gcn/simd-math-5-long.c | 1 + > .../gcc.target/gcn/simd-math-5-run-16.c | 1 + > .../gcc.target/gcn/simd-math-5-run-32.c | 1 + > .../gcc.target/gcn/simd-math-5-run-4.c | 1 + > .../gcc.target/gcn/simd-math-5-run-8.c | 1 + > .../gcc.target/gcn/simd-math-5-run.c | 1 + > .../gcc.target/gcn/simd-math-5-short-16.c | 1 + > .../gcc.target/gcn/simd-math-5-short-32.c | 1 + > .../gcc.target/gcn/simd-math-5-short-4.c | 1 + > .../gcc.target/gcn/simd-math-5-short-8.c | 1 + > .../gcc.target/gcn/simd-math-5-short-run-16.c | 1 + > .../gcc.target/gcn/simd-math-5-short-run-32.c | 1 + > .../gcc.target/gcn/simd-math-5-short-run-4.c | 1 + > .../gcc.target/gcn/simd-math-5-short-run-8.c | 1 + > .../gcc.target/gcn/simd-math-5-short-run.c | 1 + > .../gcc.target/gcn/simd-math-5-short.c | 1 + > gcc/testsuite/gcc.target/gcn/simd-math-5.c | 1 + > gcc/testsuite/gcc.target/gcn/simt-asm-g-1.c | 38 + > gcc/testsuite/gcc.target/gcn/simt-asm-r-1.c | 36 + > .../gcc.target/gcn/simt-avgpr-spill-double.c | 9 + > .../gcc.target/gcn/simt-avgpr-spill-int.c | 154 ++ > .../gcc.target/gcn/simt-avgpr-spill-long.c | 9 + > .../gcc.target/gcn/simt-avgpr-spill-run.c | 34 + > .../gcc.target/gcn/simt-avgpr-spill-short.c | 9 + > gcc/testsuite/gcc.target/gcn/simt-cond-1.c | 40 + > gcc/testsuite/gcc.target/gcn/simt-cond-10.c | 42 + > gcc/testsuite/gcc.target/gcn/simt-cond-10b.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-11.c | 45 + > gcc/testsuite/gcc.target/gcn/simt-cond-11b.c | 53 + > gcc/testsuite/gcc.target/gcn/simt-cond-12.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-12b.c | 54 + > gcc/testsuite/gcc.target/gcn/simt-cond-13.c | 41 + > gcc/testsuite/gcc.target/gcn/simt-cond-13b.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-14.c | 44 + > gcc/testsuite/gcc.target/gcn/simt-cond-14b.c | 50 + > gcc/testsuite/gcc.target/gcn/simt-cond-15.c | 44 + > gcc/testsuite/gcc.target/gcn/simt-cond-15b.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-16.c | 52 + > gcc/testsuite/gcc.target/gcn/simt-cond-17.c | 59 + > gcc/testsuite/gcc.target/gcn/simt-cond-18.c | 42 + > gcc/testsuite/gcc.target/gcn/simt-cond-1b.c | 45 + > gcc/testsuite/gcc.target/gcn/simt-cond-2.c | 44 + > gcc/testsuite/gcc.target/gcn/simt-cond-2b.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-3.c | 44 + > gcc/testsuite/gcc.target/gcn/simt-cond-3b.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-4.c | 48 + > gcc/testsuite/gcc.target/gcn/simt-cond-4b.c | 52 + > gcc/testsuite/gcc.target/gcn/simt-cond-5.c | 50 + > gcc/testsuite/gcc.target/gcn/simt-cond-5b.c | 54 + > gcc/testsuite/gcc.target/gcn/simt-cond-6.c | 51 + > gcc/testsuite/gcc.target/gcn/simt-cond-6b.c | 55 + > gcc/testsuite/gcc.target/gcn/simt-cond-7.c | 55 + > gcc/testsuite/gcc.target/gcn/simt-cond-7b.c | 59 + > gcc/testsuite/gcc.target/gcn/simt-cond-8.c | 40 + > gcc/testsuite/gcc.target/gcn/simt-cond-8b.c | 46 + > gcc/testsuite/gcc.target/gcn/simt-cond-9.c | 42 + > gcc/testsuite/gcc.target/gcn/simt-cond-9b.c | 48 + > .../gcc.target/gcn/simt-ensure-scalar-1.c | 65 + > .../gcc.target/gcn/simt-stackparams-1.c | 45 + > .../gcc.target/gcn/simt-stackparams-2.c | 52 + > .../gcc.target/gcn/simt-stackparams-3.c | 67 + > .../gcc.target/gcn/simt-stackvar-1.c | 44 + > .../gcc.target/gcn/simt-stackvar-2.c | 52 + > .../gcc.target/gcn/simt-stackvar-3.c | 59 + > gcc/testsuite/gcc.target/gcn/smax_1.c | 1 + > gcc/testsuite/gcc.target/gcn/smax_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/smin_1.c | 1 + > gcc/testsuite/gcc.target/gcn/smin_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/sram-ecc-3.c | 1 + > gcc/testsuite/gcc.target/gcn/sram-ecc-4.c | 1 + > gcc/testsuite/gcc.target/gcn/sram-ecc-7.c | 1 + > gcc/testsuite/gcc.target/gcn/sram-ecc-8.c | 1 + > gcc/testsuite/gcc.target/gcn/umax_1.c | 1 + > gcc/testsuite/gcc.target/gcn/umax_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/umin_1.c | 1 + > gcc/testsuite/gcc.target/gcn/umin_1_run.c | 1 + > gcc/testsuite/gcc.target/gcn/vcc-clobber.c | 1 + > gcc/testsuite/lib/target-supports.exp | 131 +- > gcc/timevar.def | 1 + > libgcc/config/gcn/amdgcn_veclib.h | 4 + > libgcc/config/gcn/lib2-vec_divmod-di.c | 4 + > libgcc/config/gcn/lib2-vec_divmod-hi.c | 4 + > libgcc/config/gcn/lib2-vec_divmod-qi.c | 4 + > libgcc/config/gcn/lib2-vec_divmod.c | 4 + > libgomp/libgomp.h | 8 + > 248 files changed, 8250 insertions(+), 224 deletions(-) > create mode 100644 gcc/config/gcn/gcn-passes.def > create mode 100644 gcc/config/gcn/gcn-simt-cbranch.cc > create mode 100644 gcc/config/gcn/gcn-simt-vectorize.cc > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-asm-g-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-asm-r-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-avgpr-spill-double.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-avgpr-spill-int.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-avgpr-spill-long.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-avgpr-spill-run.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-avgpr-spill-short.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-10.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-10b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-11.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-11b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-12.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-12b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-13.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-13b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-14.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-14b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-15.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-15b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-16.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-17.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-18.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-1b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-2.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-2b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-3.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-3b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-4.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-4b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-5.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-5b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-6.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-6b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-7.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-7b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-8.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-8b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-9.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-cond-9b.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-ensure-scalar-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackparams-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackparams-2.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackparams-3.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackvar-1.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackvar-2.c > create mode 100644 gcc/testsuite/gcc.target/gcn/simt-stackvar-3.c > > -- > 2.54.0 >
