1 /* A scheduling optimizer for Graphite
2 Copyright (C) 2012-2022 Free Software Foundation, Inc.
3 Contributed by Tobias Grosser <tobias@grosser.es>.
5 This file is part of GCC.
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)
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.
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/>. */
28 #include "coretypes.h"
33 #include "fold-const.h"
34 #include "gimple-iterator.h"
35 #include "tree-ssa-loop.h"
37 #include "tree-data-ref.h"
39 #include "tree-vectorizer.h"
43 /* get_schedule_for_node_st - Improve schedule for the schedule node.
44 Only Simple loop tiling is considered. */
46 static __isl_give isl_schedule_node *
47 get_schedule_for_node_st (__isl_take isl_schedule_node *node, void *user)
52 if (isl_schedule_node_get_type (node) != isl_schedule_node_band
53 || isl_schedule_node_n_children (node) != 1)
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);
63 if (type != isl_schedule_node_leaf)
66 long tile_size = param_loop_block_tile_size;
69 || !isl_schedule_node_band_get_permutable (node))
71 if (dump_file && dump_flags)
72 fprintf (dump_file, "not tiled\n");
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);
81 for (unsigned i = 0; i < dims; i++)
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);
89 node = isl_schedule_node_band_tile (node, sizes);
90 node = isl_schedule_node_child (node, 0);
95 static isl_union_set *
96 scop_get_domains (scop_p scop)
100 isl_space *space = isl_set_get_space (scop->param_context);
101 isl_union_set *res = isl_union_set_empty (space);
103 FOR_EACH_VEC_ELT (scop->pbbs, i, pbb)
104 res = isl_union_set_add_set (res, isl_set_copy (pbb->domain));
109 /* Compute the schedule for SCOP based on its parameters, domain and set of
110 constraints. Then apply the schedule to SCOP. */
113 optimize_isl (scop_p scop)
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;
119 isl_ctx_set_max_operations (scop->isl_context, max_operations);
120 isl_options_set_on_error (scop->isl_context, ISL_ON_ERROR_CONTINUE);
122 isl_union_set *domain = scop_get_domains (scop);
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));
132 /* FIXME: proximity should not be validity. */
133 isl_union_map *proximity = isl_union_map_copy (validity);
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);
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);
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);
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)
162 if (dump_enabled_p ())
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",
172 dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
173 "loop nest not optimized, ISL signalled an error\n");
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);
187 if (dump_enabled_p ())
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");
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;
205 /* Apply graphite transformations to all the basic blocks of SCOP. */
208 apply_poly_transforms (scop_p scop)
210 if (flag_loop_nest_optimize)
211 return optimize_isl (scop);
213 if (!flag_graphite_identity && !flag_loop_parallelize_all)
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);
224 #endif /* HAVE_isl */