← Back to branch summary

~mmach/netext73/mesa-haswell

« back to all changes in this revision

Viewing changes to src/intel/compiler/test_simd_selection.cpp

  • Committer: mmach
  • Date: 2022-09-22 19:56:13 UTC
  • Revision ID: netbit73@gmail.com-20220922195613-wtik9mmy20tmor0i
2022-09-22 21:17:09

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/intel/compiler/test_simd_selection.cpp
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:
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
 
 
51
 
   ~SIMDSelectionTest() {
52
 
      ralloc_free(mem_ctx);
53
 
   };
54
 
 
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:
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
 
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
 
}