From eda8fc1e909ace952c7db50239c5ccf4a25ab74d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20Sch=C3=BCrmann?= Date: Wed, 16 Jul 2025 16:45:14 +0200 Subject: [PATCH] radv: delay lowering global access Totals from 21 (0.03% of 79839) affected shaders: (Navi48) Instrs: 30258 -> 30249 (-0.03%); split: -0.05%, +0.02% CodeSize: 159660 -> 159552 (-0.07%); split: -0.07%, +0.01% Latency: 188154 -> 188131 (-0.01%); split: -0.02%, +0.00% SClause: 251 -> 252 (+0.40%) SMEM: 619 -> 598 (-3.39%) Part-of: --- src/amd/vulkan/radv_pipeline.c | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c index 0ceadc88617..652852940d1 100644 --- a/src/amd/vulkan/radv_pipeline.c +++ b/src/amd/vulkan/radv_pipeline.c @@ -558,7 +558,6 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_graphics_stat .allow_fp16 = gfx_level >= GFX9, }); - NIR_PASS(_, stage->nir, ac_nir_lower_global_access); NIR_PASS(_, stage->nir, ac_nir_lower_intrinsics_to_args, gfx_level, pdev->info.has_ls_vgpr_init_bug && gfx_state && !gfx_state->vs.has_prolog, radv_select_hw_stage(&stage->info, gfx_level), stage->info.wave_size, stage->info.workgroup_size, @@ -588,6 +587,8 @@ radv_postprocess_nir(struct radv_device *device, const struct radv_graphics_stat NIR_PASS(_, stage->nir, ac_nir_lower_mem_access_bit_sizes, gfx_level, use_llvm); } + NIR_PASS(_, stage->nir, ac_nir_lower_global_access); + radv_optimize_nir_algebraic( stage->nir, io_to_mem || lowered_ngg || stage->stage == MESA_SHADER_COMPUTE || stage->stage == MESA_SHADER_TASK, gfx_level >= GFX8, gfx_level);