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