This patch series is a first draft of 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 and are not yet finalized.
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, 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, as well as libstdc++, libgortran, etc. 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. 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
P.S. I will be giving a talk about this project at the Cauldron.
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 | 1979 +++++++++++++++
gcc/config/gcn/gcn-simt-vectorize.cc | 2225 +++++++++++++++++
gcc/config/gcn/gcn-valu.md | 189 +-
gcc/config/gcn/gcn.cc | 504 +++-
gcc/config/gcn/gcn.h | 10 +
gcc/config/gcn/gcn.md | 565 ++++-
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 | 6 +
.../gcc.target/gcn/avgpr-spill-long.c | 1 +
.../gcc.target/gcn/avgpr-spill-short.c | 10 +-
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 | 2 +
gcc/testsuite/gcc.target/gcn/sgpr-spill-int.c | 2 +
.../gcc.target/gcn/sgpr-spill-long.c | 2 +
.../gcc.target/gcn/sgpr-spill-short.c | 2 +
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, 8129 insertions(+), 227 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
