graphite: Tune parameters for OpenACC use

The default values of some parameters that restrict Graphite's
resource usage are too low for many OpenACC codes.  Furthermore,
exceeding the limits does not alwas lead to user-visible diagnostic
messages.

This commit increases the parameter values on OpenACC functions.  The
values were chosen to allow for the analysis of all "kernels" regions
in the SPEC ACCEL v1.3 benchmark suite.  Warnings about exceeded
Graphite-related limits are added to the -fopt-info-missed
output. Those warnings are phrased in a uniform way that intentionally
refers to the "data-dependence analysis" of "OpenACC loops" instead of
"a failure in Graphite" to make them easier to understand for users.

gcc/ChangeLog:

        * graphite-optimize-isl.cc (optimize_isl): Adjust
	param_max_isl_operations value for OpenACC functions and add
	special warnings if value gets exceeded.

	* graphite-scop-detection.cc (build_scops): Likewise for
	param_graphite_max_arrays_per_scop.

gcc/testsuite/ChangeLog:

        * gcc.dg/goacc/graphite-parameter-1.c: New test.
        * gcc.dg/goacc/graphite-parameter-2.c: New test.

diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 739c1b9148a7efccf87c095d399ccab74af29fbb..955e3a2222f76fa98d60626dcc9ec1f114beaa9e 100644 (file)
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,12 @@
+2021-11-16  Frederik Harwath  <frederik@codesourcery.com>
+
+       * graphite-optimize-isl.cc (optimize_isl): Adjust
+       param_max_isl_operations value for OpenACC functions and add
+       special warnings if value gets exceeded.
+
+       * graphite-scop-detection.cc (build_scops): Likewise for
+       param_graphite_max_arrays_per_scop.
+
 2021-11-16  Frederik Harwath  <frederik@codesourcery.com>
 
        * tree-ssa-pre.cc (insert): Skip any insertions in OpenACC

diff --git a/gcc/graphite-optimize-isl.cc b/gcc/graphite-optimize-isl.cc
index e786bd48c79a4493b26acf29797075dfe2fed689..9eff4e1e22a6c20cdfc3a6fe2e910ef66e29a623 100644 (file)
--- a/gcc/graphite-optimize-isl.cc
+++ b/gcc/graphite-optimize-isl.cc
@@ -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");
        }

diff --git a/gcc/graphite-scop-detection.cc b/gcc/graphite-scop-detection.cc
index 4b1d330bdb4da7e0d95f222f12f4f8cc7458688d..8656e4b1e99f1a2dd4fd40e1858a0ee19244602c 100644 (file)
--- a/gcc/graphite-scop-detection.cc
+++ b/gcc/graphite-scop-detection.cc
@@ -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;
        }

diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index c78e6cc2a1de852472d97e85cfa89ce84df4ec6c..a3553fc550462bf45dd981392f117a3c6d739be8 100644 (file)
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,8 @@
+2021-11-16  Frederik Harwath  <frederik@codesourcery.com>
+
+       * gcc.dg/goacc/graphite-parameter-1.c: New test.
+       * gcc.dg/goacc/graphite-parameter-2.c: New test.
+
 2021-11-16  Andrew Stubbs  <ams@codesourcery.com>
            Thomas Schwinge <thomas@codesourcery.com>
 

diff --git a/gcc/testsuite/gcc.dg/goacc/graphite-parameter-1.c b/gcc/testsuite/gcc.dg/goacc/graphite-parameter-1.c
new file mode 100644 (file)
index 0000000..45adbb3
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/graphite-parameter-1.c
@@ -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.*"} */

diff --git a/gcc/testsuite/gcc.dg/goacc/graphite-parameter-2.c b/gcc/testsuite/gcc.dg/goacc/graphite-parameter-2.c
new file mode 100644 (file)
index 0000000..f2830cd
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/graphite-parameter-2.c
@@ -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;
+    }
+}