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

intel/compiler: Don't allocate memory for SIMD select error handling

Browse source 
The position in the error array already indicate the SIMD in question,
so take off all the formatted printing from the errors -- which in some
cases were just not needed.  We lose a little bit of extra context but
it is all easily derivable from the message and the SIMD.

This also will remove the overhead when SIMD selection is being used to
just to find the selected dispatch width -- at a point where the shaders
were already compiled -- and the errors are not used at all.

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9849
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25336>
This commit is contained in:
Caio Oliveira 2023-09-21 13:35:42 -07:00 • committed by Marge Bot
parent e55aa87f32
commit 1cdc4be14b
5 changed files with 17 additions and 38 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -8202,7 +8202,6 @@ brw_compile_cs(const struct brw_compiler *compiler,
}
brw_simd_selection_state simd_state{
.mem_ctx = params->base.mem_ctx,
.devinfo = compiler->devinfo,
.prog_data = prog_data,
.required_width = brw_required_dispatch_width(&nir->info),
@ -8260,7 +8259,8 @@ brw_compile_cs(const struct brw_compiler *compiler,
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: %s, %s and %s.\n",
"Can't compile shader: "
"SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n",
simd_state.error[0], simd_state.error[1],
simd_state.error[2]);
return NULL;
@ -8353,7 +8353,6 @@ compile_single_bs(const struct brw_compiler *compiler,
key->base.robust_flags);
brw_simd_selection_state simd_state{
.mem_ctx = params->base.mem_ctx,
.devinfo = compiler->devinfo,
.prog_data = prog_data,
@ -8396,7 +8395,8 @@ compile_single_bs(const struct brw_compiler *compiler,
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: %s and %s.",
"Can't compile shader: "
"SIMD8 '%s' and SIMD16 '%s'.\n",
simd_state.error[0], simd_state.error[1]);
return 0;
}

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

@ -298,7 +298,6 @@ brw_compile_task(const struct brw_compiler *compiler,
BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_DRAW_ID);
brw_simd_selection_state simd_state{
.mem_ctx = params->base.mem_ctx,
.devinfo = compiler->devinfo,
.prog_data = &prog_data->base,
.required_width = brw_required_dispatch_width(&nir->info),
@ -344,7 +343,8 @@ brw_compile_task(const struct brw_compiler *compiler,
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: %s, %s and %s.\n",
"Can't compile shader: "
"SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n",
simd_state.error[0], simd_state.error[1],
simd_state.error[2]);
return NULL;
@ -1513,7 +1513,6 @@ brw_compile_mesh(const struct brw_compiler *compiler,
brw_nir_lower_mue_outputs(nir, &prog_data->map);
brw_simd_selection_state simd_state{
.mem_ctx = params->base.mem_ctx,
.devinfo = compiler->devinfo,
.prog_data = &prog_data->base,
.required_width = brw_required_dispatch_width(&nir->info),
@ -1571,9 +1570,10 @@ brw_compile_mesh(const struct brw_compiler *compiler,
if (selected_simd < 0) {
params->base.error_str =
ralloc_asprintf(params->base.mem_ctx,
"Can't compile shader: %s, %s and %s.\n",
"Can't compile shader: "
"SIMD8 '%s', SIMD16 '%s' and SIMD32 '%s'.\n",
simd_state.error[0], simd_state.error[1],
simd_state.error[2]);;
simd_state.error[2]);
return NULL;
}

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

@ -34,7 +34,6 @@ unsigned brw_required_dispatch_width(const struct shader_info *info);
static constexpr int SIMD_COUNT = 3;
struct brw_simd_selection_state {
void *mem_ctx;
const struct intel_device_info *devinfo;
std::variant<struct brw_cs_prog_data *,

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

@ -88,15 +88,12 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
if (!workgroup_size_variable) {
if (state.spilled[simd]) {
state.error[simd] = ralloc_asprintf(
state.mem_ctx, "SIMD%u skipped because would spill", width);
state.error[simd] = "Would spill";
return false;
}
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);
state.error[simd] = "Different than required dispatch width";
return false;
}
@ -109,16 +106,12 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
if (simd > 0 && state.compiled[simd - 1] &&
workgroup_size <= (width / 2)) {
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);
state.error[simd] = "Workgroup size already fits in smaller SIMD";
return false;
}
if (DIV_ROUND_UP(workgroup_size, width) > max_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);
state.error[simd] = "Would need more than max_threads to fit all invocations";
return false;
}
}
@ -129,24 +122,19 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
*/
if (width == 32) {
if (!INTEL_DEBUG(DEBUG_DO32) && (state.compiled[0] || state.compiled[1])) {
state.error[simd] = ralloc_strdup(
state.mem_ctx, "SIMD32 skipped because not required");
state.error[simd] = "SIMD32 not required (use INTEL_DEBUG=do32 to force)";
return false;
}
}
}
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);
state.error[simd] = "Ray queries not supported";
return false;
}
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);
state.error[simd] = "Bindless shader calls not supported";
return false;
}
@ -182,9 +170,7 @@ brw_simd_should_compile(brw_simd_selection_state &state, unsigned simd)
static_assert(ARRAY_SIZE(env_skip) == SIMD_COUNT);
if (unlikely(env_skip[simd])) {
state.error[simd] = ralloc_asprintf(
state.mem_ctx, "SIMD%u skipped because INTEL_DEBUG=no%u",
width, width);
state.error[simd] = "Disabled by INTEL_DEBUG environment variable";
return false;
}
@ -257,10 +243,7 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
cloned.prog_mask = 0;
cloned.prog_spilled = 0;
void *mem_ctx = ralloc_context(NULL);
brw_simd_selection_state simd_state{
.mem_ctx = mem_ctx,
.devinfo = devinfo,
.prog_data = &cloned,
};
@ -275,7 +258,5 @@ brw_simd_select_for_workgroup_size(const struct intel_device_info *devinfo,
}
}
ralloc_free(mem_ctx);
return brw_simd_select(simd_state);
}

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

@ -46,7 +46,6 @@ protected:
, 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,
}