[OG11][committed][PATCH 02/22] openacc: Move pass_oacc_device_lower after pass_graphite
Frederik Harwath
frederik@codesourcery.com
Wed Nov 17 16:03:11 GMT 2021
The OpenACC device lowering pass must run after the Graphite pass to
allow for the use of Graphite for automatic parallelization of kernels
regions in the future. Experimentation has shown that it is best,
performancewise, to run pass_oacc_device_lower together with the
related passes pass_oacc_loop_designation and pass_oacc_gimple_workers
early after pass_graphite in pass_tree_loop, at least if the other
tree loop passes are not adjusted. In particular, to enable
vectorization which is crucial for GCN offloading, device lowering
should happen before pass_vectorize. To bring the loops contained in
the offloading functions into the shape expected by the loop
vectorizer, we have to make sure that some passes that previously were
executed only once before pass_tree_loop are also executed on the
offloading functions. To ensure the execution of
pass_oacc_device_lower if pass_tree_loop does not execute (no loops,
no optimizations), we introduce two further copies of the pass to the
pipeline that run if there are no loops or if no optimization is
performed.
gcc/ChangeLog:
* omp-general.c (oacc_get_fn_dim_size): Return 0 on
missing "dims".
* omp-offload.c (pass_oacc_loop_designation::clone): New
member function.
(pass_oacc_gimple_workers::clone): Likewise.
(pass_oacc_gimple_device_lower::clone): Likewise.
* passes.c (pass_data_no_loop_optimizations): New pass_data.
(class pass_no_loop_optimizations): New pass.
(make_pass_no_loop_optimizations): New function.
* passes.def: Move pass_oacc_{loop_designation,
gimple_workers, device_lower} into tree_loop, and add
copies to pass_tree_no_loop and to new
pass_no_loop_optimizations. Add copies of passes pass_ccp,
pass_ipa_warn, pass_complete_unrolli, pass_backprop,
pass_phiprop, pass_fix_loops after the OpenACC passes
in pass_tree_loop.
* tree-ssa-loop-ivcanon.c (pass_complete_unroll::clone):
New member function.
(pass_complete_unrolli::clone): Likewise.
* tree-ssa-loop.c (pass_fix_loops::clone): Likewise.
(pass_tree_loop_init::clone): Likewise.
(pass_tree_loop_done::clone): Likewise.
* tree-ssa-phiprop.c (pass_phiprop::clone): Likewise.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Adjust
expected output to pass name changes due to the pass
reordering and cloning.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: Likewise
* testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: Likewise.
gcc/testsuite/ChangeLog:
* gcc.dg/goacc/loop-processing-1.c: Adjust expected output
* to pass name changes due to the pass reordering and cloning.
* c-c++-common/goacc/classify-kernels-unparallelized.c: Likewise.
* c-c++-common/goacc/classify-kernels.c: Likewise.
* c-c++-common/goacc/classify-parallel.c: Likewise.
* c-c++-common/goacc/classify-routine.c: Likewise.
* c-c++-common/goacc/routine-nohost-1.c: Likewise.
* c-c++-common/unroll-1.c: Likewise.
* c-c++-common/unroll-4.c: Likewise.
* gcc.dg/goacc/loop-processing-1.c: Likewise.
* gcc.dg/tree-ssa/backprop-1.c: Likewise.
* gcc.dg/tree-ssa/backprop-2.c: Likewise.
* gcc.dg/tree-ssa/backprop-3.c: Likewise.
* gcc.dg/tree-ssa/backprop-4.c: Likewise.
* gcc.dg/tree-ssa/backprop-5.c: Likewise.
* gcc.dg/tree-ssa/backprop-6.c: Likewise.
* gcc.dg/tree-ssa/cunroll-1.c: Likewise.
* gcc.dg/tree-ssa/cunroll-3.c: Likewise.
* gcc.dg/tree-ssa/cunroll-9.c: Likewise.
* gcc.dg/tree-ssa/ldist-17.c: Likewise.
* gcc.dg/tree-ssa/loop-38.c: Likewise.
* gcc.dg/tree-ssa/pr21463.c: Likewise.
* gcc.dg/tree-ssa/pr45427.c: Likewise.
* gcc.dg/tree-ssa/pr61743-1.c: Likewise.
* gcc.dg/unroll-2.c: Likewise.
* gcc.dg/unroll-3.c: Likewise.
* gcc.dg/unroll-4.c: Likewise.
* gcc.dg/unroll-5.c: Likewise.
* gcc.dg/vect/vect-profile-1.c: Likewise.
* c-c++-common/goacc/device-lowering-debug-optimization.c: New test.
* c-c++-common/goacc/device-lowering-no-loops.c: New test.
* c-c++-common/goacc/device-lowering-no-optimization.c: New test.
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
---
gcc/omp-general.c | 8 +-
gcc/omp-offload.c | 8 ++
gcc/passes.c | 42 ++++++++
gcc/passes.def | 44 ++++++++-
.../goacc/classify-kernels-unparallelized.c | 8 +-
.../c-c++-common/goacc/classify-kernels.c | 8 +-
.../c-c++-common/goacc/classify-parallel.c | 8 +-
.../c-c++-common/goacc/classify-routine.c | 8 +-
.../device-lowering-debug-optimization.c | 29 ++++++
.../goacc/device-lowering-no-loops.c | 17 ++++
.../goacc/device-lowering-no-optimization.c | 30 ++++++
.../c-c++-common/goacc/routine-nohost-1.c | 2 +-
gcc/testsuite/c-c++-common/unroll-1.c | 8 +-
gcc/testsuite/c-c++-common/unroll-4.c | 4 +-
.../gcc.dg/goacc/loop-processing-1.c | 7 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c | 6 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c | 6 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c | 6 +-
gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c | 6 +-
gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c | 2 +-
gcc/testsuite/gcc.dg/tree-ssa/loop-38.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c | 2 +-
gcc/testsuite/gcc.dg/tree-ssa/pr21463.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/pr45427.c | 4 +-
gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c | 2 +-
gcc/testsuite/gcc.dg/unroll-2.c | 2 +-
gcc/testsuite/gcc.dg/unroll-3.c | 4 +-
gcc/testsuite/gcc.dg/unroll-4.c | 4 +-
gcc/testsuite/gcc.dg/unroll-5.c | 4 +-
gcc/testsuite/gcc.dg/vect/bb-slp-59.c | 2 +-
gcc/testsuite/gcc.dg/vect/vect-profile-1.c | 2 +-
gcc/tree-pass.h | 2 +
gcc/tree-ssa-loop-ivcanon.c | 2 +
gcc/tree-ssa-loop.c | 99 +++++++++++++++++++
gcc/tree-ssa-phiprop.c | 2 +
.../libgomp.oacc-c-c++-common/pr85486-2.c | 2 +-
.../vector-length-128-1.c | 4 +-
.../vector-length-128-2.c | 4 +-
.../vector-length-128-3.c | 4 +-
.../vector-length-128-4.c | 4 +-
.../vector-length-128-5.c | 4 +-
.../vector-length-128-6.c | 4 +-
.../vector-length-128-7.c | 4 +-
48 files changed, 360 insertions(+), 86 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c
diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 5b34584ad051..694c14af7b9e 100644
--- a/gcc/omp-general.c
+++ b/gcc/omp-general.c
@@ -2954,7 +2954,13 @@ oacc_get_fn_dim_size (tree fn, int axis)
while (axis--)
dims = TREE_CHAIN (dims);
- int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+ tree v = TREE_VALUE (dims);
+ /* TODO With 'pass_oacc_device_lower' moved "later", this is necessary to
+ avoid ICE for some OpenACC 'kernels' ("parloops") constructs. */
+ if (v == NULL_TREE)
+ return 0;
+
+ int size = TREE_INT_CST_LOW (v);
return size;
}
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index b4592594ee49..bbdcc5207880 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -2594,6 +2594,8 @@ public:
return execute_oacc_loop_designation ();
}
+ opt_pass * clone () { return new pass_oacc_loop_designation (m_ctxt); }
+
}; // class pass_oacc_loop_designation
const pass_data pass_data_oacc_gimple_workers =
@@ -2628,6 +2630,8 @@ public:
{
return execute_oacc_gimple_workers ();
}
+
+ opt_pass * clone () { return new pass_oacc_gimple_workers (m_ctxt); }
}; // class pass_oacc_gimple_workers
@@ -2652,12 +2656,16 @@ public:
{}
/* opt_pass methods: */
+ /* TODO If this were gated on something like '!(fun->curr_properties &
+ PROP_gimple_oaccdevlow)', then we could easily have several instances
+ in the pass pipeline? */
virtual bool gate (function *) { return flag_openacc; };
virtual unsigned int execute (function *)
{
return execute_oacc_device_lower ();
}
+ opt_pass * clone () { return new pass_oacc_device_lower (m_ctxt); }
}; // class pass_oacc_device_lower
diff --git a/gcc/passes.c b/gcc/passes.c
index 64550b00b43c..4a1f4a4b5900 100644
--- a/gcc/passes.c
+++ b/gcc/passes.c
@@ -620,6 +620,48 @@ make_pass_all_optimizations_g (gcc::context *ctxt)
namespace {
+const pass_data pass_data_no_loop_optimizations =
+{
+ GIMPLE_PASS, /* type */
+ "*no_loop_optimizations", /* name */
+ OPTGROUP_NONE, /* optinfo_flags */
+ TV_OPTIMIZE, /* tv_id */
+ 0, /* properties_required */
+ 0, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ 0, /* todo_flags_finish */
+};
+
+/* This pass runs if loop optimizations are disabled
+ at the current optimization level. */
+
+class pass_no_loop_optimizations : public gimple_opt_pass
+{
+public:
+ pass_no_loop_optimizations (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_no_loop_optimizations, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual bool
+ gate (function *)
+ {
+ return !optimize || optimize_debug;
+ }
+
+}; // class pass_no_loop_optimizations
+
+} // anon namespace
+
+static gimple_opt_pass *
+make_pass_no_loop_optimizations (gcc::context *ctxt)
+{
+ return new pass_no_loop_optimizations (ctxt);
+}
+
+namespace {
+
const pass_data pass_data_rest_of_compilation =
{
RTL_PASS, /* type */
diff --git a/gcc/passes.def b/gcc/passes.def
index f6e99ac1f4ed..9220fdc8ca75 100644
--- a/gcc/passes.def
+++ b/gcc/passes.def
@@ -183,9 +183,6 @@ along with GCC; see the file COPYING3. If not see
INSERT_PASSES_AFTER (all_passes)
NEXT_PASS (pass_fixup_cfg);
NEXT_PASS (pass_lower_eh_dispatch);
- NEXT_PASS (pass_oacc_loop_designation);
- NEXT_PASS (pass_oacc_gimple_workers);
- NEXT_PASS (pass_oacc_device_lower);
NEXT_PASS (pass_omp_device_lower);
NEXT_PASS (pass_omp_target_link);
NEXT_PASS (pass_adjust_alignment);
@@ -288,6 +285,35 @@ along with GCC; see the file COPYING3. If not see
POP_INSERT_PASSES ()
NEXT_PASS (pass_parallelize_loops, false /* oacc_kernels_p */);
NEXT_PASS (pass_expand_omp_ssa);
+
+ /* Interrupt pass_tree_loop for OpenACC device lowering. */
+ NEXT_PASS (pass_oacc_only);
+ PUSH_INSERT_PASSES_WITHIN (pass_oacc_only)
+ NEXT_PASS (pass_tree_loop_done);
+ NEXT_PASS (pass_oacc_loop_designation);
+ NEXT_PASS (pass_oacc_gimple_workers);
+ NEXT_PASS (pass_oacc_device_lower);
+
+ NEXT_PASS (pass_oacc_functions_only);
+ PUSH_INSERT_PASSES_WITHIN (pass_oacc_functions_only)
+ /* Repeat some passes on OpenACC functions after device lowering. */
+ /* Lower complex instructions arising from OpenACC
+ reductions. */
+ NEXT_PASS (pass_lower_complex);
+ /* Those passes are necessary here to allow the loop vectorizer to
+ work on the offloading functions which is important for AMD GCN
+ offloading. */
+ NEXT_PASS (pass_ccp, true /* nonzero_p */);
+ NEXT_PASS (pass_complete_unrolli);
+ NEXT_PASS (pass_backprop);
+ NEXT_PASS (pass_phiprop);
+ NEXT_PASS (pass_fix_loops);
+ POP_INSERT_PASSES ()
+
+ /* Continue pass_tree_loop after OpenACC device lowering. */
+ NEXT_PASS (pass_tree_loop_init);
+ POP_INSERT_PASSES ()
+
NEXT_PASS (pass_ch_vect);
NEXT_PASS (pass_if_conversion);
/* pass_vectorize must immediately follow pass_if_conversion.
@@ -307,15 +333,21 @@ along with GCC; see the file COPYING3. If not see
NEXT_PASS (pass_loop_prefetch);
/* Run IVOPTs after the last pass that uses data-reference analysis
as that doesn't handle TARGET_MEM_REFs. */
+
NEXT_PASS (pass_iv_optimize);
NEXT_PASS (pass_lim);
NEXT_PASS (pass_tree_loop_done);
POP_INSERT_PASSES ()
+
+
/* Pass group that runs when pass_tree_loop is disabled or there
are no loops in the function. */
NEXT_PASS (pass_tree_no_loop);
PUSH_INSERT_PASSES_WITHIN (pass_tree_no_loop)
NEXT_PASS (pass_slp_vectorize);
+ NEXT_PASS (pass_oacc_loop_designation);
+ NEXT_PASS (pass_oacc_gimple_workers);
+ NEXT_PASS (pass_oacc_device_lower);
POP_INSERT_PASSES ()
NEXT_PASS (pass_simduid_cleanup);
NEXT_PASS (pass_lower_vector_ssa);
@@ -393,6 +425,12 @@ along with GCC; see the file COPYING3. If not see
NEXT_PASS (pass_local_pure_const);
NEXT_PASS (pass_modref);
POP_INSERT_PASSES ()
+ NEXT_PASS (pass_no_loop_optimizations);
+ PUSH_INSERT_PASSES_WITHIN (pass_no_loop_optimizations)
+ NEXT_PASS (pass_oacc_loop_designation);
+ NEXT_PASS (pass_oacc_gimple_workers);
+ NEXT_PASS (pass_oacc_device_lower);
+ POP_INSERT_PASSES ()
NEXT_PASS (pass_tm_init);
PUSH_INSERT_PASSES_WITHIN (pass_tm_init)
NEXT_PASS (pass_tm_mark);
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
index d8706b9a0a0a..7ce42a469ad3 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels-unparallelized.c
@@ -6,7 +6,7 @@
{ dg-additional-options "-fopt-info-note-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-parloops1-all" }
- { dg-additional-options "-fdump-tree-oaccloops" } */
+ { dg-additional-options "-fdump-tree-oaccloops1" } */
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -40,6 +40,6 @@ void KERNELS ()
/* Check the offloaded function's classification and compute dimensions (will
always be 1 x 1 x 1 for non-offloading compilation).
- { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
+ { dg-final { scan-tree-dump-times "(?n)Function is unparallelized OpenACC kernels offload" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
index e3dc5c01a29b..de7525e67f14 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-kernels.c
@@ -6,7 +6,7 @@
{ dg-additional-options "-fopt-info-note-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
{ dg-additional-options "-fdump-tree-parloops1-all" }
- { dg-additional-options "-fdump-tree-oaccloops" } */
+ { dg-additional-options "-fdump-tree-oaccloops1" } */
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -35,6 +35,6 @@ void KERNELS ()
/* Check the offloaded function's classification and compute dimensions (will
always be 1 x 1 x 1 for non-offloading compilation).
- { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
+ { dg-final { scan-tree-dump-times "(?n)Function is parallelized OpenACC kernels offload" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc kernels parallelized, oacc function \\(, , \\), oacc kernels, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
index 6225a4381dd4..68deb4fdfaf6 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-parallel.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
- { dg-additional-options "-fdump-tree-oaccloops" } */
+ { dg-additional-options "-fdump-tree-oaccloops1" } */
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -27,6 +27,6 @@ void PARALLEL ()
/* Check the offloaded function's classification and compute dimensions (will
always be 1 x 1 x 1 for non-offloading compilation).
- { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops" } } */
+ { dg-final { scan-tree-dump-times "(?n)Function is OpenACC parallel offload" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(1, 1, 1\\), oacc parallel, omp target entrypoint\\)\\)" 1 "oaccloops1" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index 3454771ed92b..dcd2522be1de 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -4,7 +4,7 @@
/* { dg-additional-options "-O2" }
{ dg-additional-options "-fopt-info-optimized-omp" }
{ dg-additional-options "-fdump-tree-ompexp" }
- { dg-additional-options "-fdump-tree-oaccloops" } */
+ { dg-additional-options "-fdump-tree-oaccloops1" } */
/* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
aspects of that functionality. */
@@ -29,6 +29,6 @@ void ROUTINE ()
/* Check the offloaded function's classification and compute dimensions (will
always be 1 x 1 x 1 for non-offloading compilation).
- { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops" } }
- { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccloops" } } */
+ { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccloops1" } }
+ { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccloops1" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c
new file mode 100644
index 000000000000..5bf37cc61580
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-debug-optimization.c
@@ -0,0 +1,29 @@
+/* Verify that OpenACC device lowering executes with "-Og". The actual logic in
+ the test function does not matter. */
+
+/* { dg-additional-options "-Og -fdump-tree-oaccdevlow" } */
+
+int main()
+{
+ int i, j;
+ int ina[1024], out[1024], acc;
+
+ for (j = 0; j < 32; j++)
+ for (i = 0; i < 32; i++)
+ ina[j * 32 + i] = (i == j) ? 2 : 0;
+
+ acc = 0;
+#pragma acc parallel loop copy(acc, ina, out)
+ for (j = 0; j < 32; j++)
+ {
+#pragma acc loop reduction(+:acc)
+ for (i = 0; i < 32; i++)
+ acc += ina[i];
+
+ out[j] = acc;
+ }
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow3" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c
new file mode 100644
index 000000000000..193b5620de1d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-loops.c
@@ -0,0 +1,17 @@
+/* Verify that OpenACC device lowering executes even if there are no OpenACC
+ loops. */
+
+/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow" } */
+
+int main()
+{
+ int x;
+#pragma acc parallel copy(x)
+ {
+ asm volatile("");
+ }
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow2" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c
new file mode 100644
index 000000000000..69e2b22d73ba
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/device-lowering-no-optimization.c
@@ -0,0 +1,30 @@
+/* Verify that OpenACC device lowering executes with "-O0". The actual
+ logic in the test function does not matter. */
+
+/* { dg-additional-options "-O0 -fdump-tree-oaccdevlow" } */
+
+int main()
+{
+
+ int i, j;
+ int ina[1024], out[1024], acc;
+
+ for (j = 0; j < 32; j++)
+ for (i = 0; i < 32; i++)
+ ina[j * 32 + i] = (i == j) ? 2 : 0;
+
+ acc = 0;
+#pragma acc parallel loop copy(acc, ina, out)
+ for (j = 0; j < 32; j++)
+ {
+#pragma acc loop reduction(+:acc)
+ for (i = 0; i < 32; i++)
+ acc += ina[i];
+
+ out[j] = acc;
+ }
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump ".omp_fn" "oaccdevlow3" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
index ebeaadb0b811..480c57feb05f 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -25,4 +25,4 @@ float ADD(float x, float y)
return x + y;
}
-/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccloops" } } */
+/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccloops*" } } */
diff --git a/gcc/testsuite/c-c++-common/unroll-1.c b/gcc/testsuite/c-c++-common/unroll-1.c
index fe7f4f31912c..8e57a44be231 100644
--- a/gcc/testsuite/c-c++-common/unroll-1.c
+++ b/gcc/testsuite/c-c++-common/unroll-1.c
@@ -1,5 +1,5 @@
-/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details -fdump-rtl-loop2_unroll-details" } */
+/* { dg-do compile } *
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details -fdump-rtl-loop2_unroll-details" } */
extern void bar (int);
@@ -10,12 +10,12 @@ void test (void)
#pragma GCC unroll 8
for (unsigned long i = 1; i <= 8; ++i)
bar(i);
- /* { dg-final { scan-tree-dump "11:.*: loop with 8 iterations completely unrolled" "cunrolli" } } */
+ /* { dg-final { scan-tree-dump "11:.*: loop with 8 iterations completely unrolled" "cunrolli1" } } */
#pragma GCC unroll 8
for (unsigned long i = 1; i <= 7; ++i)
bar(i);
- /* { dg-final { scan-tree-dump "16:.*: loop with 7 iterations completely unrolled" "cunrolli" } } */
+ /* { dg-final { scan-tree-dump "16:.*: loop with 7 iterations completely unrolled" "cunrolli1" } } */
#pragma GCC unroll 8
for (unsigned long i = 1; i <= 15; ++i)
diff --git a/gcc/testsuite/c-c++-common/unroll-4.c b/gcc/testsuite/c-c++-common/unroll-4.c
index 1c1988174ba7..fe7f9e10626e 100644
--- a/gcc/testsuite/c-c++-common/unroll-4.c
+++ b/gcc/testsuite/c-c++-common/unroll-4.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -funroll-all-loops -fdump-rtl-loop2_unroll-details -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O2 -funroll-all-loops -fdump-rtl-loop2_unroll-details -fdump-tree-cunrolli1-details" } */
extern void bar (int);
@@ -17,6 +17,6 @@ void test (void)
for (unsigned long i = 1; i <= j; ++i)
bar(i);
- /* { dg-final { scan-tree-dump "Not unrolling loop .: user didn't want it unrolled completely" "cunrolli" } } */
+ /* { dg-final { scan-tree-dump "Not unrolling loop .: user didn't want it unrolled completely" "cunrolli1" } } */
/* { dg-final { scan-rtl-dump-times "Not unrolling loop, user didn't want it unrolled" 2 "loop2_unroll" } } */
}
diff --git a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
index 7533c4fe0e88..6979cce71b05 100644
--- a/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
+++ b/gcc/testsuite/gcc.dg/goacc/loop-processing-1.c
@@ -1,5 +1,4 @@
-/* Make sure that OpenACC loop processing happens. */
-/* { dg-additional-options "-O2 -fdump-tree-oaccloops" } */
+/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow*" } */
extern int place ();
@@ -15,5 +14,5 @@ void vector_1 (int *ary, int size)
}
}
-/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop [0-9]{2}\(1\).*.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 68\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 68\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccloops" } } */
-
+/* { dg-final { scan-tree-dump {
+OpenACC loops.*Loop 0\(0\).*Loop [0-9]{2}\(1\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 1, 36\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 0\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 0\);.*Loop 6\(6\).*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*Head-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, 0, 2, 6\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 1\);.*Head-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_HEAD_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_FORK, \.data_dep\.[0-9_]+, 2\);.*Tail-1:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 2\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 2\);.*Tail-0:.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_TAIL_MARK, \.data_dep\.[0-9_]+, 1\);.*\.data_dep\.[0-9_]+ = \.UNIQUE \(OACC_JOIN, \.data_dep\.[0-9_]+, 1\);} "oaccdevlow*" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c
index 302fdb570b63..b6b11bf30afa 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-1.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -g -fdump-tree-backprop-details" } */
+/* { dg-options "-O -g -fdump-tree-backprop1-details" } */
/* Test a simple case of non-looping code in which both uses ignore
the sign and both definitions are sign ops. */
@@ -18,5 +18,5 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 3 "backprop" } } */
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = ABS_EXPR <x} 3 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 3 "backprop1" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = ABS_EXPR <x} 3 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c
index d54fd36e2fb3..bef921be500b 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-2.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -g -fdump-tree-backprop-details" } */
+/* { dg-options "-O -g -fdump-tree-backprop1-details" } */
/* Test a simple case of non-looping code in which both uses ignore
the sign but only one definition is a sign op. */
@@ -18,4 +18,4 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 3 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 3 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c
index a244b4af2ac2..1b76ce05cbef 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-3.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -g -fdump-tree-backprop-details" } */
+/* { dg-options "-O -g -fdump-tree-backprop1-details" } */
/* Test a simple case of non-looping code in which one use ignores
the sign but another doesn't. */
@@ -18,4 +18,4 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 0 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -x} 0 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c
index 54355009c744..02223fd9f23b 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-4.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -g -fdump-tree-backprop-details" } */
+/* { dg-options "-O -g -fdump-tree-backprop1-details" } */
/* Test a simple reduction loop in which all inputs are sign ops and
the consumer doesn't care about the sign. */
@@ -17,5 +17,5 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = __builtin_copysign} 3 "backprop" } } */
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -} 3 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = __builtin_copysign} 3 "backprop1" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -} 3 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c
index e4f0f856ff6b..9dd04408b3a8 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-5.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -g -fdump-tree-backprop-details" } */
+/* { dg-options "-O -g -fdump-tree-backprop1-details" } */
/* Test a loop that does both a multiplication and addition. The addition
should prevent any sign ops from being removed. */
@@ -17,4 +17,4 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = __builtin_copysign} 0 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = __builtin_copysign} 0 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c b/gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c
index 31f05716f149..1d17c7328036 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/backprop-6.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -fdump-tree-backprop-details" } */
+/* { dg-options "-O -fdump-tree-backprop1-details" } */
void start (void *);
void end (void *);
@@ -26,5 +26,5 @@ TEST_FUNCTION (float, f)
TEST_FUNCTION (double, )
TEST_FUNCTION (long double, l)
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -} 6 "backprop" } } */
-/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = ABS_EXPR <} 3 "backprop" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = -} 6 "backprop1" } } */
+/* { dg-final { scan-tree-dump-times {Deleting[^\n]* = ABS_EXPR <} 3 "backprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c
index bcafbfe86b50..110c6cd8635e 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-1.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O3 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O3 -fdump-tree-cunrolli1-details" } */
int a[2];
void
test(int c)
@@ -9,5 +9,5 @@ test(int c)
a[i]=5;
}
/* Array bounds says the loop will not roll much. */
-/* { dg-final { scan-tree-dump "loop with 2 iterations completely unrolled" "cunrolli"} } */
-/* { dg-final { scan-tree-dump "Last iteration exit edge was proved true." "cunrolli"} } */
+/* { dg-final { scan-tree-dump "loop with 2 iterations completely unrolled" "cunrolli1"} } */
+/* { dg-final { scan-tree-dump "Last iteration exit edge was proved true." "cunrolli1"} } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c
index e25c638ac514..f8ab47cebf08 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-3.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details" } */
int a[1];
void
test(int c)
@@ -12,4 +12,4 @@ test(int c)
}
/* If we start duplicating headers prior curoll, this loop will have 0 iterations. */
-/* { dg-final { scan-tree-dump "loop with 1 iterations completely unrolled" "cunrolli"} } */
+/* { dg-final { scan-tree-dump "loop with 1 iterations completely unrolled" "cunrolli1"} } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c
index 886dc147ad1a..f93db92ab384 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/cunroll-9.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details -fdisable-tree-evrp" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details -fdisable-tree-evrp" } */
void abort (void);
int q (void);
int a[10];
@@ -20,4 +20,4 @@ t (int n)
}
return sum;
}
-/* { dg-final { scan-tree-dump-times "Removed pointless exit:" 1 "cunrolli" } } */
+/* { dg-final { scan-tree-dump-times "Removed pointless exit:" 1 "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c b/gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c
index b3617f685a1d..86c84606ce51 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/ldist-17.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -ftree-loop-distribution -ftree-loop-distribute-patterns -fdump-tree-ldist-details -fdisable-tree-cunroll -fdisable-tree-cunrolli" } */
+/* { dg-options "-O2 -ftree-loop-distribution -ftree-loop-distribute-patterns -fdump-tree-ldist-details -fdisable-tree-cunroll -fdisable-tree-cunrolli1" } */
typedef int mad_fixed_t;
struct mad_pcm
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/loop-38.c b/gcc/testsuite/gcc.dg/tree-ssa/loop-38.c
index 7ca1e4709751..f8f04ffaa456 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/loop-38.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/loop-38.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details" } */
int a[10];
int b[11];
int q (void);
@@ -15,4 +15,4 @@ t(int n)
sum+=b[i];
return sum;
}
-/* { dg-final { scan-tree-dump "Loop 1 iterates at most 11 times" "cunrolli" } } */
+/* { dg-final { scan-tree-dump "Loop 1 iterates at most 11 times" "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c b/gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c
index d71b757fbca5..482c19ea1485 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/loopclosedphi.c
@@ -18,4 +18,4 @@ t6 (int qz, int wh)
qz = jl * wh;
}
-/* { dg-final { scan-tree-dump-times "Replacing" 2 "loopdone"} } */
+/* { dg-final { scan-tree-dump-times "Replacing" 2 "loopdone2"} } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/pr21463.c b/gcc/testsuite/gcc.dg/tree-ssa/pr21463.c
index ed0829a038c4..c6f1226d6834 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/pr21463.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/pr21463.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O -fdump-tree-phiprop-details" } */
+/* { dg-options "-O -fdump-tree-phiprop1-details" } */
struct f
{
@@ -16,4 +16,4 @@ int g(int i, int c, struct f *ff, int g)
return *t;
}
-/* { dg-final { scan-tree-dump-times "Inserting PHI for result of load" 1 "phiprop" } } */
+/* { dg-final { scan-tree-dump-times "Inserting PHI for result of load" 1 "phiprop1" } } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/pr45427.c b/gcc/testsuite/gcc.dg/tree-ssa/pr45427.c
index 2f86f02a30ce..3e8a13cd40c0 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/pr45427.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/pr45427.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details" } */
extern void abort (void);
int __attribute__((noinline,noclone))
@@ -25,4 +25,4 @@ int main()
return 0;
}
-/* { dg-final { scan-tree-dump-times "bounded by 0x0\[^0-9a-f\]" 0 "cunrolli"} } */
+/* { dg-final { scan-tree-dump-times "bounded by 0x0\[^0-9a-f\]" 0 "cunrolli1"} } */
diff --git a/gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c b/gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c
index 669d357045cb..069df138bcbe 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/pr61743-1.c
@@ -50,4 +50,4 @@ int foo1 (e_u8 a[4][N], int b1, int b2, e_u8 b[M+1][4][N])
/* { dg-final { scan-tree-dump-times "loop with 3 iterations completely unrolled" 2 "cunroll" } } */
/* { dg-final { scan-tree-dump-times "loop with 7 iterations completely unrolled" 2 "cunroll" } } */
-/* { dg-final { scan-tree-dump-not "completely unrolled" "cunrolli" } } */
+/* { dg-final { scan-tree-dump-not "completely unrolled" "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/unroll-2.c b/gcc/testsuite/gcc.dg/unroll-2.c
index 8baceaac1699..f94174f0f1d3 100644
--- a/gcc/testsuite/gcc.dg/unroll-2.c
+++ b/gcc/testsuite/gcc.dg/unroll-2.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details=stderr -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli-details=stderr -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli1" } */
/* Blank lines can occur in the output of
-fdump-tree-cunrolli-details=stderr. */
diff --git a/gcc/testsuite/gcc.dg/unroll-3.c b/gcc/testsuite/gcc.dg/unroll-3.c
index 10bf59b9a2e7..0284378b9c5c 100644
--- a/gcc/testsuite/gcc.dg/unroll-3.c
+++ b/gcc/testsuite/gcc.dg/unroll-3.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunrolli=foo -fenable-tree-cunrolli=foo" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunrolli1=foo -fenable-tree-cunrolli1=foo" } */
unsigned a[100], b[100];
inline void bar()
@@ -28,4 +28,4 @@ int foo2(void)
return 1;
}
-/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli" } } */
+/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/unroll-4.c b/gcc/testsuite/gcc.dg/unroll-4.c
index 17f194212279..d62e2e7afa0a 100644
--- a/gcc/testsuite/gcc.dg/unroll-4.c
+++ b/gcc/testsuite/gcc.dg/unroll-4.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli=foo -fdisable-tree-cunrolli=foo2" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli1=foo -fdisable-tree-cunrolli1=foo2" } */
unsigned a[100], b[100];
inline void bar()
@@ -28,4 +28,4 @@ int foo2(void)
return 1;
}
-/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli" } } */
+/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/unroll-5.c b/gcc/testsuite/gcc.dg/unroll-5.c
index f3bdebe9882f..c81467cd4202 100644
--- a/gcc/testsuite/gcc.dg/unroll-5.c
+++ b/gcc/testsuite/gcc.dg/unroll-5.c
@@ -1,5 +1,5 @@
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli=foo2 -fdisable-tree-cunrolli=foo" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details -fno-peel-loops -fno-tree-vrp -fdisable-tree-cunroll -fenable-tree-cunrolli1=foo2 -fdisable-tree-cunrolli1=foo" } */
unsigned a[100], b[100];
inline void bar()
@@ -28,4 +28,4 @@ int foo2(void)
return 1;
}
-/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli" } } */
+/* { dg-final { scan-tree-dump-times "loop with 2 iterations completely unrolled" 1 "cunrolli1" } } */
diff --git a/gcc/testsuite/gcc.dg/vect/bb-slp-59.c b/gcc/testsuite/gcc.dg/vect/bb-slp-59.c
index 815b44e1f7cf..2f7c17d803eb 100644
--- a/gcc/testsuite/gcc.dg/vect/bb-slp-59.c
+++ b/gcc/testsuite/gcc.dg/vect/bb-slp-59.c
@@ -22,5 +22,5 @@ void foo (void)
/* We should be able to vectorize the cycle in one SLP attempt including
both load groups and do only one permutation. */
/* { dg-final { scan-tree-dump-times "transform load" 2 "slp1" } } */
-/* { dg-final { scan-tree-dump-times "VEC_PERM_EXPR" 1 "loopdone" } } */
+/* { dg-final { scan-tree-dump-times "VEC_PERM_EXPR" 1 "loopdone2" } } */
/* { dg-final { scan-tree-dump-times "optimized: basic block" 1 "slp1" } } */
diff --git a/gcc/testsuite/gcc.dg/vect/vect-profile-1.c b/gcc/testsuite/gcc.dg/vect/vect-profile-1.c
index 922f965806f9..a8b3ffb87d06 100644
--- a/gcc/testsuite/gcc.dg/vect/vect-profile-1.c
+++ b/gcc/testsuite/gcc.dg/vect/vect-profile-1.c
@@ -1,6 +1,6 @@
/* { dg-do compile } */
/* { dg-require-effective-target vect_int } */
-/* { dg-additional-options "-fdump-tree-vect-details-blocks -fdisable-tree-cunrolli" } */
+/* { dg-additional-options "-fdump-tree-vect-details-blocks -fdisable-tree-cunrolli1" } */
/* At least one of these should correspond to a full vector. */
diff --git a/gcc/tree-pass.h b/gcc/tree-pass.h
index 4a575a54c045..5ffd2799ac2c 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -482,6 +482,8 @@ extern gimple_opt_pass *make_pass_vtable_verify (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_ubsan (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_sanopt (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_oacc_kernels (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_oacc_only (gcc::context *ctxt);
+extern gimple_opt_pass *make_pass_oacc_functions_only (gcc::context *ctxt);
extern simple_ipa_opt_pass *make_pass_ipa_oacc (gcc::context *ctxt);
extern simple_ipa_opt_pass *make_pass_ipa_oacc_kernels (gcc::context *ctxt);
extern gimple_opt_pass *make_pass_warn_nonnull_compare (gcc::context *ctxt);
diff --git a/gcc/tree-ssa-loop-ivcanon.c b/gcc/tree-ssa-loop-ivcanon.c
index 3f9e9d0869f2..430c0520736a 100644
--- a/gcc/tree-ssa-loop-ivcanon.c
+++ b/gcc/tree-ssa-loop-ivcanon.c
@@ -1587,6 +1587,7 @@ public:
/* opt_pass methods: */
virtual unsigned int execute (function *);
+ opt_pass * clone () { return new pass_complete_unroll (m_ctxt); }
}; // class pass_complete_unroll
@@ -1646,6 +1647,7 @@ public:
/* opt_pass methods: */
virtual bool gate (function *) { return optimize >= 2; }
virtual unsigned int execute (function *);
+ opt_pass * clone () { return new pass_complete_unrolli (m_ctxt); }
}; // class pass_complete_unrolli
diff --git a/gcc/tree-ssa-loop.c b/gcc/tree-ssa-loop.c
index 957ac0f3baab..21961200db66 100644
--- a/gcc/tree-ssa-loop.c
+++ b/gcc/tree-ssa-loop.c
@@ -70,6 +70,8 @@ public:
virtual bool gate (function *) { return flag_tree_loop_optimize; }
virtual unsigned int execute (function *fn);
+
+ opt_pass * clone () { return new pass_fix_loops (m_ctxt); }
}; // class pass_fix_loops
unsigned int
@@ -136,6 +138,8 @@ public:
/* opt_pass methods: */
virtual bool gate (function *fn) { return gate_loop (fn); }
+
+ opt_pass * clone () { return new pass_tree_loop (m_ctxt); }
}; // class pass_tree_loop
} // anon namespace
@@ -201,6 +205,97 @@ make_pass_oacc_kernels (gcc::context *ctxt)
{
return new pass_oacc_kernels (ctxt);
}
+/* A superpass that runs its subpasses on OpenACC functions only. */
+
+namespace {
+
+const pass_data pass_data_oacc_functions_only =
+{
+ GIMPLE_PASS, /* type */
+ "*oacc_fns_only", /* name */
+ OPTGROUP_LOOP, /* optinfo_flags */
+ TV_TREE_LOOP, /* tv_id */
+ 0, /* properties_required */
+ 0, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ 0, /* todo_flags_finish */
+};
+
+class pass_oacc_functions_only: public gimple_opt_pass
+{
+public:
+ pass_oacc_functions_only (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_oacc_functions_only, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual bool gate (function *fn) {
+ if (!flag_openacc)
+ return false;
+
+ if (!oacc_get_fn_attrib (fn->decl))
+ return false;
+
+ return true;
+ }
+
+}; // class pass_oacc_functions_only
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_oacc_functions_only (gcc::context *ctxt)
+{
+ return new pass_oacc_functions_only (ctxt);
+}
+
+/* A superpass that runs its subpasses only if compiling for OpenACC. */
+
+namespace {
+
+const pass_data pass_data_oacc_only =
+{
+ GIMPLE_PASS, /* type */
+ "*oacc_only", /* name */
+ OPTGROUP_LOOP, /* optinfo_flags */
+ TV_TREE_LOOP, /* tv_id */
+ 0, /* properties_required */
+ 0, /* properties_provided */
+ 0, /* properties_destroyed */
+ 0, /* todo_flags_start */
+ 0, /* todo_flags_finish */
+};
+
+class pass_oacc_only: public gimple_opt_pass
+{
+public:
+ pass_oacc_only (gcc::context *ctxt)
+ : gimple_opt_pass (pass_data_oacc_only, ctxt)
+ {}
+
+ /* opt_pass methods: */
+ virtual bool gate (function *fn) {
+ if (!flag_openacc)
+ return false;
+
+ if (!oacc_get_fn_attrib (fn->decl))
+ return false;
+
+ return true;
+ }
+
+}; // class pass_oacc_only
+
+} // anon namespace
+
+gimple_opt_pass *
+make_pass_oacc_only (gcc::context *ctxt)
+{
+ return new pass_oacc_only (ctxt);
+}
+
+
/* The ipa oacc superpass. */
@@ -344,6 +439,8 @@ public:
/* opt_pass methods: */
virtual unsigned int execute (function *);
+ opt_pass * clone () { return new pass_tree_loop_init (m_ctxt); }
+
}; // class pass_tree_loop_init
unsigned int
@@ -558,6 +655,8 @@ public:
/* opt_pass methods: */
virtual unsigned int execute (function *) { return tree_ssa_loop_done (); }
+ opt_pass * clone () { return new pass_tree_loop_done (m_ctxt); }
+
}; // class pass_tree_loop_done
} // anon namespace
diff --git a/gcc/tree-ssa-phiprop.c b/gcc/tree-ssa-phiprop.c
index 64d6eda5f6c2..2b727ed0d013 100644
--- a/gcc/tree-ssa-phiprop.c
+++ b/gcc/tree-ssa-phiprop.c
@@ -479,6 +479,8 @@ public:
virtual bool gate (function *) { return flag_tree_phiprop; }
virtual unsigned int execute (function *);
+ opt_pass * clone () { return new pass_phiprop (m_ctxt); }
+
}; // class pass_phiprop
unsigned int
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
index d45326488cd8..bc55d158a81f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85486-2.c
@@ -7,5 +7,5 @@
#include "pr85486.c"
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c
index 18d77cc5ecb1..22891a243e14 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c
@@ -1,5 +1,5 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -34,5 +34,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c
index 8b5b2a4a92d5..30418f378f93 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c
@@ -1,6 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-fopenacc-dim=::128" } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -35,5 +35,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=128" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
index 59be37a7c27e..754964d60100 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c
@@ -1,5 +1,5 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* We default to warp size 32 for the vector length, so the GOMP_OPENACC_DIM has
no effect. */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" "::128" } */
@@ -38,5 +38,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=1, vectors=32" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c
index e5d1df09b8a3..44364cbc51a7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c
@@ -1,5 +1,5 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -36,5 +36,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c
index e60f1c28db4a..5e387c6ced61 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c
@@ -1,6 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-fopenacc-dim=:2:128" } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -37,5 +37,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c
index a1f67622f84d..d32f4e4417ab 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c
@@ -1,6 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-set-target-env-var "GOMP_OPENACC_DIM" ":2:" } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -37,5 +37,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=2, vectors=128" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c
index c419f6499b53..df5cb09df712 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c
@@ -1,5 +1,5 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
-/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
+/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow1" } */
/* { dg-set-target-env-var "GOMP_DEBUG" "1" } */
#include <stdlib.h>
@@ -36,5 +36,5 @@ main (void)
return 0;
}
-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccdevlow1" } } */
/* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */
--
2.33.0
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955
More information about the Gcc-patches
mailing list