1 /* A scheduling optimizer for Graphite 2 Copyright (C) 2012-2020 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 USES_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 */ 225