From 67fc7a1763c4878315da7456206f70477291edde Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Wed, 5 Jan 2022 13:51:50 +0000 Subject: [PATCH] nir/uniform_atomics: fix is_atomic_already_optimized without workgroups MIME-Version: 1.0 Content-Type: text/plain; charset=utf8 Content-Transfer-Encoding: 8bit dims_needed would have been zero, so this would always returned true for non-compute stages. Also fix this for variable workgroup sizes. Improves Shadow of the Tomb Raider RX 6800 performance by 10.6%, 11.5% and 4.5% (day_of_dead, jungle and paititi scenes). radv_perf before and after: {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'day_of_dead', 'avg_fps': '62.913333333333334', 'min_fps': '62.81', 'max_fps': '62.98', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'jungle', 'avg_fps': '64.02666666666666', 'min_fps': '63.93', 'max_fps': '64.11', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'paititi', 'avg_fps': '74.81666666666666', 'min_fps': '74.72', 'max_fps': '74.88', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'day_of_dead', 'avg_fps': '69.57', 'min_fps': '69.52', 'max_fps': '69.63', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'jungle', 'avg_fps': '71.41000000000001', 'min_fps': '71.31', 'max_fps': '71.5', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'paititi', 'avg_fps': '78.16666666666667', 'min_fps': '78.07', 'max_fps': '78.23', 'interations': '3'} Performance now seems slightly better than AMDVLK 2021.Q4.3: {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'day_of_dead', 'avg_fps': '68.02666666666666', 'min_fps': '67.95', 'max_fps': '68.16', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'jungle', 'avg_fps': '70.24666666666667', 'min_fps': '69.83', 'max_fps': '70.51', 'interations': '3'} {'app': 'SotTR', 'resolution': '3840x2160', 'preset': 'VeryHigh', 'antialiasing': 'off', 'scene': 'paititi', 'avg_fps': '77.19', 'min_fps': '77.18', 'max_fps': '77.2', 'interations': '3'} fossil-db (Sienna Cichlid): Totals from 40 (0.03% of 134621) affected shaders: CodeSize: 62676 -> 65996 (+5.30%) Instrs: 11372 -> 12111 (+6.50%) Latency: 144122 -> 142848 (-0.88%); split: -1.09%, +0.21% InvThroughput: 19686 -> 19847 (+0.82%); split: -0.06%, +0.87% VClause: 304 -> 306 (+0.66%) SClause: 603 -> 604 (+0.17%); split: -0.83%, +1.00% Copies: 780 -> 858 (+10.00%) Branches: 235 -> 329 (+40.00%) PreSGPRs: 1072 -> 1083 (+1.03%); split: -0.37%, +1.40% Signed-off-by: Rhys Perry Reviewed-by: Bas Nieuwenhuizen Reviewed-by: Daniel Schürmann Part-of: --- src/compiler/nir/nir_opt_uniform_atomics.c | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/src/compiler/nir/nir_opt_uniform_atomics.c b/src/compiler/nir/nir_opt_uniform_atomics.c index cdd1319..6746a94 100644 --- a/src/compiler/nir/nir_opt_uniform_atomics.c +++ b/src/compiler/nir/nir_opt_uniform_atomics.c @@ -167,11 +167,16 @@ is_atomic_already_optimized(nir_shader *shader, nir_intrinsic_instr *instr) } } - unsigned dims_needed = 0; - for (unsigned i = 0; i < 3; i++) - dims_needed |= (shader->info.workgroup_size[i] > 1) << i; + if (gl_shader_stage_uses_workgroup(shader->info.stage)) { + unsigned dims_needed = 0; + for (unsigned i = 0; i < 3; i++) + dims_needed |= (shader->info.workgroup_size_variable || + shader->info.workgroup_size[i] > 1) << i; + if ((dims & dims_needed) == dims_needed) + return true; + } - return (dims & dims_needed) == dims_needed || dims & 0x8; + return dims & 0x8; } /* Perform a reduction and/or exclusive scan. */ -- 2.7.4