Mesa (main): nir: Fix local_invocation_index upper bound for non-compute-like stages.

[GitLab Mirror]

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

Author: Timur Kristóf <timur.kristof at gmail.com>
Date:   Thu Aug 26 13:32:51 2021 +0200

nir: Fix local_invocation_index upper bound for non-compute-like stages.

The lowered LS and NGG stages use local_invocation_index and they
can benefit from the unsigned upper bound because they can emit a
less expensive integer multiplication instruction.
This was working in the past, but accidentally borked by a refactor.

Fossil DB changes on Sienna Cichlid:

Totals from 956 (0.74% of 128647) affected shaders:
CodeSize: 2354172 -> 2344712 (-0.40%)
Instrs: 434359 -> 434327 (-0.01%)
Latency: 1883949 -> 1876814 (-0.38%)
InvThroughput: 762638 -> 757405 (-0.69%)

Fossil DB changes on Sienna Cichlid (with NGGC enabled):

Totals from 57873 (44.99% of 128647) affected shaders:
CodeSize: 155844192 -> 155607064 (-0.15%)
Instrs: 29799184 -> 29799152 (-0.00%)
Latency: 130959764 -> 130814224 (-0.11%); split: -0.11%, +0.00%
InvThroughput: 21100300 -> 20928635 (-0.81%); split: -0.81%, +0.00%

Fixes: 8af6766062044167fb3b61950ddbc7d67e4c3e48
Signed-off-by: Timur Kristóf <timur.kristof at gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02 at gmail.com>
Reviewed-by: Jason Ekstrand <jason at jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12558>

---

 src/compiler/nir/nir_range_analysis.c | 10 +++++++++-
 1 file changed, 9 insertions(+), 1 deletion(-)

diff --git a/src/compiler/nir/nir_range_analysis.c b/src/compiler/nir/nir_range_analysis.c
index f45f6db2fd7..a50e43651bb 100644
--- a/src/compiler/nir/nir_range_analysis.c
+++ b/src/compiler/nir/nir_range_analysis.c
@@ -1292,7 +1292,15 @@ nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
       nir_intrinsic_instr *intrin = nir_instr_as_intrinsic(scalar.def->parent_instr);
       switch (intrin->intrinsic) {
       case nir_intrinsic_load_local_invocation_index:
-         if (shader->info.workgroup_size_variable) {
+         /* The local invocation index is used under the hood by RADV for
+          * some non-compute-like shaders (eg. LS and NGG). These technically
+          * run in workgroups on the HW, even though this fact is not exposed
+          * by the API.
+          * They can safely use the same code path here as variable sized
+          * compute-like shader stages.
+          */
+         if (!gl_shader_stage_uses_workgroup(shader->info.stage) ||
+             shader->info.workgroup_size_variable) {
             res = config->max_workgroup_invocations - 1;
          } else {
             res = (shader->info.workgroup_size[0] *