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

nir/lower_system_values: optimize global ID

Browse source 
for drivers where we need to lower a base_workgroup_id but not global IDs.
rather than lowering the whole global ID to stick the base workgroup ID in
there, just add the workgroup offset to the final thread position.

Elden ring fossils:

Totals from 52 (1.62% of 3206) affected shaders:
Instrs: 48355 -> 48233 (-0.25%); split: -0.31%, +0.06%
CodeSize: 331912 -> 331148 (-0.23%); split: -0.28%, +0.05%
ALU: 30853 -> 30674 (-0.58%); split: -0.70%, +0.12%
FSCIB: 30853 -> 30674 (-0.58%); split: -0.70%, +0.12%
IC: 9054 -> 8958 (-1.06%)
GPRs: 4184 -> 4216 (+0.76%)
Uniforms: 6703 -> 6677 (-0.39%); split: -1.61%, +1.22%

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35909>
This commit is contained in:
Alyssa Rosenzweig 2025-06-23 16:02:16 -04:00 • committed by Marge Bot
parent 96d124092f
commit b992703477
1 changed files with 11 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

13
src/compiler/nir/nir_lower_system_values.c
View file

@ -686,8 +686,7 @@ lower_compute_system_value_instr(nir_builder *b,
}
case nir_intrinsic_load_global_invocation_id: {
if ((options && options->has_base_workgroup_id) ||
!b->shader->options->has_cs_global_id) {
if (!b->shader->options->has_cs_global_id) {
nir_def *group_size = nir_load_workgroup_size(b);
nir_def *group_id = nir_load_workgroup_id(b);
nir_def *base_group_id = nir_load_base_workgroup_id(b, bit_size);
@ -695,6 +694,16 @@ lower_compute_system_value_instr(nir_builder *b,
return nir_iadd(b, nir_imul(b, nir_iadd(b, nir_u2uN(b, group_id, bit_size), base_group_id), nir_u2uN(b, group_size, bit_size)),
nir_u2uN(b, local_id, bit_size));
} else if (options && options->has_base_workgroup_id &&
_mesa_set_search(state->lower_once_list, instr) == NULL) {
nir_def *global_id = nir_load_global_invocation_id(b, bit_size);
nir_def *group_size = nir_u2uN(b, nir_load_workgroup_size(b), bit_size);
nir_def *base_group_id = nir_load_base_workgroup_id(b, bit_size);
_mesa_set_add(state->lower_once_list, global_id->parent_instr);
return nir_iadd(b, global_id, nir_imul(b, base_group_id, group_size));
} else if (options && options->global_id_is_32bit && bit_size > 32) {
return nir_u2uN(b, nir_load_global_invocation_id(b, 32), bit_size);
} else {