1 /* A scheduling optimizer for Graphite
2 Copyright (C) 2012-2021 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 *
get_schedule_for_node_st(__isl_take isl_schedule_node * node,void * user)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 *
scop_get_domains(scop_p scop)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
optimize_isl(scop_p scop)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
apply_poly_transforms(scop_p scop)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