fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 00:58:05 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

intel/compiler: Create a struct to hold SIMD selection state

Browse source 
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 <lionel.g.landwerlin@intel.com>
Reviewed-by: Ivan Briano <ivan.briano@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
This commit is contained in:
Caio Oliveira 2022-11-08 01:47:50 -08:00 • committed by Marge Bot
parent 8cda6cd774
commit a0580dadfd
5 changed files with 208 additions and 192 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

21
src/intel/compiler/brw_fs.cpp
View file

@ -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<fs_visitor> 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;
}

42
src/intel/compiler/brw_mesh.cpp
View file

@ -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<fs_visitor> 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<fs_visitor> 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;
}

24
src/intel/compiler/brw_private.h
View file

@ -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,

80
src/intel/compiler/brw_simd_selection.cpp
View file

@ -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<struct brw_cs_prog_data *>(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);
}

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

@ -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);
}