From 719b48f85d6a5f5d9129e99d2e9a0783c03ecec4 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Thu, 11 Nov 2021 14:26:00 +0000 Subject: [PATCH] nir/lower_system_values: replace local_invocation_id components with zero fossil-db (Sienna Cichlid): Totals from 360 (0.28% of 128647) affected shaders: VGPRs: 7912 -> 7272 (-8.09%); split: -8.59%, +0.51% CodeSize: 542456 -> 544688 (+0.41%); split: -0.32%, +0.73% MaxWaves: 10866 -> 10952 (+0.79%) Instrs: 95973 -> 96010 (+0.04%); split: -0.34%, +0.38% Latency: 4366023 -> 4344664 (-0.49%); split: -0.90%, +0.41% InvThroughput: 19656659 -> 18297185 (-6.92%); split: -6.92%, +0.00% VClause: 3242 -> 3116 (-3.89%); split: -4.04%, +0.15% SClause: 3422 -> 3504 (+2.40%); split: -0.20%, +2.60% Copies: 8854 -> 9376 (+5.90%); split: -0.89%, +6.79% Branches: 2329 -> 2326 (-0.13%); split: -0.39%, +0.26% PreSGPRs: 7620 -> 7841 (+2.90%); split: -0.43%, +3.33% PreVGPRs: 5765 -> 5504 (-4.53%) Signed-off-by: Rhys Perry Reviewed-by: Daniel-schuermann Part-of: --- src/compiler/nir/nir_lower_system_values.c | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index 4d270578da6..5a60b42fbc0 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -387,6 +387,21 @@ lower_compute_system_value_instr(nir_builder *b, return nir_vec3(b, x, y, z); } + + /* If a workgroup size dimension is 1, then the local invocation id must be zero. */ + nir_component_mask_t is_zero = 0; + is_zero |= b->shader->info.workgroup_size[0] == 1 ? 0x1 : 0x0; + is_zero |= b->shader->info.workgroup_size[1] == 1 ? 0x2 : 0x0; + is_zero |= b->shader->info.workgroup_size[2] == 1 ? 0x4 : 0x0; + if (!b->shader->info.workgroup_size_variable && is_zero) { + nir_ssa_def *defs[3]; + for (unsigned i = 0; i < 3; i++) { + defs[i] = is_zero & (1 << i) ? nir_imm_zero(b, 1, 32) : + nir_channel(b, &intrin->dest.ssa, i); + } + return nir_vec(b, defs, 3); + } + return NULL; case nir_intrinsic_load_local_invocation_index: