From c80d81140329456d43ec19c82bfd5e57253f440c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timur=20Krist=C3=B3f?= Date: Wed, 7 Sep 2022 15:23:46 +0200 Subject: [PATCH] nir/lower_system_values: Add shortcut for 1D workgroups. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 Reviewed-by: Alyssa Rosenzweig Reviewed-by: Marcin Ślusarz Part-of: --- 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); }