From 3072e6e0da5d226700a8d6d0683636b5c7e2f237 Mon Sep 17 00:00:00 2001 From: Caio Marcelo de Oliveira Filho Date: Thu, 7 Oct 2021 20:18:39 -0700 Subject: [PATCH] intel/compiler: Don't use SIMD larger than needed for workgroup 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 Part-of: --- src/intel/compiler/brw_simd_selection.c | 8 +++++++- src/intel/compiler/test_simd_selection.cpp | 12 +++++------- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/src/intel/compiler/brw_simd_selection.c b/src/intel/compiler/brw_simd_selection.c index 2ecf568ad57..0c9bf1f4fd2 100644 --- a/src/intel/compiler/brw_simd_selection.c +++ b/src/intel/compiler/brw_simd_selection.c @@ -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( diff --git a/src/intel/compiler/test_simd_selection.cpp b/src/intel/compiler/test_simd_selection.cpp index 7344c57aca4..6cbf9240a27 100644 --- a/src/intel/compiler/test_simd_selection.cpp +++ b/src/intel/compiler/test_simd_selection.cpp @@ -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);