]>
Commit | Line | Data |
---|---|---|
b60cc080 | 1 | /* A scheduling optimizer for Graphite |
7adcbafe | 2 | Copyright (C) 2012-2022 Free Software Foundation, Inc. |
b60cc080 TG |
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 | ||
deda4625 | 21 | #define INCLUDE_ISL |
4d776011 | 22 | |
b60cc080 TG |
23 | #include "config.h" |
24 | ||
eae1a5d4 | 25 | #ifdef HAVE_isl |
4d776011 DE |
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" | |
4d776011 | 38 | #include "dumpfile.h" |
4d6e2f33 | 39 | #include "tree-vectorizer.h" |
cf98f0f4 | 40 | #include "graphite.h" |
b60cc080 | 41 | |
5affe17f AZ |
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) | |
b60cc080 | 48 | { |
5affe17f AZ |
49 | if (user) |
50 | return node; | |
b60cc080 | 51 | |
5affe17f AZ |
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 | ||
028d4092 | 66 | long tile_size = param_loop_block_tile_size; |
c1f80eef RB |
67 | if (dims <= 1 |
68 | || tile_size == 0 | |
69 | || !isl_schedule_node_band_get_permutable (node)) | |
5affe17f AZ |
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); | |
5affe17f AZ |
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); | |
b60cc080 | 91 | |
5affe17f | 92 | return node; |
b60cc080 TG |
93 | } |
94 | ||
adba512d AK |
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); | |
5affe17f | 102 | |
adba512d AK |
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. */ | |
5affe17f | 111 | |
adba512d AK |
112 | static bool |
113 | optimize_isl (scop_p scop) | |
5affe17f | 114 | { |
871a0725 | 115 | int old_err = isl_options_get_on_error (scop->isl_context); |
adba512d | 116 | int old_max_operations = isl_ctx_get_max_operations (scop->isl_context); |
028d4092 | 117 | int max_operations = param_max_isl_operations; |
adba512d AK |
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); | |
5affe17f | 121 | |
adba512d AK |
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); | |
adba512d | 155 | |
871a0725 | 156 | isl_options_set_on_error (scop->isl_context, old_err); |
adba512d AK |
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 | |
871a0725 | 160 | || isl_ctx_last_error (scop->isl_context) != isl_error_none) |
adba512d | 161 | { |
bbeeac91 DM |
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 | } | |
adba512d AK |
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 | { | |
bbeeac91 DM |
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 | } | |
adba512d | 195 | if (dump_file) |
871a0725 | 196 | print_schedule_ast (dump_file, scop->original_schedule, scop); |
adba512d AK |
197 | isl_schedule_free (scop->transformed_schedule); |
198 | scop->transformed_schedule = isl_schedule_copy (scop->original_schedule); | |
d2552094 | 199 | return flag_graphite_identity || flag_loop_parallelize_all; |
adba512d AK |
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; | |
5affe17f | 222 | } |
adba512d | 223 | |
ec62c373 | 224 | #endif /* HAVE_isl */ |