@@ -38,6 +38,7 @@ along with GCC; see the file COPYING3. If not see
#include "dumpfile.h"
#include "tree-vectorizer.h"
#include "graphite.h"
+#include "graphite-oacc.h"
/* get_schedule_for_node_st - Improve schedule for the schedule node.
@@ -115,6 +116,14 @@ optimize_isl (scop_p scop, bool oacc_enabled_graphite)
int old_err = isl_options_get_on_error (scop->isl_context);
int old_max_operations = isl_ctx_get_max_operations (scop->isl_context);
int max_operations = param_max_isl_operations;
+
+ /* The default value for param_max_isl_operations is easily exceeded
+ by "kernels" loops in existing OpenACC codes. Raise the values
+ significantly since analyzing those loops is crucial. */
+ if (param_max_isl_operations == 350000 /* default value */
+ && oacc_function_p (cfun))
+ max_operations = 2000000;
+
if (max_operations)
isl_ctx_set_max_operations (scop->isl_context, max_operations);
isl_options_set_on_error (scop->isl_context, ISL_ON_ERROR_CONTINUE);
@@ -164,11 +173,27 @@ optimize_isl (scop_p scop, bool oacc_enabled_graphite)
dump_user_location_t loc = find_loop_location
(scop->scop_info->region.entry->dest->loop_father);
if (isl_ctx_last_error (scop->isl_context) == isl_error_quota)
- dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
- "loop nest not optimized, optimization timed out "
- "after %d operations [--param max-isl-operations]\n",
- max_operations);
- else
+ {
+ if (oacc_function_p (cfun))
+ {
+ /* Special casing for OpenACC to unify diagnostic messages
+ here and in graphite-scop-detection.c. */
+ dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
+ "data-dependence analysis of OpenACC loop "
+ "nest "
+ "failed; try increasing the value of "
+ "--param="
+ "max-isl-operations=%d.\n",
+ max_operations);
+ }
+ else
+ dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
+ "loop nest not optimized, optimization timed "
+ "out after %d operations [--param "
+ "max-isl-operations]\n",
+ max_operations);
+ }
+ else
dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
"loop nest not optimized, ISL signalled an error\n");
}
@@ -2053,6 +2053,9 @@ determine_openacc_reductions (scop_p scop)
}
}
+
+extern dump_user_location_t find_loop_location (class loop *);
+
/* Find Static Control Parts (SCoP) in the current function and pushes
them to SCOPS. */
@@ -2106,6 +2109,11 @@ build_scops (vec<scop_p> *scops)
}
unsigned max_arrays = param_graphite_max_arrays_per_scop;
+
+ if (oacc_function_p (cfun)
+ && param_graphite_max_arrays_per_scop == 100 /* default value */)
+ max_arrays = 200;
+
if (max_arrays > 0
&& scop->drs.length () >= max_arrays)
{
@@ -2113,7 +2121,16 @@ build_scops (vec<scop_p> *scops)
<< scop->drs.length ()
<< " is larger than --param graphite-max-arrays-per-scop="
<< max_arrays << ".\n");
- free_scop (scop);
+
+ if (dump_enabled_p () && oacc_function_p (cfun))
+ dump_printf_loc (MSG_MISSED_OPTIMIZATION,
+ find_loop_location (s->entry->dest->loop_father),
+ "data-dependence analysis of OpenACC loop nest "
+ "failed; try increasing the value of --param="
+ "graphite-max-arrays-per-scop=%d.\n",
+ max_arrays);
+
+ free_scop (scop);
continue;
}
@@ -2126,6 +2143,15 @@ build_scops (vec<scop_p> *scops)
<< scop_nb_params (scop)
<< " larger than --param graphite-max-nb-scop-params="
<< max_dim << ".\n");
+
+ if (dump_enabled_p () && oacc_function_p (cfun))
+ dump_printf_loc (MSG_MISSED_OPTIMIZATION,
+ find_loop_location (s->entry->dest->loop_father),
+ "data-dependence analysis of OpenACC loop nest "
+ "failed; try increasing the value of --param="
+ "graphite-max-nb-scop-params=%d.\n",
+ max_dim);
+
free_scop (scop);
continue;
}
new file mode 100644
@@ -0,0 +1,21 @@
+/* Verify that a warning about an exceeded Graphite parameter gets
+ output as optimization information and not only as a dump message
+ for OpenACC functions. */
+
+/* { dg-additional-options "-O2 -fopt-info-missed --param=graphite-max-arrays-per-scop=1" } */
+
+extern int a[1000];
+extern int b[1000];
+
+void test ()
+{
+#pragma acc parallel loop auto
+/* { dg-missed {data-dependence analysis of OpenACC loop nest failed\; try increasing the value of --param=graphite-max-arrays-per-scop=1.} "" { target *-*-* } .-1 } */
+/* { dg-missed {'auto' loop has not been analyzed \(cf. 'graphite' dumps for more information\).} "" { target *-*-* } .-2 } */
+/* { dg-missed {.*not inlinable.*} "" { target *-*-* } .-3 } */
+ for (int i = 1; i < 995; i++)
+ a[i] = b[i + 5] + b[i - 1];
+}
+
+
+/* { dg-prune-output ".*not inlinable.*"} */
new file mode 100644
@@ -0,0 +1,23 @@
+/* Verify that a warning about an exceeded Graphite parameter gets
+ output as optimization information and not only as a dump message
+ for OpenACC functions. */
+
+/* { dg-additional-options "-O2 -fopt-info-missed --param=max-isl-operations=1" } */
+
+void test (int* restrict a, int *restrict b)
+{
+ int i = 1;
+ int j = 1;
+ int m = 0;
+
+#pragma acc parallel loop auto copyin(b) copyout(a) reduction(max:m)
+/* { dg-missed {data-dependence analysis of OpenACC loop nest failed; try increasing the value of --param=max-isl-operations=1.} "" { target *-*-* } .-1 } */
+/* { dg-missed {'auto' loop has not been analyzed \(cf. 'graphite' dumps for more information\).} "" { target *-*-* } .-2 } */
+/* { dg-missed {.*not inlinable.*} "" { target *-*-* } .-3 } */
+ for (i = 1; i < 995; i++)
+ {
+ int x = b[i] * 2;
+ for (j = 1; j < 995; j++)
+ m = m + a[i] + x;
+ }
+}