From patchwork Wed Nov 17 16:03:27 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 47832 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 930FA385C41C for ; Wed, 17 Nov 2021 16:20:51 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 8EBFE3858410 for ; Wed, 17 Nov 2021 16:04:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 8EBFE3858410 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 6covh08JQfELSHc1NuMPaGAYppv8apj8SCQU6XQynHkwcZr1Wouqo5fgIYbKP4hIhKfArzfR4u F51LHruNB5igNOiFOUhkrrsKCYvudfjaOJAMG/CEoTokriS5aVbdH8Lh6t/xgtDaRLQkquklsL 9ds1DyUtwFogyHowx0h1535F3+KJL4Rp9dcR6fnzhQ4ZEDdbrvXxB2ZG82XHyROlU3cUuYGUIq YkIAMsqAmD4paA//xWm+JKT3faEEewxeKlkRZORPMvPO4fXrxNGzgp4EIE5axRUjrpn4iemTz0 dTXYbu7No76gjYtVvv9XOALg X-IronPort-AV: E=Sophos;i="5.87,241,1631606400"; d="scan'208";a="68445393" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 17 Nov 2021 08:04:55 -0800 IronPort-SDR: QbhpEhKgCUfSsG4ofjBJ1JrTClrXaZMFFNhQc4XRouY1BNAWSlIniUOPL7+3nqrOVysbqxJMDj ZlPcjqTALlrlJpE5Xgvf6EWXaNQPnXuPQW664sO1ZWb8sI7WCYU+bUqyxbr4EJGfZqwukBoeul cBPXPCu9jSfYUKuPv5FIRQ00r26Hjhd5ihjGKW+qYqAzMQlLVrlKWcqIZNA8w22Q+vY9GF4jqW DpR63HkHbEKWc4O2h4KZlN0yDznaeel79WHG7B3y73rk9C49CXsgcM6hB7pDA3Ux3mCGYIkxjO p5s= From: Frederik Harwath To: Subject: [OG11][committed][PATCH 19/22] graphite: Tune parameters for OpenACC use Date: Wed, 17 Nov 2021 17:03:27 +0100 Message-ID: <20211117160330.20029-19-frederik@codesourcery.com> X-Mailer: git-send-email 2.33.0 In-Reply-To: <20211117160330.20029-1-frederik@codesourcery.com> References: <20211117160330.20029-1-frederik@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.6 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" 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.c (optimize_isl): Adjust param_max_isl_operations value for OpenACC functions and add special warnings if value gets exceeded. * graphite-scop-detection.c (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. --- gcc/graphite-optimize-isl.c | 35 ++++++++++++++++--- gcc/graphite-scop-detection.c | 28 ++++++++++++++- .../gcc.dg/goacc/graphite-parameter-1.c | 21 +++++++++++ .../gcc.dg/goacc/graphite-parameter-2.c | 23 ++++++++++++ 4 files changed, 101 insertions(+), 6 deletions(-) create mode 100644 gcc/testsuite/gcc.dg/goacc/graphite-parameter-1.c create mode 100644 gcc/testsuite/gcc.dg/goacc/graphite-parameter-2.c -- 2.33.0 ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 diff --git a/gcc/graphite-optimize-isl.c b/gcc/graphite-optimize-isl.c index 019452700a49..4eecbd20b740 100644 --- a/gcc/graphite-optimize-isl.c +++ b/gcc/graphite-optimize-isl.c @@ -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.c b/gcc/graphite-scop-detection.c index 8b41044bce5e..afc955cc97eb 100644 --- a/gcc/graphite-scop-detection.c +++ b/gcc/graphite-scop-detection.c @@ -2056,6 +2056,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. */ @@ -2109,6 +2112,11 @@ build_scops (vec *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) { @@ -2116,7 +2124,16 @@ build_scops (vec *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; } @@ -2129,6 +2146,15 @@ build_scops (vec *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/gcc.dg/goacc/graphite-parameter-1.c b/gcc/testsuite/gcc.dg/goacc/graphite-parameter-1.c new file mode 100644 index 000000000000..45adbb3f0e85 --- /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 index 000000000000..f2830cd62db0 --- /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; + } +}