KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
151,762 Commits 1 Branch 0 Tags 356 MiB
Commit Graph

13 Commits

Author SHA1 Message Date
Rhys Perry 67fc7a1763 nir/uniform_atomics: fix is_atomic_already_optimized without workgroups 
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 <pendingchaos02@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14407>
 2022-01-10 19:57:38 +00:00
Filip Gawin f32dcb6fe1 nir: assert that variables in optimize_atomic are initialized 
If you gonna view context of function parse_atomic_op,
then you gonna know that index for array (data_src)
can be unitialized. Imho this approach is cleaner
than doing stuff inside parse_atomic_op.

Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12995>
 2021-11-08 15:10:07 +00:00
Caio Marcelo de Oliveira Filho 8af6766062 nir: Move workgroup_size and workgroup_variable_size into common shader_info 
Move it out the "cs" sub-struct, since these will be used for other
shader stages in the future.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225>
 2021-06-08 09:23:55 -07:00
Caio Marcelo de Oliveira Filho 430d2206da compiler: Rename local_size to workgroup_size 
Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
 2021-06-07 22:34:42 +00:00
Jason Ekstrand 117668b811 nir: Make nir_ssa_def_rewrite_uses take an SSA value 
This commit replaces the new_src parameter of nir_ssa_def_rewrite_uses()
with an SSA def, removes nir_ssa_def_rewrite_uses_ssa(), and rewrites
all the users as needed.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Alyssa Rosenzweig <alyssa@collabora.com>
Acked-By: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9383>
 2021-03-08 16:59:55 +00:00
Rhys Perry ed9c3c4f19 nir: add nir_ssa_def_is_unused() 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8784>
 2021-03-01 17:38:10 +00:00
Rhys Perry f8072c133d nir/opt_uniform_atomics: fix elect detection 
fossil-db (GFX10.3):
Totals from 30 (0.02% of 139391) affected shaders:
SGPRs: 1736 -> 1712 (-1.38%)
CodeSize: 262116 -> 254728 (-2.82%)
Instrs: 50341 -> 48857 (-2.95%)
Cycles: 486384 -> 477556 (-1.82%)
VMEM: 4821 -> 4589 (-4.81%)
Copies: 5013 -> 4890 (-2.45%)
Branches: 2108 -> 1983 (-5.93%)
PreSGPRs: 1444 -> 1418 (-1.80%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8654>
 2021-01-25 21:04:52 +00:00
Rhys Perry eb70c52abe nir/opt_uniform_atomics: recognize more complicated invocation comparisons 
For example, gl_LocalInvocationID.x + gl_LocalInvocationID.y * 8.

fossil-db (GFX10.3):
Totals from 8 (0.01% of 139391) affected shaders:
CodeSize: 15224 -> 14800 (-2.79%)
Instrs: 2880 -> 2798 (-2.85%)
Cycles: 44556 -> 44204 (-0.79%)
VMEM: 407 -> 473 (+16.22%); split: +17.69%, -1.47%
Copies: 491 -> 483 (-1.63%)
Branches: 200 -> 192 (-4.00%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8654>
 2021-01-25 21:04:52 +00:00
Rhys Perry 9b040737d3 nir: use intrinsic builders 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6587>
 2020-11-26 17:50:38 +00:00
Rhys Perry 21422b1ff2 nir/opt_uniform_atomics: remove useless returns 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7117>
 2020-10-14 09:53:34 +00:00
Rhys Perry a7114f3f46 nir/opt_uniform_atomics: don't optimize atomics twice 
Applications sometimes already do this optimization themselves.

fossil-db (Navi):
Totals from 51 (0.04% of 135946) affected shaders:
CodeSize: 507484 -> 501860 (-1.11%)
Instrs: 99635 -> 98471 (-1.17%)
Cycles: 2421944 -> 2414780 (-0.30%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6558>
 2020-10-13 12:47:21 +00:00
Rhys Perry bc43650522 nir/opt_uniform_atomics: optimize image atomics 
fossil-db (Navi):
Totals from 65 (0.05% of 135946) affected shaders:
SGPRs: 3792 -> 3784 (-0.21%)
VGPRs: 2784 -> 2716 (-2.44%)
CodeSize: 707492 -> 713080 (+0.79%)
MaxWaves: 873 -> 887 (+1.60%)
Instrs: 133376 -> 134524 (+0.86%)
Cycles: 3004772 -> 3011440 (+0.22%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6558>
 2020-10-13 12:47:21 +00:00
Rhys Perry f83bc5beb8 nir: add pass to optimize uniform atomics 
This optimizes atomics with a uniform offset so that only one atomic
operation is done in the subgroup.

For shaders which do a very large amount of atomics, this can
significantly improve performance.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6558>
 2020-10-13 12:47:21 +00:00