From a0580dadfd20f4b27852bc10e2f38191cae05197 Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Tue, 8 Nov 2022 01:47:50 -0800 Subject: [PATCH] intel/compiler: Create a struct to hold SIMD selection state This is a preparation to decouple the storage of what SIMDs compiled/spilled from the cs_prog_data. This will allow reuse of SIMD selection code by Bindless Shaders. And since we have a struct now, move the error array there so reduce the boilerplate of the users. Reviewed-by: Lionel Landwerlin Reviewed-by: Ivan Briano Part-of: --- src/intel/compiler/brw_fs.cpp | 21 +- src/intel/compiler/brw_mesh.cpp | 42 ++-- src/intel/compiler/brw_private.h | 24 ++- src/intel/compiler/brw_simd_selection.cpp | 80 +++---- src/intel/compiler/test_simd_selection.cpp | 233 ++++++++++----------- 5 files changed, 208 insertions(+), 192 deletions(-) diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index f52f8cabf63..c661a8c2a14 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/src/intel/compiler/brw_fs.cpp @@ -7806,15 +7806,17 @@ brw_compile_cs(const struct brw_compiler *compiler, prog_data->local_size[2] = nir->info.workgroup_size[2]; } - const unsigned required_dispatch_width = - brw_required_dispatch_width(&nir->info); + brw_simd_selection_state simd_state{ + .mem_ctx = mem_ctx, + .devinfo = compiler->devinfo, + .prog_data = prog_data, + .required_width = brw_required_dispatch_width(&nir->info), + }; std::unique_ptr v[3]; - const char *error[3] = {0}; for (unsigned simd = 0; simd < 3; simd++) { - if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, prog_data, - required_dispatch_width, &error[simd])) + if (!brw_simd_should_compile(simd_state, simd)) continue; const unsigned dispatch_width = 8u << simd; @@ -7847,9 +7849,9 @@ brw_compile_cs(const struct brw_compiler *compiler, if (v[simd]->run_cs(allow_spilling)) { cs_fill_push_const_info(compiler->devinfo, prog_data); - brw_simd_mark_compiled(simd, prog_data, v[simd]->spilled_any_registers); + brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers); } else { - error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); + simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); if (simd > 0) { brw_shader_perf_log(compiler, params->log_data, "SIMD%u shader failed to compile: %s\n", @@ -7858,10 +7860,11 @@ brw_compile_cs(const struct brw_compiler *compiler, } } - const int selected_simd = brw_simd_select(prog_data); + const int selected_simd = brw_simd_select(simd_state); if (selected_simd < 0) { params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n", - error[0], error[1], error[2]);; + simd_state.error[0], simd_state.error[1], + simd_state.error[2]); return NULL; } diff --git a/src/intel/compiler/brw_mesh.cpp b/src/intel/compiler/brw_mesh.cpp index f7c4b9e17fe..64bcb52a015 100644 --- a/src/intel/compiler/brw_mesh.cpp +++ b/src/intel/compiler/brw_mesh.cpp @@ -265,15 +265,17 @@ brw_compile_task(const struct brw_compiler *compiler, prog_data->uses_drawid = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID); - const unsigned required_dispatch_width = - brw_required_dispatch_width(&nir->info); + brw_simd_selection_state simd_state{ + .mem_ctx = mem_ctx, + .devinfo = compiler->devinfo, + .prog_data = &prog_data->base, + .required_width = brw_required_dispatch_width(&nir->info), + }; std::unique_ptr v[3]; - const char *error[3] = {0}; for (unsigned simd = 0; simd < 3; simd++) { - if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, &prog_data->base, - required_dispatch_width, &error[simd])) + if (!brw_simd_should_compile(simd_state, simd)) continue; const unsigned dispatch_width = 8 << simd; @@ -301,15 +303,16 @@ brw_compile_task(const struct brw_compiler *compiler, const bool allow_spilling = !prog_data->base.prog_mask; if (v[simd]->run_task(allow_spilling)) - brw_simd_mark_compiled(simd, &prog_data->base, v[simd]->spilled_any_registers); + brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers); else - error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); + simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); } - int selected_simd = brw_simd_select(&prog_data->base); + int selected_simd = brw_simd_select(simd_state); if (selected_simd < 0) { params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n", - error[0], error[1], error[2]);; + simd_state.error[0], simd_state.error[1], + simd_state.error[2]); return NULL; } @@ -761,15 +764,17 @@ brw_compile_mesh(const struct brw_compiler *compiler, brw_compute_mue_map(nir, &prog_data->map); brw_nir_lower_mue_outputs(nir, &prog_data->map); - const unsigned required_dispatch_width = - brw_required_dispatch_width(&nir->info); + brw_simd_selection_state simd_state{ + .mem_ctx = mem_ctx, + .devinfo = compiler->devinfo, + .prog_data = &prog_data->base, + .required_width = brw_required_dispatch_width(&nir->info), + }; std::unique_ptr v[3]; - const char *error[3] = {0}; for (int simd = 0; simd < 3; simd++) { - if (!brw_simd_should_compile(mem_ctx, simd, compiler->devinfo, &prog_data->base, - required_dispatch_width, &error[simd])) + if (!brw_simd_should_compile(simd_state, simd)) continue; const unsigned dispatch_width = 8 << simd; @@ -809,15 +814,16 @@ brw_compile_mesh(const struct brw_compiler *compiler, const bool allow_spilling = !prog_data->base.prog_mask; if (v[simd]->run_mesh(allow_spilling)) - brw_simd_mark_compiled(simd, &prog_data->base, v[simd]->spilled_any_registers); + brw_simd_mark_compiled(simd_state, simd, v[simd]->spilled_any_registers); else - error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); + simd_state.error[simd] = ralloc_strdup(mem_ctx, v[simd]->fail_msg); } - int selected_simd = brw_simd_select(&prog_data->base); + int selected_simd = brw_simd_select(simd_state); if (selected_simd < 0) { params->error_str = ralloc_asprintf(mem_ctx, "Can't compile shader: %s, %s and %s.\n", - error[0], error[1], error[2]);; + simd_state.error[0], simd_state.error[1], + simd_state.error[2]);; return NULL; } diff --git a/src/intel/compiler/brw_private.h b/src/intel/compiler/brw_private.h index 70b6fd93179..6f1374e53b5 100644 --- a/src/intel/compiler/brw_private.h +++ b/src/intel/compiler/brw_private.h @@ -29,18 +29,22 @@ unsigned brw_required_dispatch_width(const struct shader_info *info); -bool brw_simd_should_compile(void *mem_ctx, - unsigned simd, - const struct intel_device_info *devinfo, - struct brw_cs_prog_data *prog_data, - unsigned required_dispatch_width, - const char **error); +struct brw_simd_selection_state { + void *mem_ctx; + const struct intel_device_info *devinfo; -void brw_simd_mark_compiled(unsigned simd, - struct brw_cs_prog_data *prog_data, - bool spilled); + struct brw_cs_prog_data *prog_data; -int brw_simd_select(const struct brw_cs_prog_data *prog_data); + unsigned required_width; + + const char *error[3]; +}; + +bool brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd); + +void brw_simd_mark_compiled(brw_simd_selection_state &state, unsigned simd, bool spilled); + +int brw_simd_select(const brw_simd_selection_state &state); int brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo, const struct brw_cs_prog_data *prog_data, diff --git a/src/intel/compiler/brw_simd_selection.cpp b/src/intel/compiler/brw_simd_selection.cpp index 078ff0a9ced..aa16674cac4 100644 --- a/src/intel/compiler/brw_simd_selection.cpp +++ b/src/intel/compiler/brw_simd_selection.cpp @@ -47,16 +47,11 @@ test_bit(unsigned mask, unsigned bit) { } bool -brw_simd_should_compile(void *mem_ctx, - unsigned simd, - const struct intel_device_info *devinfo, - struct brw_cs_prog_data *prog_data, - unsigned required, - const char **error) - +brw_simd_should_compile(brw_simd_selection_state &state, + unsigned simd) { + struct brw_cs_prog_data *prog_data = state.prog_data; assert(!test_bit(prog_data->prog_mask, simd)); - assert(error); const unsigned width = 8u << simd; @@ -68,8 +63,8 @@ brw_simd_should_compile(void *mem_ctx, if (!workgroup_size_variable) { if (test_bit(prog_data->prog_spilled, simd)) { - *error = ralloc_asprintf( - mem_ctx, "SIMD%u skipped because would spill", width); + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u skipped because would spill", width); return false; } @@ -77,26 +72,26 @@ brw_simd_should_compile(void *mem_ctx, prog_data->local_size[1] * prog_data->local_size[2]; - unsigned max_threads = devinfo->max_cs_workgroup_threads; + unsigned max_threads = state.devinfo->max_cs_workgroup_threads; - if (required && required != width) { - *error = ralloc_asprintf( - mem_ctx, "SIMD%u skipped because required dispatch width is %u", - width, required); + if (state.required_width && state.required_width != width) { + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u skipped because required dispatch width is %u", + width, state.required_width); return false; } 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", + state.error[simd] = ralloc_asprintf( + state.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( - mem_ctx, "SIMD%u can't fit all %u invocations in %u threads", + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u can't fit all %u invocations in %u threads", width, workgroup_size, max_threads); return false; } @@ -107,23 +102,23 @@ brw_simd_should_compile(void *mem_ctx, */ if (width == 32) { if (!INTEL_DEBUG(DEBUG_DO32) && prog_data->prog_mask) { - *error = ralloc_strdup( - mem_ctx, "SIMD32 skipped because not required"); + state.error[simd] = ralloc_strdup( + state.mem_ctx, "SIMD32 skipped because not required"); return false; } } } if (width == 32 && prog_data->base.ray_queries > 0) { - *error = ralloc_asprintf( - mem_ctx, "SIMD%u skipped because of ray queries", + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u skipped because of ray queries", width); return false; } if (width == 32 && prog_data->uses_btd_stack_ids) { - *error = ralloc_asprintf( - mem_ctx, "SIMD%u skipped because of bindless shader calls", + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u skipped because of bindless shader calls", width); return false; } @@ -135,8 +130,8 @@ brw_simd_should_compile(void *mem_ctx, }; if (unlikely(env_skip[simd])) { - *error = ralloc_asprintf( - mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u", + state.error[simd] = ralloc_asprintf( + state.mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u", width, width); return false; } @@ -145,8 +140,9 @@ brw_simd_should_compile(void *mem_ctx, } void -brw_simd_mark_compiled(unsigned simd, struct brw_cs_prog_data *prog_data, bool spilled) +brw_simd_mark_compiled(brw_simd_selection_state &state, unsigned simd, bool spilled) { + struct brw_cs_prog_data *prog_data = state.prog_data; assert(!test_bit(prog_data->prog_mask, simd)); prog_data->prog_mask |= 1u << simd; @@ -159,8 +155,9 @@ brw_simd_mark_compiled(unsigned simd, struct brw_cs_prog_data *prog_data, bool s } int -brw_simd_select(const struct brw_cs_prog_data *prog_data) +brw_simd_select(const struct brw_simd_selection_state &state) { + const struct brw_cs_prog_data *prog_data = state.prog_data; assert((prog_data->prog_mask & ~0x7u) == 0); const unsigned not_spilled_mask = prog_data->prog_mask & ~prog_data->prog_spilled; @@ -182,10 +179,12 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo, { if (!sizes || (prog_data->local_size[0] == sizes[0] && prog_data->local_size[1] == sizes[1] && - prog_data->local_size[2] == sizes[2])) - return brw_simd_select(prog_data); - - void *mem_ctx = ralloc_context(NULL); + prog_data->local_size[2] == sizes[2])) { + const brw_simd_selection_state simd_state{ + .prog_data = const_cast(prog_data), + }; + return brw_simd_select(simd_state); + } struct brw_cs_prog_data cloned = *prog_data; for (unsigned i = 0; i < 3; i++) @@ -194,20 +193,25 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo, cloned.prog_mask = 0; cloned.prog_spilled = 0; - const char *error[3] = {0}; + void *mem_ctx = ralloc_context(NULL); + + brw_simd_selection_state simd_state{ + .mem_ctx = mem_ctx, + .devinfo = devinfo, + .prog_data = &cloned, + }; for (unsigned simd = 0; simd < 3; simd++) { /* We are not recompiling, so use original results of prog_mask and * prog_spilled as they will already contain all possible compilations. */ - if (brw_simd_should_compile(mem_ctx, simd, devinfo, &cloned, - 0 /* required_dispatch_width */, &error[simd]) && + if (brw_simd_should_compile(simd_state, simd) && test_bit(prog_data->prog_mask, simd)) { - brw_simd_mark_compiled(simd, &cloned, test_bit(prog_data->prog_spilled, simd)); + brw_simd_mark_compiled(simd_state, simd, test_bit(prog_data->prog_spilled, simd)); } } ralloc_free(mem_ctx); - return brw_simd_select(&cloned); + return brw_simd_select(simd_state); } diff --git a/src/intel/compiler/test_simd_selection.cpp b/src/intel/compiler/test_simd_selection.cpp index 5783c328772..c8a6f02c521 100644 --- a/src/intel/compiler/test_simd_selection.cpp +++ b/src/intel/compiler/test_simd_selection.cpp @@ -41,27 +41,26 @@ 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() + : mem_ctx(ralloc_context(NULL)) + , devinfo(rzalloc(mem_ctx, intel_device_info)) + , prog_data(rzalloc(mem_ctx, struct brw_cs_prog_data)) + , simd_state{ + .mem_ctx = mem_ctx, + .devinfo = devinfo, + .prog_data = prog_data, + } + { } ~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; + brw_simd_selection_state simd_state; }; class SIMDSelectionCS : public SIMDSelectionTest { @@ -78,13 +77,13 @@ protected: 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD16); + ASSERT_EQ(brw_simd_select(simd_state), SIMD16); } TEST_F(SIMDSelectionCS, TooBigFor16) @@ -93,12 +92,12 @@ TEST_F(SIMDSelectionCS, TooBigFor16) 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_FALSE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, spilled); - ASSERT_EQ(brw_simd_select(prog_data), SIMD32); + ASSERT_EQ(brw_simd_select(simd_state), SIMD32); } TEST_F(SIMDSelectionCS, WorkgroupSize1) @@ -107,12 +106,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSize1) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD8); + ASSERT_EQ(brw_simd_select(simd_state), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSize8) @@ -121,12 +120,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSize8) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD8); + ASSERT_EQ(brw_simd_select(simd_state), SIMD8); } TEST_F(SIMDSelectionCS, WorkgroupSizeVariable) @@ -135,12 +134,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariable) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32); @@ -160,12 +159,12 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32); @@ -185,11 +184,11 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32); @@ -209,11 +208,11 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD16) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD32); @@ -233,10 +232,10 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32); @@ -252,118 +251,118 @@ TEST_F(SIMDSelectionCS, WorkgroupSizeVariableNoSIMD8NoSIMD16) 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD8); + ASSERT_EQ(brw_simd_select(simd_state), 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD8); + ASSERT_EQ(brw_simd_select(simd_state), 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); - ASSERT_EQ(brw_simd_select(prog_data), SIMD32); + ASSERT_EQ(brw_simd_select(simd_state), 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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, spilled); - ASSERT_EQ(brw_simd_select(prog_data), SIMD16); + ASSERT_EQ(brw_simd_select(simd_state), SIMD16); } TEST_F(SIMDSelectionCS, Require8) { - required_dispatch_width = 8; + simd_state.required_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_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + brw_simd_mark_compiled(simd_state, SIMD8, not_spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD8); + ASSERT_EQ(brw_simd_select(simd_state), SIMD8); } TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile) { - required_dispatch_width = 8; + simd_state.required_width = 8; - ASSERT_TRUE(should_compile(SIMD8)); - ASSERT_FALSE(should_compile(SIMD16)); - ASSERT_FALSE(should_compile(SIMD32)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), -1); + ASSERT_EQ(brw_simd_select(simd_state), -1); } TEST_F(SIMDSelectionCS, Require16) { - required_dispatch_width = 16; + simd_state.required_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_FALSE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + brw_simd_mark_compiled(simd_state, SIMD16, not_spilled); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), SIMD16); + ASSERT_EQ(brw_simd_select(simd_state), SIMD16); } TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile) { - required_dispatch_width = 16; + simd_state.required_width = 16; - ASSERT_FALSE(should_compile(SIMD8)); - ASSERT_TRUE(should_compile(SIMD16)); - ASSERT_FALSE(should_compile(SIMD32)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), -1); + ASSERT_EQ(brw_simd_select(simd_state), -1); } TEST_F(SIMDSelectionCS, Require32) { - required_dispatch_width = 32; + simd_state.required_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_FALSE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); + brw_simd_mark_compiled(simd_state, SIMD32, not_spilled); - ASSERT_EQ(brw_simd_select(prog_data), SIMD32); + ASSERT_EQ(brw_simd_select(simd_state), SIMD32); } TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile) { - required_dispatch_width = 32; + simd_state.required_width = 32; - ASSERT_FALSE(should_compile(SIMD8)); - ASSERT_FALSE(should_compile(SIMD16)); - ASSERT_TRUE(should_compile(SIMD32)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD8)); + ASSERT_FALSE(brw_simd_should_compile(simd_state, SIMD16)); + ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD32)); - ASSERT_EQ(brw_simd_select(prog_data), -1); + ASSERT_EQ(brw_simd_select(simd_state), -1); }