fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-02-21 14:20:29 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 4

intel/compiler: Don't use SIMD larger than needed for workgroup

Browse source 
Unless we are combining multiple workgroups in the same HW thread,
there's no advantage of using SIMD16 when SIMD8 already fits the
entire workgroup.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
This commit is contained in:
Caio Marcelo de Oliveira Filho 2021-10-07 20:18:39 -07:00 • committed by Marge Bot
parent 4e7b71e00c
commit 3072e6e0da
2 changed files with 12 additions and 8 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
src/intel/compiler/brw_simd_selection.c
View file

@ -93,7 +93,13 @@ brw_simd_should_compile(void *mem_ctx,
return false;
}
/* TODO: Ignore SIMD larger than workgroup if previous SIMD already passed. */
if (simd > 0 && test_bit(prog_data->prog_mask, simd - 1) &&
workgroup_size <= (width / 2)) {
*error = ralloc_asprintf(
mem_ctx, "SIMD%u skipped because workgroup size %u already fits in SIMD%u",
width, workgroup_size, width / 2);
return false;
}
if (DIV_ROUND_UP(workgroup_size, width) > max_threads) {
*error = ralloc_asprintf(

12
src/intel/compiler/test_simd_selection.cpp
View file

@ -109,11 +109,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSize1)
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(SIMD16));
ASSERT_FALSE(should_compile(SIMD32));
ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
}
TEST_F(SIMDSelectionCS, WorkgroupSize8)
@ -124,11 +123,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSize8)
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(SIMD16));
ASSERT_FALSE(should_compile(SIMD32));
ASSERT_EQ(brw_simd_select(prog_data), SIMD16);
ASSERT_EQ(brw_simd_select(prog_data), SIMD8);
}
TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
@ -147,7 +145,7 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
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), SIMD16);
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);