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

nir/lower_system_values: Add shortcut for 1D workgroups.

Browse source 
When the workgroup is 1 dimensional, simply use	a vec3
filled with zeroes and the local invocation index.
This is is better than lower_id_to_index + constant folding,
because this way we don't leave behind extra ALU instrs.

Note, this is relevant to mesh shaders on RDNA2 because
it enables us to better detect cross-invocation output
access.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marcin Ślusarz <marcin.slusarz@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18464>
This commit is contained in:
Timur Kristóf 2022-09-07 15:23:46 +02:00 • committed by Marge Bot
parent 9fdfe43240
commit c80d811403
1 changed files with 24 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -353,6 +353,30 @@ lower_compute_system_value_instr(nir_builder *b,
if (b->shader->options->lower_cs_local_id_to_index ||
(options && options->lower_cs_local_id_to_index)) {
nir_ssa_def *local_index = nir_load_local_invocation_index(b);
if (!b->shader->info.workgroup_size_variable) {
/* Shortcut for 1 dimensional workgroups:
* Use local_invocation_index directly, which is better than
* lower_id_to_index + constant folding, because
* this way we don't leave behind extra ALU instrs.
*/
/* size_x = 1, size_y = 1, therefore Z = local index */
if (b->shader->info.workgroup_size[0] == 1 &&
b->shader->info.workgroup_size[1] == 1)
return nir_vec3(b, nir_imm_int(b, 0), nir_imm_int(b, 0), local_index);
/* size_x = 1, size_z = 1, therefore Y = local index */
if (b->shader->info.workgroup_size[0] == 1 &&
b->shader->info.workgroup_size[2] == 1)
return nir_vec3(b, nir_imm_int(b, 0), local_index, nir_imm_int(b, 0));
/* size_y = 1, size_z = 1, therefore X = local index */
if (b->shader->info.workgroup_size[1] == 1 &&
b->shader->info.workgroup_size[2] == 1)
return nir_vec3(b, local_index, nir_imm_int(b, 0), nir_imm_int(b, 0));
}
nir_ssa_def *local_size = nir_load_workgroup_size(b);
return lower_id_to_index(b, local_index, local_size, bit_size);
}