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

intel/compiler: Prepare SIMD selection helpers to handle different prog_datas

Browse source 
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 04:34:55 -08:00 • committed by Marge Bot
parent 6ffa597bcf
commit 6c194ddd18
2 changed files with 46 additions and 25 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -27,6 +27,8 @@
#include "brw_compiler.h"
#include <variant>
unsigned brw_required_dispatch_width(const struct shader_info *info);
static constexpr int SIMD_COUNT = 3;
@ -35,7 +37,7 @@ struct brw_simd_selection_state {
void *mem_ctx;
const struct intel_device_info *devinfo;
struct brw_cs_prog_data *prog_data;
std::variant<struct brw_cs_prog_data *> prog_data;
unsigned required_width;

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

@ -46,20 +46,33 @@ test_bit(unsigned mask, unsigned bit) {
return mask & (1u << bit);
}
namespace {
struct brw_cs_prog_data *
get_cs_prog_data(brw_simd_selection_state &state)
{
if (std::holds_alternative<struct brw_cs_prog_data *>(state.prog_data))
return std::get<struct brw_cs_prog_data *>(state.prog_data);
else
return nullptr;
}
}
bool
brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
{
assert(simd < SIMD_COUNT);
assert(!state.compiled[simd]);
struct brw_cs_prog_data *prog_data = state.prog_data;
const auto cs_prog_data = get_cs_prog_data(state);
const unsigned width = 8u << simd;
/* For shaders with variable size workgroup, in most cases we can compile
* all the variants (exceptions are bindless dispatch & ray queries), since
* the choice will happen only at dispatch time.
*/
const bool workgroup_size_variable = prog_data->local_size[0] == 0;
const bool workgroup_size_variable = cs_prog_data && cs_prog_data->local_size[0] == 0;
if (!workgroup_size_variable) {
if (state.spilled[simd]) {
@ -68,12 +81,6 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
return false;
}
const unsigned workgroup_size = prog_data->local_size[0] *
prog_data->local_size[1] *
prog_data->local_size[2];
unsigned max_threads = state.devinfo->max_cs_workgroup_threads;
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",
@ -81,6 +88,13 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
return false;
}
if (cs_prog_data) {
const unsigned workgroup_size = cs_prog_data->local_size[0] *
cs_prog_data->local_size[1] *
cs_prog_data->local_size[2];
unsigned max_threads = state.devinfo->max_cs_workgroup_threads;
if (simd > 0 && state.compiled[simd - 1] &&
workgroup_size <= (width / 2)) {
state.error[simd] = ralloc_asprintf(
@ -95,6 +109,7 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
width, workgroup_size, max_threads);
return false;
}
}
/* The SIMD32 is only enabled for cases it is needed unless forced.
*
@ -109,14 +124,14 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
}
}
if (width == 32 && prog_data->base.ray_queries > 0) {
if (width == 32 && cs_prog_data && cs_prog_data->base.ray_queries > 0) {
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) {
if (width == 32 && cs_prog_data && cs_prog_data->uses_btd_stack_ids) {
state.error[simd] = ralloc_asprintf(
state.mem_ctx, "SIMD%u skipped because of bindless shader calls",
width);
@ -147,14 +162,18 @@ brw_simd_mark_compiled(brw_simd_selection_state &state, unsigned simd, bool spil
assert(simd < SIMD_COUNT);
assert(!state.compiled[simd]);
auto cs_prog_data = get_cs_prog_data(state);
state.compiled[simd] = true;
state.prog_data->prog_mask |= 1u << simd;
if (cs_prog_data)
cs_prog_data->prog_mask |= 1u << simd;
/* If a SIMD spilled, all the larger ones would spill too. */
if (spilled) {
for (unsigned i = simd; i < SIMD_COUNT; i++) {
state.spilled[i] = true;
state.prog_data->prog_spilled |= 1u << i;
if (cs_prog_data)
cs_prog_data->prog_spilled |= 1u << i;
}
}
}