Module: Mesa Branch: main Commit: c80d81140329456d43ec19c82bfd5e57253f440c URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=c80d81140329456d43ec19c82bfd5e57253f440c
Author: Timur Kristóf <[email protected]> Date: Wed Sep 7 15:23:46 2022 +0200 nir/lower_system_values: Add shortcut for 1D workgroups. 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 <[email protected]> Reviewed-by: Alyssa Rosenzweig <[email protected]> Reviewed-by: Marcin Ślusarz <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18464> --- src/compiler/nir/nir_lower_system_values.c | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index e4698965edc..aed2f73e25d 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -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); }
