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 33enum { 34 SIMD8 = 0, 35 SIMD16 = 1, 36 SIMD32 = 2, 37}; 38 39const bool spilled = true; 40const bool not_spilled = false; 41 42class SIMDSelectionTest : public ::testing::Test { 43protected: 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 67class SIMDSelectionCS : public SIMDSelectionTest { 68protected: 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 79TEST_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 90TEST_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 104TEST_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 118TEST_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 132TEST_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 157TEST_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 182TEST_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 206TEST_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 230TEST_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 253TEST_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 263TEST_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 274TEST_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 288TEST_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 302TEST_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 314TEST_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 325TEST_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 337TEST_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 348TEST_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 360TEST_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