/* * Copyright © 2021 Intel Corporation * * Permission is hereby granted, free of charge, to any person obtaining a * copy of this software and associated documentation files (the "Software"), * to deal in the Software without restriction, including without limitation * the rights to use, copy, modify, merge, publish, distribute, sublicense, * and/or sell copies of the Software, and to permit persons to whom the * Software is furnished to do so, subject to the following conditions: * * The above copyright notice and this permission notice (including the next * paragraph) shall be included in all copies or substantial portions of the * Software. * * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS * IN THE SOFTWARE. */ #include "brw_private.h" #include "compiler/shader_info.h" #include "intel/dev/intel_debug.h" #include "intel/dev/intel_device_info.h" #include "util/ralloc.h" #include enum { SIMD8 = 0, SIMD16 = 1, SIMD32 = 2, }; const bool spilled = true; const bool not_spilled = false; class SIMDSelectionTest : public ::testing::Test { protected: SIMDSelectionTest() : error{NULL, NULL, NULL} { mem_ctx = ralloc_context(NULL); devinfo = rzalloc(mem_ctx, intel_device_info); prog_data = rzalloc(mem_ctx, struct brw_cs_prog_data); required_dispatch_width = 0; } ~SIMDSelectionTest() { ralloc_free(mem_ctx); }; bool should_compile(unsigned simd) { return brw_simd_should_compile(mem_ctx, simd, devinfo, prog_data, required_dispatch_width, &error[simd]); } void *mem_ctx; intel_device_info *devinfo; struct brw_cs_prog_data *prog_data; const char *error[3]; unsigned required_dispatch_width; }; class SIMDSelectionCS : public SIMDSelectionTest { protected: SIMDSelectionCS() { prog_data->base.stage = MESA_SHADER_COMPUTE; prog_data->local_size[0] = 32; prog_data->local_size[1] = 1; prog_data->local_size[2] = 1; devinfo->max_cs_workgroup_threads = 64; } }; TEST_F(SIMDSelectionCS, DefaultsToSIMD16) { ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD16); } TEST_F(SIMDSelectionCS, TooBigFor16) { prog_data->local_size[0] = devinfo->max_cs_workgroup_threads; prog_data->local_size[1] = 32; prog_data->local_size[2] = 1; ASSERT_FALSE(should_compile(SIMD8)); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, spilled); ASSERT_EQ(brw_simd_select(prog_data), SIMD32); } TEST_F(SIMDSelectionCS, WorkgroupSize1) { prog_data->local_size[0] = 1; prog_data->local_size[1] = 1; prog_data->local_size[2] = 1; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSize8) { prog_data->local_size[0] = 8; prog_data->local_size[1] = 1; prog_data->local_size[2] = 1; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariable) { prog_data->local_size[0] = 0; prog_data->local_size[1] = 0; prog_data->local_size[2] = 0; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32); const unsigned wg_8_1_1[] = { 8, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8); const unsigned wg_16_1_1[] = { 16, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16); const unsigned wg_32_1_1[] = { 32, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled) { prog_data->local_size[0] = 0; prog_data->local_size[1] = 0; prog_data->local_size[2] = 0; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, spilled); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32); const unsigned wg_8_1_1[] = { 8, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8); const unsigned wg_16_1_1[] = { 16, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8); const unsigned wg_32_1_1[] = { 32, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8) { prog_data->local_size[0] = 0; prog_data->local_size[1] = 0; prog_data->local_size[2] = 0; ASSERT_TRUE(should_compile(SIMD8)); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32); const unsigned wg_8_1_1[] = { 8, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD16); const unsigned wg_16_1_1[] = { 16, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD16); const unsigned wg_32_1_1[] = { 32, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD16); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD16) { prog_data->local_size[0] = 0; prog_data->local_size[1] = 0; prog_data->local_size[2] = 0; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD32); const unsigned wg_8_1_1[] = { 8, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8); const unsigned wg_16_1_1[] = { 16, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD8); const unsigned wg_32_1_1[] = { 32, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16) { prog_data->local_size[0] = 0; prog_data->local_size[1] = 0; prog_data->local_size[2] = 0; ASSERT_TRUE(should_compile(SIMD8)); ASSERT_TRUE(should_compile(SIMD16)); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32); const unsigned wg_8_1_1[] = { 8, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD32); const unsigned wg_16_1_1[] = { 16, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), SIMD32); const unsigned wg_32_1_1[] = { 32, 1, 1 }; ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_32_1_1), SIMD32); } TEST_F(SIMDSelectionCS, SpillAtSIMD8) { ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, spilled); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD8); } TEST_F(SIMDSelectionCS, SpillAtSIMD16) { ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, spilled); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD8); } TEST_F(SIMDSelectionCS, EnvironmentVariable32) { intel_debug |= DEBUG_DO32; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(brw_simd_select(prog_data), SIMD32); } TEST_F(SIMDSelectionCS, EnvironmentVariable32ButSpills) { intel_debug |= DEBUG_DO32; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, spilled); ASSERT_EQ(brw_simd_select(prog_data), SIMD16); } TEST_F(SIMDSelectionCS, Require8) { required_dispatch_width = 8; ASSERT_TRUE(should_compile(SIMD8)); brw_simd_mark_compiled(SIMD8, prog_data, not_spilled); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD8); } TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile) { required_dispatch_width = 8; ASSERT_TRUE(should_compile(SIMD8)); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), -1); } TEST_F(SIMDSelectionCS, Require16) { required_dispatch_width = 16; ASSERT_FALSE(should_compile(SIMD8)); ASSERT_TRUE(should_compile(SIMD16)); brw_simd_mark_compiled(SIMD16, prog_data, not_spilled); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), SIMD16); } TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile) { required_dispatch_width = 16; ASSERT_FALSE(should_compile(SIMD8)); ASSERT_TRUE(should_compile(SIMD16)); ASSERT_FALSE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), -1); } TEST_F(SIMDSelectionCS, Require32) { required_dispatch_width = 32; ASSERT_FALSE(should_compile(SIMD8)); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_TRUE(should_compile(SIMD32)); brw_simd_mark_compiled(SIMD32, prog_data, not_spilled); ASSERT_EQ(brw_simd_select(prog_data), SIMD32); } TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile) { required_dispatch_width = 32; ASSERT_FALSE(should_compile(SIMD8)); ASSERT_FALSE(should_compile(SIMD16)); ASSERT_TRUE(should_compile(SIMD32)); ASSERT_EQ(brw_simd_select(prog_data), -1); }