From 2d07bcad661c5a4e9f7df1fb9eb630deb1fa3fe4 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Thu, 11 Nov 2021 14:26:22 +0000 Subject: [PATCH] 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 Reviewed-by: Daniel-schuermann Part-of: --- 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;