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
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
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
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
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
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
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
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
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
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
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
Caio Marcelo de Oliveira Filho
Jason Ekstrand
Rhys Perry