OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
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: 8af6766062
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12558>
This commit is contained in:
Timur Kristóf
2021-08-26 13:32:51 +02:00
committed by Marge Bot
parent a25fd1787a
commit 548b383310
1 changed files with 9 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
src/compiler/nir/nir_range_analysis.c
View File

@@ -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] *