1/* A scheduling optimizer for Graphite
2 Copyright (C) 2012-2023 Free Software Foundation, Inc.
3 Contributed by Tobias Grosser <tobias@grosser.es>.
4
5This file is part of GCC.
6
7GCC is free software; you can redistribute it and/or modify
8it under the terms of the GNU General Public License as published by
9the Free Software Foundation; either version 3, or (at your option)
10any later version.
11
12GCC is distributed in the hope that it will be useful,
13but WITHOUT ANY WARRANTY; without even the implied warranty of
14MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
15GNU General Public License for more details.
16
17You should have received a copy of the GNU General Public License
18along 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
46static __isl_give isl_schedule_node *
47get_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
95static isl_union_set *
96scop_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
112static bool
113optimize_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
207bool
208apply_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

source code of gcc/graphite-optimize-isl.cc