1 | /* A scheduling optimizer for Graphite |
2 | Copyright (C) 2012-2023 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 | 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 * |
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); |
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 |
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 | |