gcc/graphite-optimize-isl.cc

   1 /* A scheduling optimizer for Graphite
   2    Copyright (C) 2012-2023 Free Software Foundation, Inc.
   3    Contributed by Tobias Grosser <tobias@grosser.es>.
   4
   5 This file is part of GCC.
   6
   7 GCC is free software; you can redistribute it and/or modify
   8 it under the terms of the GNU General Public License as published by
   9 the Free Software Foundation; either version 3, or (at your option)
  10 any later version.
  11
  12 GCC is distributed in the hope that it will be useful,
  13 but WITHOUT ANY WARRANTY; without even the implied warranty of
  14 MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
  15 GNU General Public License for more details.
  16
  17 You should have received a copy of the GNU General Public License
  18 along with GCC; see the file COPYING3.  If not see
  19 <http://www.gnu.org/licenses/>.  */
  20
  21 #define INCLUDE_ISL
  22
  23 #include "config.h"
  24
  25 #ifdef HAVE_isl
  26
  27 #include "system.h"
  28 #include "coretypes.h"
  29 #include "backend.h"
  30 #include "cfghooks.h"
  31 #include "tree.h"
  32 #include "gimple.h"
  33 #include "fold-const.h"
  34 #include "gimple-iterator.h"
  35 #include "tree-ssa-loop.h"
  36 #include "cfgloop.h"
  37 #include "tree-data-ref.h"
  38 #include "dumpfile.h"
  39 #include "tree-vectorizer.h"
  40 #include "graphite.h"
  41
  42
  43 /* get_schedule_for_node_st - Improve schedule for the schedule node.
  44    Only Simple loop tiling is considered.  */
  45
  46 static __isl_give isl_schedule_node *
  47 get_schedule_for_node_st (__isl_take isl_schedule_node *node, void *user)
  48 {
  49   if (user)
  50     return node;
  51
  52   if (isl_schedule_node_get_type (node) != isl_schedule_node_band
  53       || isl_schedule_node_n_children (node) != 1)
  54     return node;
  55
  56   isl_space *space = isl_schedule_node_band_get_space (node);
  57   unsigned dims = isl_space_dim (space, isl_dim_set);
  58   isl_schedule_node *child = isl_schedule_node_get_child (node, 0);
  59   isl_schedule_node_type type = isl_schedule_node_get_type (child);
  60   isl_space_free (space);
  61   isl_schedule_node_free (child);
  62
  63   if (type != isl_schedule_node_leaf)
  64     return node;
  65
  66   long tile_size = param_loop_block_tile_size;
  67   if (dims <= 1
  68       || tile_size == 0
  69       || !isl_schedule_node_band_get_permutable (node))
  70     {
  71       if (dump_file && dump_flags)
  72         fprintf (dump_file, "not tiled\n");
  73       return node;
  74     }
  75
  76   /* Tile loops.  */
  77   space = isl_schedule_node_band_get_space (node);
  78   isl_multi_val *sizes = isl_multi_val_zero (space);
  79   isl_ctx *ctx = isl_schedule_node_get_ctx (node);
  80
  81   for (unsigned i = 0; i < dims; i++)
  82     {
  83       sizes = isl_multi_val_set_val (sizes, i,
  84                                      isl_val_int_from_si (ctx, tile_size));
  85       if (dump_file && dump_flags)
  86         fprintf (dump_file, "tiled by %ld\n", tile_size);
  87     }
  88
  89   node = isl_schedule_node_band_tile (node, sizes);
  90   node = isl_schedule_node_child (node, 0);
  91
  92   return node;
  93 }
  94
  95 static isl_union_set *
  96 scop_get_domains (scop_p scop)
  97 {
  98   int i;
  99   poly_bb_p pbb;
 100   isl_space *space = isl_set_get_space (scop->param_context);
 101   isl_union_set *res = isl_union_set_empty (space);
 102
 103   FOR_EACH_VEC_ELT (scop->pbbs, i, pbb)
 104     res = isl_union_set_add_set (res, isl_set_copy (pbb->domain));
 105
 106   return res;
 107 }
 108
 109 /* Compute the schedule for SCOP based on its parameters, domain and set of
 110    constraints.  Then apply the schedule to SCOP.  */
 111
 112 static bool
 113 optimize_isl (scop_p scop)
 114 {
 115   int old_err = isl_options_get_on_error (scop->isl_context);
 116   int old_max_operations = isl_ctx_get_max_operations (scop->isl_context);
 117   int max_operations = param_max_isl_operations;
 118   if (max_operations)
 119     isl_ctx_set_max_operations (scop->isl_context, max_operations);
 120   isl_options_set_on_error (scop->isl_context, ISL_ON_ERROR_CONTINUE);
 121
 122   isl_union_set *domain = scop_get_domains (scop);
 123
 124   /* Simplify the dependences on the domain.  */
 125   scop_get_dependences (scop);
 126   isl_union_map *dependences
 127     = isl_union_map_gist_domain (isl_union_map_copy (scop->dependence),
 128                                  isl_union_set_copy (domain));
 129   isl_union_map *validity
 130     = isl_union_map_gist_range (dependences, isl_union_set_copy (domain));
 131
 132   /* FIXME: proximity should not be validity.  */
 133   isl_union_map *proximity = isl_union_map_copy (validity);
 134
 135   isl_schedule_constraints *sc = isl_schedule_constraints_on_domain (domain);
 136   sc = isl_schedule_constraints_set_proximity (sc, proximity);
 137   sc = isl_schedule_constraints_set_validity (sc, isl_union_map_copy (validity));
 138   sc = isl_schedule_constraints_set_coincidence (sc, validity);
 139
 140   isl_options_set_schedule_serialize_sccs (scop->isl_context, 0);
 141   isl_options_set_schedule_maximize_band_depth (scop->isl_context, 1);
 142   isl_options_set_schedule_max_constant_term (scop->isl_context, 20);
 143   isl_options_set_schedule_max_coefficient (scop->isl_context, 20);
 144   isl_options_set_tile_scale_tile_loops (scop->isl_context, 0);
 145   /* Generate loop upper bounds that consist of the current loop iterator, an
 146      operator (< or <=) and an expression not involving the iterator.  If this
 147      option is not set, then the current loop iterator may appear several times
 148      in the upper bound.  See the isl manual for more details.  */
 149   isl_options_set_ast_build_atomic_upper_bound (scop->isl_context, 1);
 150
 151   scop->transformed_schedule = isl_schedule_constraints_compute_schedule (sc);
 152   scop->transformed_schedule =
 153     isl_schedule_map_schedule_node_bottom_up (scop->transformed_schedule,
 154                                               get_schedule_for_node_st, NULL);
 155
 156   isl_options_set_on_error (scop->isl_context, old_err);
 157   isl_ctx_reset_operations (scop->isl_context);
 158   isl_ctx_set_max_operations (scop->isl_context, old_max_operations);
 159   if (!scop->transformed_schedule
 160       || isl_ctx_last_error (scop->isl_context) != isl_error_none)
 161     {
 162       if (dump_enabled_p ())
 163         {
 164           dump_user_location_t loc = find_loop_location
 165             (scop->scop_info->region.entry->dest->loop_father);
 166           if (isl_ctx_last_error (scop->isl_context) == isl_error_quota)
 167             dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
 168                              "loop nest not optimized, optimization timed out "
 169                              "after %d operations [--param max-isl-operations]\n",
 170                              max_operations);
 171           else
 172             dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
 173                              "loop nest not optimized, ISL signalled an error\n");
 174         }
 175       return false;
 176     }
 177
 178   gcc_assert (scop->original_schedule);
 179   isl_union_map *original = isl_schedule_get_map (scop->original_schedule);
 180   isl_union_map *transformed = isl_schedule_get_map (scop->transformed_schedule);
 181   bool same_schedule = isl_union_map_is_equal (original, transformed);
 182   isl_union_map_free (original);
 183   isl_union_map_free (transformed);
 184
 185   if (same_schedule)
 186     {
 187       if (dump_enabled_p ())
 188         {
 189           dump_user_location_t loc = find_loop_location
 190             (scop->scop_info->region.entry->dest->loop_father);
 191           dump_printf_loc (MSG_NOTE, loc,
 192                            "loop nest not optimized, optimized schedule is "
 193                            "identical to original schedule\n");
 194         }
 195       if (dump_file)
 196         print_schedule_ast (dump_file, scop->original_schedule, scop);
 197       isl_schedule_free (scop->transformed_schedule);
 198       scop->transformed_schedule = isl_schedule_copy (scop->original_schedule);
 199       return flag_graphite_identity || flag_loop_parallelize_all;
 200     }
 201
 202   return true;
 203 }
 204
 205 /* Apply graphite transformations to all the basic blocks of SCOP.  */
 206
 207 bool
 208 apply_poly_transforms (scop_p scop)
 209 {
 210   if (flag_loop_nest_optimize)
 211     return optimize_isl (scop);
 212
 213   if (!flag_graphite_identity && !flag_loop_parallelize_all)
 214     return false;
 215
 216   /* Generate code even if we did not apply any real transformation.
 217      This also allows to check the performance for the identity
 218      transformation: GIMPLE -> GRAPHITE -> GIMPLE.  */
 219   gcc_assert (scop->original_schedule);
 220   scop->transformed_schedule = isl_schedule_copy (scop->original_schedule);
 221   return true;
 222 }
 223
 224 #endif /* HAVE_isl */