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;
}
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 =
{
return execute_oacc_gimple_workers ();
}
+
+ opt_pass * clone () { return new pass_oacc_gimple_workers (m_ctxt); }
}; // class pass_oacc_gimple_workers
{}
/* 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
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 */
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);
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.
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);
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);
{ 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. */
/* 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" } } */
{ 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. */
/* 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" } } */
/* { 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. */
/* 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" } } */
/* { 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. */
/* 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" } } */
--- /dev/null
+/* 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" } } */
--- /dev/null
+/* 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" } } */
--- /dev/null
+/* 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" } } */
return x + y;
}
-/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccloops" } } */
+/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccloops*" } } */
-/* { 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);
#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)
/* { 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);
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" } } */
}
-/* Make sure that OpenACC loop processing happens. */
-/* { dg-additional-options "-O2 -fdump-tree-oaccloops" } */
+/* { dg-additional-options "-O2 -fdump-tree-oaccdevlow*" } */
extern int place ();
}
}
-/* { 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*" } } */
/* { 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. */
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" } } */
/* { 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. */
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" } } */
/* { 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. */
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" } } */
/* { 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. */
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" } } */
/* { 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. */
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" } } */
/* { dg-do compile } */
-/* { dg-options "-O -fdump-tree-backprop-details" } */
+/* { dg-options "-O -fdump-tree-backprop1-details" } */
void start (void *);
void end (void *);
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" } } */
/* { dg-do compile } */
-/* { dg-options "-O3 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O3 -fdump-tree-cunrolli1-details" } */
int a[2];
void
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"} } */
/* { dg-do compile } */
-/* { dg-options "-O2 -fdump-tree-cunrolli-details" } */
+/* { dg-options "-O2 -fdump-tree-cunrolli1-details" } */
int a[1];
void
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"} } */
/* { 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];
}
return sum;
}
-/* { dg-final { scan-tree-dump-times "Removed pointless exit:" 1 "cunrolli" } } */
+/* { dg-final { scan-tree-dump-times "Removed pointless exit:" 1 "cunrolli1" } } */
/* { 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
/* { 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);
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" } } */
qz = jl * wh;
}
-/* { dg-final { scan-tree-dump-times "Replacing" 2 "loopdone"} } */
+/* { dg-final { scan-tree-dump-times "Replacing" 2 "loopdone2"} } */
/* { dg-do compile } */
-/* { dg-options "-O -fdump-tree-phiprop-details" } */
+/* { dg-options "-O -fdump-tree-phiprop1-details" } */
struct f
{
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" } } */
/* { 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))
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"} } */
\r
/* { dg-final { scan-tree-dump-times "loop with 3 iterations completely unrolled" 2 "cunroll" } } */\r
/* { dg-final { scan-tree-dump-times "loop with 7 iterations completely unrolled" 2 "cunroll" } } */\r
-/* { dg-final { scan-tree-dump-not "completely unrolled" "cunrolli" } } */\r
+/* { dg-final { scan-tree-dump-not "completely unrolled" "cunrolli1" } } */\r
/* { 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. */
/* { 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()
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" } } */
/* { 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()
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" } } */
/* { 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()
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" } } */
/* 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" } } */
/* { 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. */
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);
/* opt_pass methods: */
virtual unsigned int execute (function *);
+ opt_pass * clone () { return new pass_complete_unroll (m_ctxt); }
}; // class pass_complete_unroll
/* 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
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
/* 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
{
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. */
/* 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
/* 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
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
#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" } */
/* { 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>
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" } */
/* { 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>
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" } */
/* { 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" } */
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" } */
/* { 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>
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" } */
/* { 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>
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" } */
/* { 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>
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" } */
/* { 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>
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" } */
projects / thirdparty / gcc.git / commitdiff
