fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-11 01:30:25 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radv: lower load_local_invocation_index with 1D workgroups

Browse source 
For 1D workgroups, we can just load from an input VGPR.

fossil-db (Sienna Cichlid):
Totals from 226 (0.18% of 128647) affected shaders:
CodeSize: 1200476 -> 1195696 (-0.40%); split: -0.49%, +0.09%
Instrs: 223817 -> 223328 (-0.22%); split: -0.29%, +0.07%
Latency: 2552394 -> 2549606 (-0.11%); split: -0.15%, +0.04%
InvThroughput: 533989 -> 532670 (-0.25%); split: -0.27%, +0.02%
VClause: 5191 -> 5188 (-0.06%)
SClause: 7637 -> 7636 (-0.01%)
Copies: 18165 -> 18182 (+0.09%); split: -0.22%, +0.31%
Branches: 10446 -> 10442 (-0.04%)
PreSGPRs: 8049 -> 8041 (-0.10%); split: -0.17%, +0.07%
PreVGPRs: 7785 -> 7767 (-0.23%); split: -0.32%, +0.09%

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel-schuermann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13757>
This commit is contained in:
Rhys Perry 2021-11-11 14:26:22 +00:00 • committed by Marge Bot
parent 719b48f85d
commit 2d07bcad66
1 changed files with 6 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

7
src/amd/vulkan/radv_shader.c
View file

@ -615,7 +615,12 @@ radv_shader_compile_to_nir(struct radv_device *device, struct vk_shader_module *
}
NIR_PASS_V(nir, nir_lower_system_values);
NIR_PASS_V(nir, nir_lower_compute_system_values, NULL);
nir_lower_compute_system_values_options csv_options = {
.lower_local_invocation_index = ((nir->info.workgroup_size[0] == 1) +
(nir->info.workgroup_size[1] == 1) +
(nir->info.workgroup_size[2] == 1)) == 2,
};
NIR_PASS_V(nir, nir_lower_compute_system_values, &csv_options);
/* Vulkan uses the separate-shader linking model */
nir->info.separate_shader = true;