Mesa (main): nir/lower_system_values: replace local_invocation_id components with zero

[GitLab Mirror]

Module: Mesa
Branch: main
Commit: 719b48f85d6a5f5d9129e99d2e9a0783c03ecec4
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=719b48f85d6a5f5d9129e99d2e9a0783c03ecec4

Author: Rhys Perry <pendingchaos02 at gmail.com>
Date:   Thu Nov 11 14:26:00 2021 +0000

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 <pendingchaos02 at gmail.com>
Reviewed-by: Daniel-schuermann <daniel at schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13757>

---

 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: