fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-23 04:18:14 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1
mesa/src/intel/compiler/brw/test_simd_selection.cpp

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

380 lines
12 KiB
C++
Raw Normal View History

intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
/*
* Copyright © 2021 Intel Corporation
intel/compiler: Use SPDX annotations Minor adjustments to formatting of the copyright line, but keep dates and holders. "Authors" entries that could be obtained via Git logs were also removed. The license in brw_disasm.c and elk_disasm.c don't match directly any SPDX pattern I could find, so kept as is. Reviewed-by: Ian Romanick <ian.d.romanick@intel.com> Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39503>
2026-01-23 17:03:58 -08:00
* SPDX-License-Identifier: MIT
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
*/
#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 <gtest/gtest.h>
enum {
SIMD8 = 0,
SIMD16 = 1,
SIMD32 = 2,
};
const bool spilled = true;
const bool not_spilled = false;
class SIMDSelectionTest : public ::testing::Test {
protected:
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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{
.devinfo = devinfo,
.prog_data = prog_data,
}
{
intel: Remove brw_ prefix from process debug function Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27644>
2024-01-19 18:56:18 -08:00
process_intel_debug_variable();
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
~SIMDSelectionTest() {
ralloc_free(mem_ctx);
};
void *mem_ctx;
intel_device_info *devinfo;
struct brw_cs_prog_data *prog_data;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
brw_simd_selection_state simd_state;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
};
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)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
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;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, WorkgroupSize1)
{
prog_data->local_size[0] = 1;
prog_data->local_size[1] = 1;
prog_data->local_size[2] = 1;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, WorkgroupSize8)
{
prog_data->local_size[0] = 8;
prog_data->local_size[1] = 1;
prog_data->local_size[2] = 1;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, WorkgroupSizeVariable)
{
prog_data->local_size[0] = 0;
prog_data->local_size[1] = 0;
prog_data->local_size[2] = 0;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
const unsigned wg_8_1_1[] = { 8, 1, 1 };
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 <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-07 20:18:39 -07:00
ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), SIMD8);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, WorkgroupSizeVariableSpilled)
{
prog_data->local_size[0] = 0;
prog_data->local_size[1] = 0;
prog_data->local_size[2] = 0;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
ASSERT_EQ(prog_data->prog_mask, 1u << SIMD8 | 1u << SIMD16 | 1u << SIMD32);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
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;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
ASSERT_EQ(prog_data->prog_mask, 1u << SIMD16 | 1u << SIMD32);
const unsigned wg_8_1_1[] = { 8, 1, 1 };
intel/brw/xe3+: Optimize CS/TASK/MESH compile time optimistically assuming SIMD32. This is similar in principle to the previous commit "intel/brw/xe3+: brw_compile_fs() implementation for Xe3+." but applied to compute-like shader stages. It changes the implementation of brw_compile_cs/task/mesh() to reduce compile time and take advantage of wider dispatch modes more aggressively than the original logic, since as of Xe3 SIMD32 builds succeed without spills in most cases thanks to VRT. The new "optimistic" SIMD selection logic starts with the SIMD width that is potentially highest performance and only compiles additional narrower variants if that fails (typically due to spilling), while the old "pessimistic" logic did the opposite: It started with the narrowest SIMD width and compiled additional variants with increasing register pressure until one of them failed to compile. In typical non-spilling cases where we formerly compiled SIMD16 and SIMD32 variants of the same compute shader, this change will halve the number of backend compilations required to build it. XXX - Possibly don't do this in cases with variable workgroup size until effect on runtime performance can be measured directly. Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com> v2: Don't do this for now in cases with variable workgroup size, still compile every possible variant in such cases. Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32664>
2024-09-24 16:53:25 -07:00
ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), -1);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
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;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
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;
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
ASSERT_EQ(prog_data->prog_mask, 1u << SIMD32);
const unsigned wg_8_1_1[] = { 8, 1, 1 };
intel/brw/xe3+: Optimize CS/TASK/MESH compile time optimistically assuming SIMD32. This is similar in principle to the previous commit "intel/brw/xe3+: brw_compile_fs() implementation for Xe3+." but applied to compute-like shader stages. It changes the implementation of brw_compile_cs/task/mesh() to reduce compile time and take advantage of wider dispatch modes more aggressively than the original logic, since as of Xe3 SIMD32 builds succeed without spills in most cases thanks to VRT. The new "optimistic" SIMD selection logic starts with the SIMD width that is potentially highest performance and only compiles additional narrower variants if that fails (typically due to spilling), while the old "pessimistic" logic did the opposite: It started with the narrowest SIMD width and compiled additional variants with increasing register pressure until one of them failed to compile. In typical non-spilling cases where we formerly compiled SIMD16 and SIMD32 variants of the same compute shader, this change will halve the number of backend compilations required to build it. XXX - Possibly don't do this in cases with variable workgroup size until effect on runtime performance can be measured directly. Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com> v2: Don't do this for now in cases with variable workgroup size, still compile every possible variant in such cases. Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32664>
2024-09-24 16:53:25 -07:00
ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_8_1_1), -1);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
const unsigned wg_16_1_1[] = { 16, 1, 1 };
intel/brw/xe3+: Optimize CS/TASK/MESH compile time optimistically assuming SIMD32. This is similar in principle to the previous commit "intel/brw/xe3+: brw_compile_fs() implementation for Xe3+." but applied to compute-like shader stages. It changes the implementation of brw_compile_cs/task/mesh() to reduce compile time and take advantage of wider dispatch modes more aggressively than the original logic, since as of Xe3 SIMD32 builds succeed without spills in most cases thanks to VRT. The new "optimistic" SIMD selection logic starts with the SIMD width that is potentially highest performance and only compiles additional narrower variants if that fails (typically due to spilling), while the old "pessimistic" logic did the opposite: It started with the narrowest SIMD width and compiled additional variants with increasing register pressure until one of them failed to compile. In typical non-spilling cases where we formerly compiled SIMD16 and SIMD32 variants of the same compute shader, this change will halve the number of backend compilations required to build it. XXX - Possibly don't do this in cases with variable workgroup size until effect on runtime performance can be measured directly. Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com> v2: Don't do this for now in cases with variable workgroup size, still compile every possible variant in such cases. Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32664>
2024-09-24 16:53:25 -07:00
ASSERT_EQ(brw_simd_select_for_workgroup_size(devinfo, prog_data, wg_16_1_1), -1);
intel/compiler: Use SIMD selection helpers for variable workgroup size Variable workgroup size works by compiling as much SIMD variants as possible and then selecting the right one during dispatch (when the actual workgroup size is passed to us). Instead of replicating the logic in a separate function, reuse the same logic for regular SIMD selection. And move function for that together with the remaining simd selection functions. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-11 07:49:40 -07:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, SpillAtSIMD8)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, SpillAtSIMD16)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, EnvironmentVariable32)
{
intel: Switch uint64_t intel_debug to a bitset We are reaching our limit of adding flags to intel_debug (apporaching 64 flags). Switch intel_debug to a bitset, which gives us almost "unlimited" bits to use in the future. v2(Michael Cheng): Fixed a few ci errors Signed-off-by: Michael Cheng <michael.cheng@intel.com> Reviewed-by: Casey Bowman <casey.g.bowman@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34596>
2025-04-18 14:48:08 -07:00
BITSET_SET(intel_debug, DEBUG_DO32);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, EnvironmentVariable32ButSpills)
{
intel: Switch uint64_t intel_debug to a bitset We are reaching our limit of adding flags to intel_debug (apporaching 64 flags). Switch intel_debug to a bitset, which gives us almost "unlimited" bits to use in the future. v2(Michael Cheng): Fixed a few ci errors Signed-off-by: Michael Cheng <michael.cheng@intel.com> Reviewed-by: Casey Bowman <casey.g.bowman@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34596>
2025-04-18 14:48:08 -07:00
BITSET_SET(intel_debug, DEBUG_DO32);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require8)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 8;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD8);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require8ErrorWhenNotCompile)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 8;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), -1);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require16)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 16;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD16);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require16ErrorWhenNotCompile)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 16;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), -1);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require32)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 32;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), SIMD32);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
TEST_F(SIMDSelectionCS, Require32ErrorWhenNotCompile)
{
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
simd_state.required_width = 32;
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
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));
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
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 <lionel.g.landwerlin@intel.com> Reviewed-by: Ivan Briano <ivan.briano@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19601>
2022-11-08 01:47:50 -08:00
ASSERT_EQ(brw_simd_select(simd_state), -1);
intel/compiler: Add helpers to select SIMD for compute shaders Clean up the logic and move it to functions that work with prog_data attributes to select the right SIMD. This shouldn't change any behavior compared to the original. Having it extracted will allow reuse by Task/Mesh and make it easier to write tests. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13249>
2021-10-06 22:37:42 -07:00
}
intel/compiler: Keep track of compiled/spilled in brw_simd_selection_state We still update the cs_prog_data, but don't rely on it for this state anymore. This will allow use the SIMD selector with shaders that don't use cs_prog_data. 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>
2022-11-08 03:38:18 -08:00
TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD8)
{
ASSERT_TRUE(brw_simd_should_compile(simd_state, SIMD8));
brw_simd_mark_compiled(simd_state, SIMD8, not_spilled);
ASSERT_TRUE(brw_simd_any_compiled(simd_state));
ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD8);
}
TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD16)
{
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_any_compiled(simd_state));
ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD16);
}
TEST_F(SIMDSelectionCS, FirstCompiledIsSIMD32)
{
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_TRUE(brw_simd_any_compiled(simd_state));
ASSERT_EQ(brw_simd_first_compiled(simd_state), SIMD32);
}
Reference in a new issue Copy permalink