From 548b383310da7e5e29e7c24525a816a45afc26b8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timur=20Krist=C3=B3f?= Date: Thu, 26 Aug 2021 13:32:51 +0200 Subject: [PATCH] nir: Fix local_invocation_index upper bound for non-compute-like stages. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 Reviewed-by: Rhys Perry Reviewed-by: Jason Ekstrand Part-of: --- 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] *