Mesa (main): radv: lower load_local_invocation_index with 1D workgroups

[GitLab Mirror]

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

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

radv: lower load_local_invocation_index with 1D workgroups

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

---

 src/amd/vulkan/radv_shader.c | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/src/amd/vulkan/radv_shader.c b/src/amd/vulkan/radv_shader.c
index 2e903ba4862..f4ea6265c47 100644
--- a/src/amd/vulkan/radv_shader.c
+++ b/src/amd/vulkan/radv_shader.c
@@ -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;