[14/40] openacc: Move pass_oacc_device_lower after pass_graphite

diff --git a/gcc/omp-general.c b/gcc/omp-general.c
index 445275524134..27a1bc8092c8 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-oacc-neuter-broadcast.cc b/gcc/omp-oacc-neuter-broadcast.cc
index e43338f3abf2..94ecdc4d4e9a 100644
--- a/gcc/omp-oacc-neuter-broadcast.cc
+++ b/gcc/omp-oacc-neuter-broadcast.cc
@@ -1992,6 +1992,8 @@  public:
       return execute_omp_oacc_neuter_broadcast ();
     }

+  opt_pass * clone () { return new pass_omp_oacc_neuter_broadcast (m_ctxt); }
+
 }; // class pass_omp_oacc_neuter_broadcast

 } // anon namespace
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 833f7ddea58f..e99aaac0e515 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -2444,6 +2444,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_device_lower =
@@ -2467,12 +2469,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 0f541454e7f1..5b9bb422d281 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_omp_oacc_neuter_broadcast);
-  NEXT_PASS (pass_oacc_device_lower);
   NEXT_PASS (pass_omp_device_lower);
   NEXT_PASS (pass_omp_target_link);
   NEXT_PASS (pass_adjust_alignment);
@@ -292,6 +289,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_omp_oacc_neuter_broadcast);
+           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.
@@ -311,15 +337,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_omp_oacc_neuter_broadcast);
+         NEXT_PASS (pass_oacc_device_lower);
       POP_INSERT_PASSES ()
       NEXT_PASS (pass_simduid_cleanup);
       NEXT_PASS (pass_lower_vector_ssa);
@@ -397,6 +429,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_omp_oacc_neuter_broadcast);
+      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 e391184f403d..338676aa20ff 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-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.  */
@@ -39,6 +39,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 779e2b0a24db..37e2a57455d1 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-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 9056aa69dad6..82b70ae280cd 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 f7f0454009bf..cd539370dbbf 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,14 +29,14 @@  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)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccloops" { target c } } }
-   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
-   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
-   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccloops" { target c } } }
-   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && { ! offloading_enabled } } } } }
-   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccloops" { target { c++ && offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccloops1" } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' doesn't have 'nohost' clause" 1 "oaccloops1" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops1" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' doesn't have 'nohost' clause" 1 "oaccloops1" { target { c++ && offloading_enabled } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE' not discarded" 1 "oaccloops1" { target c } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'void ROUTINE\\(\\)' not discarded" 1 "oaccloops1" { target { c++ && { ! offloading_enabled } } } } }
+   { dg-final { scan-tree-dump-times "(?n)OpenACC routine 'ROUTINE\\(\\)' not discarded" 1 "oaccloops1" { target { c++ && offloading_enabled } } } }
    TODO See PR101551 for 'offloading_enabled' differences.
-   { 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)void ROUTINE \\(\\)" 1 "oaccloops" } } */
+   { 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" } }
+   { dg-final { scan-tree-dump-times "(?n)void ROUTINE \\(\\)" 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 59ebb2bc5a9f..4f9a3a333570 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-nohost-1.c
@@ -13,7 +13,7 @@  int THREE(void)
 #pragma acc routine nohost
 extern int THREE(void);

-/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*THREE[^']*' has 'nohost' clause\.$} 1 "oaccloops*" } } */


 #pragma acc routine nohost
@@ -30,7 +30,7 @@  extern void NOTHING(void);

 #pragma acc routine (NOTHING) nohost

-/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*NOTHING[^']*' has 'nohost' clause\.$} 1 "oaccloops*" } } */


 extern float ADD(float, float);
@@ -47,4 +47,4 @@  extern float ADD(float, float);

 #pragma acc routine (ADD) nohost

-/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 oaccloops } } */
+/* { dg-final { scan-tree-dump-times {(?n)^OpenACC routine '[^']*ADD[^']*' has 'nohost' clause\.$} 1 "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 78b9aed89beb..c191125b7951 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,4 +14,4 @@  void vector_1 (int *ary, int size)
   }
 }

-/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(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\);} "oaccloops" } } */
+/* { dg-final { scan-tree-dump {OpenACC loops.*Loop 0\(0\).*Loop 24\(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\);} "oaccloops*" } } */
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/pr59597.c b/gcc/testsuite/gcc.dg/tree-ssa/pr59597.c
index 0f66aae87bba..98d639bc24dd 100644
--- a/gcc/testsuite/gcc.dg/tree-ssa/pr59597.c
+++ b/gcc/testsuite/gcc.dg/tree-ssa/pr59597.c
@@ -1,5 +1,5 @@ 
 /* { dg-do compile } */
-/* { dg-options "-Ofast -fdisable-tree-cunrolli -fdump-tree-threadfull1-details" } */
+/* { dg-options "-Ofast -fdisable-tree-cunrolli1 -fdump-tree-threadfull1-details" } */

 typedef unsigned short u16;
 typedef unsigned char u8;
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 e807ad855efd..ebaa3c86694f 100644
--- a/gcc/tree-pass.h
+++ b/gcc/tree-pass.h
@@ -489,6 +489,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 be533b03a85b..2d4145ce5b8e 100644
--- a/gcc/tree-ssa-loop-ivcanon.c
+++ b/gcc/tree-ssa-loop-ivcanon.c
@@ -1583,6 +1583,7 @@  public:

   /* opt_pass methods: */
   virtual unsigned int execute (function *);
+  opt_pass * clone () { return new pass_complete_unroll (m_ctxt); }

 }; // class pass_complete_unroll

@@ -1642,6 +1643,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 1bbf2f1fb2c8..8d5572033f7b 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
@@ -200,6 +204,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.  */

@@ -343,6 +438,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
@@ -556,6 +653,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 78b0461c839d..f138f766286b 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 17cc9bd663e5..4438f6c24fed 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\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops1" } } */
 /* { 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 5158bb5eb89e..c0a29c7556f9 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
@@ -34,5 +34,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops*" } } */
 /* { 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 a3e44ebfbcb4..326f6d8dc31a 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,7 @@ 
 /* { dg-do run { target openacc_nvidia_accel_selected } } */
 /* { dg-additional-options "-fopenacc-dim=::128" } */
 /* { dg-additional-options "-foffload=-fdump-tree-oaccloops" } */
+>>>>>>> adfd567486a0 (Move pass_oacc_device_lower after pass_graphite)
 /* { dg-set-target-env-var "GOMP_DEBUG" "1" } */

 #include <stdlib.h>
@@ -35,5 +36,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 128\\)" "oaccloops*" } } */
 /* { 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 a85400d09c50..efc9297acdee 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
@@ -38,5 +38,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 1, 32\\)" "oaccloops*" } } */
 /* { 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 24c078f377c3..1c83ec0cc18d 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
@@ -36,5 +36,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops*" } } */
 /* { 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 fcca9f593bb2..f2391dca7272 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
@@ -37,5 +37,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 2, 128\\)" "oaccloops*" } } */
 /* { 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 0807eab7eee4..8ddaaf592cc1 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
@@ -37,5 +37,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops*" } } */
 /* { 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 4a8c1bf549e9..97abbfc20986 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
@@ -36,5 +36,5 @@  main (void)
   return 0;
 }

-/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops" } } */
+/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function \\(1, 0, 128\\)" "oaccloops*" } } */
 /* { dg-output "nvptx_exec: kernel main\\\$_omp_fn\\\$0: launch gangs=1, workers=8, vectors=128" } */