Branch data Line data Source code
1 : : /* A scheduling optimizer for Graphite
2 : : Copyright (C) 2012-2024 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 : 1271 : get_schedule_for_node_st (__isl_take isl_schedule_node *node, void *user)
48 : : {
49 : 1271 : if (user)
50 : : return node;
51 : :
52 : 1271 : if (isl_schedule_node_get_type (node) != isl_schedule_node_band
53 : 1271 : || isl_schedule_node_n_children (node) != 1)
54 : 1020 : return node;
55 : :
56 : 251 : isl_space *space = isl_schedule_node_band_get_space (node);
57 : 251 : unsigned dims = isl_space_dim (space, isl_dim_set);
58 : 251 : isl_schedule_node *child = isl_schedule_node_get_child (node, 0);
59 : 251 : isl_schedule_node_type type = isl_schedule_node_get_type (child);
60 : 251 : isl_space_free (space);
61 : 251 : isl_schedule_node_free (child);
62 : :
63 : 251 : if (type != isl_schedule_node_leaf)
64 : : return node;
65 : :
66 : 83 : long tile_size = param_loop_block_tile_size;
67 : 83 : if (dims <= 1
68 : 83 : || tile_size == 0
69 : 83 : || !isl_schedule_node_band_get_permutable (node))
70 : : {
71 : 32 : if (dump_file && dump_flags)
72 : 5 : fprintf (dump_file, "not tiled\n");
73 : 32 : return node;
74 : : }
75 : :
76 : : /* Tile loops. */
77 : 51 : space = isl_schedule_node_band_get_space (node);
78 : 51 : isl_multi_val *sizes = isl_multi_val_zero (space);
79 : 51 : isl_ctx *ctx = isl_schedule_node_get_ctx (node);
80 : :
81 : 159 : for (unsigned i = 0; i < dims; i++)
82 : : {
83 : 108 : sizes = isl_multi_val_set_val (sizes, i,
84 : : isl_val_int_from_si (ctx, tile_size));
85 : 108 : if (dump_file && dump_flags)
86 : 53 : fprintf (dump_file, "tiled by %ld\n", tile_size);
87 : : }
88 : :
89 : 51 : node = isl_schedule_node_band_tile (node, sizes);
90 : 51 : node = isl_schedule_node_child (node, 0);
91 : :
92 : 51 : return node;
93 : : }
94 : :
95 : : static isl_union_set *
96 : 109 : scop_get_domains (scop_p scop)
97 : : {
98 : 109 : int i;
99 : 109 : poly_bb_p pbb;
100 : 109 : isl_space *space = isl_set_get_space (scop->param_context);
101 : 109 : isl_union_set *res = isl_union_set_empty (space);
102 : :
103 : 578 : FOR_EACH_VEC_ELT (scop->pbbs, i, pbb)
104 : 360 : res = isl_union_set_add_set (res, isl_set_copy (pbb->domain));
105 : :
106 : 109 : 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 : 109 : optimize_isl (scop_p scop)
114 : : {
115 : 109 : int old_err = isl_options_get_on_error (scop->isl_context);
116 : 109 : int old_max_operations = isl_ctx_get_max_operations (scop->isl_context);
117 : 109 : int max_operations = param_max_isl_operations;
118 : 109 : if (max_operations)
119 : 109 : isl_ctx_set_max_operations (scop->isl_context, max_operations);
120 : 109 : isl_options_set_on_error (scop->isl_context, ISL_ON_ERROR_CONTINUE);
121 : :
122 : 109 : isl_union_set *domain = scop_get_domains (scop);
123 : :
124 : : /* Simplify the dependences on the domain. */
125 : 109 : scop_get_dependences (scop);
126 : 109 : isl_union_map *dependences
127 : 109 : = isl_union_map_gist_domain (isl_union_map_copy (scop->dependence),
128 : : isl_union_set_copy (domain));
129 : 109 : isl_union_map *validity
130 : 109 : = isl_union_map_gist_range (dependences, isl_union_set_copy (domain));
131 : :
132 : : /* FIXME: proximity should not be validity. */
133 : 109 : isl_union_map *proximity = isl_union_map_copy (validity);
134 : :
135 : 109 : isl_schedule_constraints *sc = isl_schedule_constraints_on_domain (domain);
136 : 109 : sc = isl_schedule_constraints_set_proximity (sc, proximity);
137 : 109 : sc = isl_schedule_constraints_set_validity (sc, isl_union_map_copy (validity));
138 : 109 : sc = isl_schedule_constraints_set_coincidence (sc, validity);
139 : :
140 : 109 : isl_options_set_schedule_serialize_sccs (scop->isl_context, 0);
141 : 109 : isl_options_set_schedule_maximize_band_depth (scop->isl_context, 1);
142 : 109 : isl_options_set_schedule_max_constant_term (scop->isl_context, 20);
143 : 109 : isl_options_set_schedule_max_coefficient (scop->isl_context, 20);
144 : 109 : 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 : 109 : isl_options_set_ast_build_atomic_upper_bound (scop->isl_context, 1);
150 : :
151 : 109 : scop->transformed_schedule = isl_schedule_constraints_compute_schedule (sc);
152 : 218 : scop->transformed_schedule =
153 : 109 : isl_schedule_map_schedule_node_bottom_up (scop->transformed_schedule,
154 : : get_schedule_for_node_st, NULL);
155 : :
156 : 109 : isl_options_set_on_error (scop->isl_context, old_err);
157 : 109 : isl_ctx_reset_operations (scop->isl_context);
158 : 109 : isl_ctx_set_max_operations (scop->isl_context, old_max_operations);
159 : 109 : if (!scop->transformed_schedule
160 : 109 : || isl_ctx_last_error (scop->isl_context) != isl_error_none)
161 : : {
162 : 2 : if (dump_enabled_p ())
163 : : {
164 : 1 : dump_user_location_t loc = find_loop_location
165 : 1 : (scop->scop_info->region.entry->dest->loop_father);
166 : 1 : if (isl_ctx_last_error (scop->isl_context) == isl_error_quota)
167 : 1 : 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 : 0 : dump_printf_loc (MSG_MISSED_OPTIMIZATION, loc,
173 : : "loop nest not optimized, ISL signalled an error\n");
174 : : }
175 : 2 : return false;
176 : : }
177 : :
178 : 107 : gcc_assert (scop->original_schedule);
179 : 107 : isl_union_map *original = isl_schedule_get_map (scop->original_schedule);
180 : 107 : isl_union_map *transformed = isl_schedule_get_map (scop->transformed_schedule);
181 : 107 : bool same_schedule = isl_union_map_is_equal (original, transformed);
182 : 107 : isl_union_map_free (original);
183 : 107 : isl_union_map_free (transformed);
184 : :
185 : 107 : if (same_schedule)
186 : : {
187 : 6 : if (dump_enabled_p ())
188 : : {
189 : 1 : dump_user_location_t loc = find_loop_location
190 : 1 : (scop->scop_info->region.entry->dest->loop_father);
191 : 1 : dump_printf_loc (MSG_NOTE, loc,
192 : : "loop nest not optimized, optimized schedule is "
193 : : "identical to original schedule\n");
194 : : }
195 : 6 : if (dump_file)
196 : 1 : print_schedule_ast (dump_file, scop->original_schedule, scop);
197 : 6 : isl_schedule_free (scop->transformed_schedule);
198 : 6 : scop->transformed_schedule = isl_schedule_copy (scop->original_schedule);
199 : 6 : 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 : 193 : apply_poly_transforms (scop_p scop)
209 : : {
210 : 193 : if (flag_loop_nest_optimize)
211 : 109 : return optimize_isl (scop);
212 : :
213 : 84 : 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 : 57 : gcc_assert (scop->original_schedule);
220 : 57 : scop->transformed_schedule = isl_schedule_copy (scop->original_schedule);
221 : 57 : return true;
222 : : }
223 : :
224 : : #endif /* HAVE_isl */
|
