|
111
|
1 /* A scheduling optimizer for Graphite
|
|
145
|
2 Copyright (C) 2012-2020 Free Software Foundation, Inc.
|
|
111
|
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
|
|
145
|
66 long tile_size = param_loop_block_tile_size;
|
|
111
|
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);
|
|
145
|
117 int max_operations = param_max_isl_operations;
|
|
111
|
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 {
|
|
145
|
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 }
|
|
111
|
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 {
|
|
145
|
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 }
|
|
111
|
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 */
|
