From d04e1ea02dc8165dcf36d27850ea59b7ed5e48d8 Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Tue, 28 Jan 2025 14:35:20 +0000 Subject: [PATCH] radv: move nir_opt_shrink_vectors later This seems to be helpful with shaders which use NGG culling. fossil-db (navi21): Totals from 3529 (4.45% of 79377) affected shaders: MaxWaves: 81490 -> 82066 (+0.71%) Instrs: 2868872 -> 2863476 (-0.19%); split: -0.22%, +0.04% CodeSize: 14949540 -> 14927580 (-0.15%); split: -0.18%, +0.03% VGPRs: 165440 -> 164144 (-0.78%) SpillSGPRs: 578 -> 405 (-29.93%) Latency: 15388119 -> 15151882 (-1.54%); split: -1.74%, +0.20% InvThroughput: 2935873 -> 2929736 (-0.21%); split: -0.25%, +0.04% VClause: 70192 -> 68904 (-1.83%); split: -2.17%, +0.33% SClause: 67678 -> 67679 (+0.00%); split: -0.10%, +0.10% Copies: 265824 -> 261458 (-1.64%); split: -1.96%, +0.32% Branches: 75084 -> 75088 (+0.01%); split: -0.02%, +0.02% PreSGPRs: 165962 -> 165716 (-0.15%) PreVGPRs: 135122 -> 134724 (-0.29%) VALU: 1681747 -> 1677134 (-0.27%); split: -0.32%, +0.05% SALU: 436003 -> 435915 (-0.02%); split: -0.03%, +0.01% Signed-off-by: Rhys Perry Reviewed-by: Georg Lehmann Part-of: --- src/amd/vulkan/radv_pipeline.c | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c index 2812e4d8fa1..2afefc56ee8 100644 --- a/src/amd/vulkan/radv_pipeline.c +++ b/src/amd/vulkan/radv_pipeline.c @@ -406,10 +406,6 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_graphics_stat /* TODO: vectorize loads after this to vectorize loading adjacent descriptors */ NIR_PASS_V(stage->nir, radv_nir_apply_pipeline_layout, device, stage); - if (!stage->key.optimisations_disabled) { - NIR_PASS(_, stage->nir, nir_opt_shrink_vectors, true); - } - NIR_PASS(_, stage->nir, nir_lower_alu_width, opt_vectorize_callback, device); nir_move_options sink_opts = nir_move_const_undef | nir_move_copies; @@ -517,6 +513,11 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_graphics_stat radv_select_hw_stage(&stage->info, gfx_level), stage->info.wave_size, stage->info.workgroup_size, &stage->args.ac); NIR_PASS_V(stage->nir, radv_nir_lower_abi, gfx_level, stage, gfx_state, pdev->info.address32_hi); + + if (!stage->key.optimisations_disabled) { + NIR_PASS(_, stage->nir, nir_opt_shrink_vectors, true); + } + radv_optimize_nir_algebraic( stage->nir, io_to_mem || lowered_ngg || stage->stage == MESA_SHADER_COMPUTE || stage->stage == MESA_SHADER_TASK, gfx_level >= GFX8);