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() : error{NULL, NULL, NULL} {
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 required_dispatch_width = 0;
49 }
50
~SIMDSelectionTest()51 ~SIMDSelectionTest() {
52 ralloc_free(mem_ctx);
53 };
54
should_compile(unsigned simd)55 bool should_compile(unsigned simd) {
56 return brw_simd_should_compile(mem_ctx, simd, devinfo, prog_data,
57 required_dispatch_width, &error[simd]);
58 }
59
60 void *mem_ctx;
61 intel_device_info *devinfo;
62 struct brw_cs_prog_data *prog_data;
63 const char *error[3];
64 unsigned required_dispatch_width;
65 };
66
67 class SIMDSelectionCS : public SIMDSelectionTest {
68 protected:
SIMDSelectionCS()69 SIMDSelectionCS() {
70 prog_data->base.stage = MESA_SHADER_COMPUTE;
71 prog_data->local_size[0] = 32;
72 prog_data->local_size[1] = 1;
73 prog_data->local_size[2] = 1;
74
75 devinfo->max_cs_workgroup_threads = 64;
76 }
77 };
78
TEST_F(SIMDSelectionCS,DefaultsToSIMD16)79 TEST_F(SIMDSelectionCS, DefaultsToSIMD16)
80 {
81 ASSERT_TRUE(should_compile(SIMD8));
82 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
83 ASSERT_TRUE(should_compile(SIMD16));
84 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
85 ASSERT_FALSE(should_compile(SIMD32));
86
87 ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
88 }
89
TEST_F(SIMDSelectionCS,TooBigFor16)90 TEST_F(SIMDSelectionCS, TooBigFor16)
91 {
92 prog_data->local_size[0] = devinfo->max_cs_workgroup_threads;
93 prog_data->local_size[1] = 32;
94 prog_data->local_size[2] = 1;
95
96 ASSERT_FALSE(should_compile(SIMD8));
97 ASSERT_FALSE(should_compile(SIMD16));
98 ASSERT_TRUE(should_compile(SIMD32));
99 brw_simd_mark_compiled(SIMD32, prog_data, spilled);
100
101 ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
102 }
103
TEST_F(SIMDSelectionCS,WorkgroupSize1)104 TEST_F(SIMDSelectionCS, WorkgroupSize1)
105 {
106 prog_data->local_size[0] = 1;
107 prog_data->local_size[1] = 1;
108 prog_data->local_size[2] = 1;
109
110 ASSERT_TRUE(should_compile(SIMD8));
111 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
112 ASSERT_FALSE(should_compile(SIMD16));
113 ASSERT_FALSE(should_compile(SIMD32));
114
115 ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
116 }
117
TEST_F(SIMDSelectionCS,WorkgroupSize8)118 TEST_F(SIMDSelectionCS, WorkgroupSize8)
119 {
120 prog_data->local_size[0] = 8;
121 prog_data->local_size[1] = 1;
122 prog_data->local_size[2] = 1;
123
124 ASSERT_TRUE(should_compile(SIMD8));
125 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
126 ASSERT_FALSE(should_compile(SIMD16));
127 ASSERT_FALSE(should_compile(SIMD32));
128
129 ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
130 }
131
TEST_F(SIMDSelectionCS,WorkgroupSizeVariable)132 TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
133 {
134 prog_data->local_size[0] = 0;
135 prog_data->local_size[1] = 0;
136 prog_data->local_size[2] = 0;
137
138 ASSERT_TRUE(should_compile(SIMD8));
139 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
140 ASSERT_TRUE(should_compile(SIMD16));
141 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
142 ASSERT_TRUE(should_compile(SIMD32));
143 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
144
145 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
146
147 const unsigned wg_8_1_1[] = { 8, 1, 1 };
148 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
149
150 const unsigned wg_16_1_1[] = { 16, 1, 1 };
151 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16);
152
153 const unsigned wg_32_1_1[] = { 32, 1, 1 };
154 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16);
155 }
156
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableSpilled)157 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled)
158 {
159 prog_data->local_size[0] = 0;
160 prog_data->local_size[1] = 0;
161 prog_data->local_size[2] = 0;
162
163 ASSERT_TRUE(should_compile(SIMD8));
164 brw_simd_mark_compiled(SIMD8, prog_data, spilled);
165 ASSERT_TRUE(should_compile(SIMD16));
166 brw_simd_mark_compiled(SIMD16, prog_data, spilled);
167 ASSERT_TRUE(should_compile(SIMD32));
168 brw_simd_mark_compiled(SIMD32, prog_data, spilled);
169
170 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
171
172 const unsigned wg_8_1_1[] = { 8, 1, 1 };
173 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
174
175 const unsigned wg_16_1_1[] = { 16, 1, 1 };
176 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8);
177
178 const unsigned wg_32_1_1[] = { 32, 1, 1 };
179 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8);
180 }
181
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD8)182 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8)
183 {
184 prog_data->local_size[0] = 0;
185 prog_data->local_size[1] = 0;
186 prog_data->local_size[2] = 0;
187
188 ASSERT_TRUE(should_compile(SIMD8));
189 ASSERT_TRUE(should_compile(SIMD16));
190 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
191 ASSERT_TRUE(should_compile(SIMD32));
192 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
193
194 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32);
195
196 const unsigned wg_8_1_1[] = { 8, 1, 1 };
197 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD16);
198
199 const unsigned wg_16_1_1[] = { 16, 1, 1 };
200 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16);
201
202 const unsigned wg_32_1_1[] = { 32, 1, 1 };
203 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16);
204 }
205
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD16)206 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD16)
207 {
208 prog_data->local_size[0] = 0;
209 prog_data->local_size[1] = 0;
210 prog_data->local_size[2] = 0;
211
212 ASSERT_TRUE(should_compile(SIMD8));
213 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
214 ASSERT_TRUE(should_compile(SIMD16));
215 ASSERT_TRUE(should_compile(SIMD32));
216 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
217
218 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD32);
219
220 const unsigned wg_8_1_1[] = { 8, 1, 1 };
221 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
222
223 const unsigned wg_16_1_1[] = { 16, 1, 1 };
224 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8);
225
226 const unsigned wg_32_1_1[] = { 32, 1, 1 };
227 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8);
228 }
229
TEST_F(SIMDSelectionCS,WorkgroupSizeVariableNoSIMD8NoSIMD16)230 TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16)
231 {
232 prog_data->local_size[0] = 0;
233 prog_data->local_size[1] = 0;
234 prog_data->local_size[2] = 0;
235
236 ASSERT_TRUE(should_compile(SIMD8));
237 ASSERT_TRUE(should_compile(SIMD16));
238 ASSERT_TRUE(should_compile(SIMD32));
239 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
240
241 ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32);
242
243 const unsigned wg_8_1_1[] = { 8, 1, 1 };
244 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD32);
245
246 const unsigned wg_16_1_1[] = { 16, 1, 1 };
247 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD32);
248
249 const unsigned wg_32_1_1[] = { 32, 1, 1 };
250 ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD32);
251 }
252
TEST_F(SIMDSelectionCS,SpillAtSIMD8)253 TEST_F(SIMDSelectionCS, SpillAtSIMD8)
254 {
255 ASSERT_TRUE(should_compile(SIMD8));
256 brw_simd_mark_compiled(SIMD8, prog_data, spilled);
257 ASSERT_FALSE(should_compile(SIMD16));
258 ASSERT_FALSE(should_compile(SIMD32));
259
260 ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
261 }
262
TEST_F(SIMDSelectionCS,SpillAtSIMD16)263 TEST_F(SIMDSelectionCS, SpillAtSIMD16)
264 {
265 ASSERT_TRUE(should_compile(SIMD8));
266 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
267 ASSERT_TRUE(should_compile(SIMD16));
268 brw_simd_mark_compiled(SIMD16, prog_data, spilled);
269 ASSERT_FALSE(should_compile(SIMD32));
270
271 ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
272 }
273
TEST_F(SIMDSelectionCS,EnvironmentVariable32)274 TEST_F(SIMDSelectionCS, EnvironmentVariable32)
275 {
276 intel_debug |= DEBUG_DO32;
277
278 ASSERT_TRUE(should_compile(SIMD8));
279 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
280 ASSERT_TRUE(should_compile(SIMD16));
281 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
282 ASSERT_TRUE(should_compile(SIMD32));
283 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
284
285 ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
286 }
287
TEST_F(SIMDSelectionCS,EnvironmentVariable32ButSpills)288 TEST_F(SIMDSelectionCS, EnvironmentVariable32ButSpills)
289 {
290 intel_debug |= DEBUG_DO32;
291
292 ASSERT_TRUE(should_compile(SIMD8));
293 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
294 ASSERT_TRUE(should_compile(SIMD16));
295 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
296 ASSERT_TRUE(should_compile(SIMD32));
297 brw_simd_mark_compiled(SIMD32, prog_data, spilled);
298
299 ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
300 }
301
TEST_F(SIMDSelectionCS,Require8)302 TEST_F(SIMDSelectionCS, Require8)
303 {
304 required_dispatch_width = 8;
305
306 ASSERT_TRUE(should_compile(SIMD8));
307 brw_simd_mark_compiled(SIMD8, prog_data, not_spilled);
308 ASSERT_FALSE(should_compile(SIMD16));
309 ASSERT_FALSE(should_compile(SIMD32));
310
311 ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
312 }
313
TEST_F(SIMDSelectionCS,Require8ErrorWhenNotCompile)314 TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile)
315 {
316 required_dispatch_width = 8;
317
318 ASSERT_TRUE(should_compile(SIMD8));
319 ASSERT_FALSE(should_compile(SIMD16));
320 ASSERT_FALSE(should_compile(SIMD32));
321
322 ASSERT_EQ(brw_simd_select(prog_data), -1);
323 }
324
TEST_F(SIMDSelectionCS,Require16)325 TEST_F(SIMDSelectionCS, Require16)
326 {
327 required_dispatch_width = 16;
328
329 ASSERT_FALSE(should_compile(SIMD8));
330 ASSERT_TRUE(should_compile(SIMD16));
331 brw_simd_mark_compiled(SIMD16, prog_data, not_spilled);
332 ASSERT_FALSE(should_compile(SIMD32));
333
334 ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
335 }
336
TEST_F(SIMDSelectionCS,Require16ErrorWhenNotCompile)337 TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile)
338 {
339 required_dispatch_width = 16;
340
341 ASSERT_FALSE(should_compile(SIMD8));
342 ASSERT_TRUE(should_compile(SIMD16));
343 ASSERT_FALSE(should_compile(SIMD32));
344
345 ASSERT_EQ(brw_simd_select(prog_data), -1);
346 }
347
TEST_F(SIMDSelectionCS,Require32)348 TEST_F(SIMDSelectionCS, Require32)
349 {
350 required_dispatch_width = 32;
351
352 ASSERT_FALSE(should_compile(SIMD8));
353 ASSERT_FALSE(should_compile(SIMD16));
354 ASSERT_TRUE(should_compile(SIMD32));
355 brw_simd_mark_compiled(SIMD32, prog_data, not_spilled);
356
357 ASSERT_EQ(brw_simd_select(prog_data), SIMD32);
358 }
359
TEST_F(SIMDSelectionCS,Require32ErrorWhenNotCompile)360 TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile)
361 {
362 required_dispatch_width = 32;
363
364 ASSERT_FALSE(should_compile(SIMD8));
365 ASSERT_FALSE(should_compile(SIMD16));
366 ASSERT_TRUE(should_compile(SIMD32));
367
368 ASSERT_EQ(brw_simd_select(prog_data), -1);
369 }
370