1 /*
2 * Copyright © 2021 Intel Corporation
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice (including the next
12 * paragraph) shall be included in all copies or substantial portions of the
13 * Software.
14 *
15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
21 * IN THE SOFTWARE.
22 */
23
24
25 #include "brw_private.h"
26 #include "compiler/shader_info.h"
27 #include "intel/dev/intel_debug.h"
28 #include "intel/dev/intel_device_info.h"
29 #include "util/ralloc.h"
30
31 #include <gtest/gtest.h>
32
33 enum {
34 SIMD8 = 0,
35 SIMD16 = 1,
36 SIMD32 = 2,
37 };
38
39 const bool spilled = true;
40 const bool not_spilled = false;
41
42 class SIMDSelectionTest : public ::testing::Test {
43 protected:
SIMDSelectionTest()44 SIMDSelectionTest()
45 : mem_ctx(ralloc_context(NULL))
46 , devinfo(rzalloc(mem_ctx, intel_device_info))
47 , prog_data(rzalloc(mem_ctx, struct brw_cs_prog_data))
48 , simd_state{
49 .devinfo = devinfo,
50 .prog_data = prog_data,
51 }
52 {
53 process_intel_debug_variable();
54 }
55
~SIMDSelectionTest()56 ~SIMDSelectionTest() {
57 ralloc_free(mem_ctx);
58 };
59
60 void *mem_ctx;
61 intel_device_info *devinfo;
62 struct brw_cs_prog_data *prog_data;
63 brw_simd_selection_state simd_state;
64 };
65
66 class SIMDSelectionCS : public SIMDSelectionTest {
67 protected:
SIMDSelectionCS()68 SIMDSelectionCS() {
69 prog_data->base.stage = MESA_SHADER_COMPUTE;
70 prog_data->local_size[0] = 32;
71 prog_data->local_size[1] = 1;
72 prog_data->local_size[2] = 1;
73
74 devinfo->max_cs_workgroup_threads = 64;
75 }
76 };
77
TEST_F(SIMDSelectionCS,DefaultsToSIMD16)78 TEST_F(SIMDSelectionCS, DefaultsToSIMD16)
79 {
80 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
81 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
82 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
83 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
84 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
85
86 ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
87 }
88
TEST_F(SIMDSelectionCS,TooBigFor16)89 TEST_F(SIMDSelectionCS, TooBigFor16)
90 {
91 prog_data->local_size[0] = devinfo->max_cs_workgroup_threads;
92 prog_data->local_size[1] = 32;
93 prog_data->local_size[2] = 1;
94
95 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
96 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
97 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
98 brw_simd_mark_compiled(simd_state, SIMD32, spilled);
99
100 ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
101 }
102
TEST_F(SIMDSelectionCS,WorkgroupSize1)103 TEST_F(SIMDSelectionCS, WorkgroupSize1)
104 {
105 prog_data->local_size[0] = 1;
106 prog_data->local_size[1] = 1;
107 prog_data->local_size[2] = 1;
108
109 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
110 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
111 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
112 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
113
114 ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
115 }
116
TEST_F(SIMDSelectionCS,WorkgroupSize8)117 TEST_F(SIMDSelectionCS, WorkgroupSize8)
118 {
119 prog_data->local_size[0] = 8;
120 prog_data->local_size[1] = 1;
121 prog_data->local_size[2] = 1;
122
123 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
124 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
125 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
126 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
127
128 ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
129 }
130
TEST_F(SIMDSelectionCS,WorkgroupSizeVariable)131 TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
132 {
133 prog_data->local_size[0] = 0;
134 prog_data->local_size[1] = 0;
135 prog_data->local_size[2] = 0;
136
137 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
138 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
139 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
140 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
141 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
142 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
143
144 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
145
146 const unsigned wg_8_1_1[] = { 8, 1, 1 };
147 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
148
149 const unsigned wg_16_1_1[] = { 16, 1, 1 };
150 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16);
151
152 const unsigned wg_32_1_1[] = { 32, 1, 1 };
153 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16);
154 }
155
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableSpilled)156 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled)
157 {
158 prog_data->local_size[0] = 0;
159 prog_data->local_size[1] = 0;
160 prog_data->local_size[2] = 0;
161
162 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
163 brw_simd_mark_compiled(simd_state, SIMD8, spilled);
164 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
165 brw_simd_mark_compiled(simd_state, SIMD16, spilled);
166 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
167 brw_simd_mark_compiled(simd_state, SIMD32, spilled);
168
169 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
170
171 const unsigned wg_8_1_1[] = { 8, 1, 1 };
172 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
173
174 const unsigned wg_16_1_1[] = { 16, 1, 1 };
175 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8);
176
177 const unsigned wg_32_1_1[] = { 32, 1, 1 };
178 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8);
179 }
180
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD8)181 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8)
182 {
183 prog_data->local_size[0] = 0;
184 prog_data->local_size[1] = 0;
185 prog_data->local_size[2] = 0;
186
187 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
188 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
189 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
190 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
191 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
192
193 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32);
194
195 const unsigned wg_8_1_1[] = { 8, 1, 1 };
196 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD16);
197
198 const unsigned wg_16_1_1[] = { 16, 1, 1 };
199 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16);
200
201 const unsigned wg_32_1_1[] = { 32, 1, 1 };
202 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16);
203 }
204
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD16)205 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD16)
206 {
207 prog_data->local_size[0] = 0;
208 prog_data->local_size[1] = 0;
209 prog_data->local_size[2] = 0;
210
211 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
212 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
213 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
214 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
215 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
216
217 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD32);
218
219 const unsigned wg_8_1_1[] = { 8, 1, 1 };
220 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
221
222 const unsigned wg_16_1_1[] = { 16, 1, 1 };
223 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8);
224
225 const unsigned wg_32_1_1[] = { 32, 1, 1 };
226 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8);
227 }
228
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD8NoSIMD16)229 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16)
230 {
231 prog_data->local_size[0] = 0;
232 prog_data->local_size[1] = 0;
233 prog_data->local_size[2] = 0;
234
235 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
236 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
237 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
238 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
239
240 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32);
241
242 const unsigned wg_8_1_1[] = { 8, 1, 1 };
243 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD32);
244
245 const unsigned wg_16_1_1[] = { 16, 1, 1 };
246 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD32);
247
248 const unsigned wg_32_1_1[] = { 32, 1, 1 };
249 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD32);
250 }
251
TEST_F(SIMDSelectionCS,SpillAtSIMD8)252 TEST_F(SIMDSelectionCS, SpillAtSIMD8)
253 {
254 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
255 brw_simd_mark_compiled(simd_state, SIMD8, spilled);
256 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
257 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
258
259 ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
260 }
261
TEST_F(SIMDSelectionCS,SpillAtSIMD16)262 TEST_F(SIMDSelectionCS, SpillAtSIMD16)
263 {
264 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
265 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
266 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
267 brw_simd_mark_compiled(simd_state, SIMD16, spilled);
268 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
269
270 ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
271 }
272
TEST_F(SIMDSelectionCS,EnvironmentVariable32)273 TEST_F(SIMDSelectionCS, EnvironmentVariable32)
274 {
275 intel_debug |= DEBUG_DO32;
276
277 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
278 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
279 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
280 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
281 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
282 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
283
284 ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
285 }
286
TEST_F(SIMDSelectionCS,EnvironmentVariable32ButSpills)287 TEST_F(SIMDSelectionCS, EnvironmentVariable32ButSpills)
288 {
289 intel_debug |= DEBUG_DO32;
290
291 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
292 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
293 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
294 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
295 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
296 brw_simd_mark_compiled(simd_state, SIMD32, spilled);
297
298 ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
299 }
300
TEST_F(SIMDSelectionCS,Require8)301 TEST_F(SIMDSelectionCS, Require8)
302 {
303 simd_state.required_width = 8;
304
305 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
306 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
307 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
308 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
309
310 ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
311 }
312
TEST_F(SIMDSelectionCS,Require8ErrorWhenNotCompile)313 TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile)
314 {
315 simd_state.required_width = 8;
316
317 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
318 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
319 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
320
321 ASSERT_EQ(brw_simd_select(simd_state), -1);
322 }
323
TEST_F(SIMDSelectionCS,Require16)324 TEST_F(SIMDSelectionCS, Require16)
325 {
326 simd_state.required_width = 16;
327
328 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
329 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
330 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
331 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
332
333 ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
334 }
335
TEST_F(SIMDSelectionCS,Require16ErrorWhenNotCompile)336 TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile)
337 {
338 simd_state.required_width = 16;
339
340 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
341 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
342 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32));
343
344 ASSERT_EQ(brw_simd_select(simd_state), -1);
345 }
346
TEST_F(SIMDSelectionCS,Require32)347 TEST_F(SIMDSelectionCS, Require32)
348 {
349 simd_state.required_width = 32;
350
351 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
352 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
353 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
354 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
355
356 ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
357 }
358
TEST_F(SIMDSelectionCS,Require32ErrorWhenNotCompile)359 TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile)
360 {
361 simd_state.required_width = 32;
362
363 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8));
364 ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16));
365 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
366
367 ASSERT_EQ(brw_simd_select(simd_state), -1);
368 }
369
TEST_F(SIMDSelectionCS,FirstCompiledIsSIMD8)370 TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD8)
371 {
372 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
373 brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
374
375 ASSERT_TRUE(brw_simd_any_compiled(simd_state));
376 ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD8);
377 }
378
TEST_F(SIMDSelectionCS,FirstCompiledIsSIMD16)379 TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD16)
380 {
381 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
382 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
383 brw_simd_mark_compiled(simd_state, SIMD16, not_spilled);
384
385 ASSERT_TRUE(brw_simd_any_compiled(simd_state));
386 ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD16);
387 }
388
TEST_F(SIMDSelectionCS,FirstCompiledIsSIMD32)389 TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD32)
390 {
391 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
392 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16));
393 ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32));
394 brw_simd_mark_compiled(simd_state, SIMD32, not_spilled);
395
396 ASSERT_TRUE(brw_simd_any_compiled(simd_state));
397 ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD32);
398 }
399