Konstantin Seurer
de32f9275f
treewide: add & use parent instr helpers
...
We add a bunch of new helpers to avoid the need to touch >parent_instr,
including the full set of:
* nir_def_is_*
* nir_def_as_*_or_null
* nir_def_as_* [assumes the right instr type]
* nir_src_is_*
* nir_src_as_*
* nir_scalar_is_*
* nir_scalar_as_*
Plus nir_def_instr() where there's no more suitable helper.
Also an existing helper is renamed to unify all the names, while we're
churning the tree:
* nir_src_as_alu_instr -> nir_src_as_alu
..and then we port the tree to use the helpers as much as possible, using
nir_def_instr() where that does not work.
Acked-by: Marek Olšák <maraeo@gmail.com>
---
To eliminate nir_def::parent_instr we need to churn the tree anyway, so I'm
taking this opportunity to clean up a lot of NIR patterns.
Co-authored-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38313 >
2025-11-12 21:22:13 +00:00
Georg Lehmann
142079759b
nir/opt_uniform_atomics: optimize xchg with uniform address and data
...
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37822 >
2025-10-14 06:24:16 +00:00
Georg Lehmann
b0d3db3733
nir: add atomic isub
...
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37702 >
2025-10-07 14:07:56 +00:00
Qiang Yu
4847e0b380
all: rename gl_shader_stage_uses_workgroup to mesa_shader_stage_uses_workgroup
...
Acked-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Yonggang Luo <luoyonggang@gmail.com>
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36569 >
2025-08-06 10:28:41 +08:00
Alyssa Rosenzweig
cc6e3b84cb
treewide: use nir_def_as_*
...
Via Coccinelle patch:
@@
expression definition;
@@
-nir_instr_as_alu(definition->parent_instr)
+nir_def_as_alu(definition)
@@
expression definition;
@@
-nir_instr_as_intrinsic(definition->parent_instr)
+nir_def_as_intrinsic(definition)
@@
expression definition;
@@
-nir_instr_as_phi(definition->parent_instr)
+nir_def_as_phi(definition)
@@
expression definition;
@@
-nir_instr_as_load_const(definition->parent_instr)
+nir_def_as_load_const(definition)
@@
expression definition;
@@
-nir_instr_as_deref(definition->parent_instr)
+nir_def_as_deref(definition)
@@
expression definition;
@@
-nir_instr_as_tex(definition->parent_instr)
+nir_def_as_tex(definition)
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Reviewed-by: Marek Olšák <maraeo@gmail.com>
Acked-by: Karol Herbst <kherbst@redhat.com>
Acked-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36489 >
2025-08-01 15:34:24 +00:00
Alyssa Rosenzweig
9a58a8257e
treewide: Switch to nir_progress
...
Via the Coccinelle patch at the end of the commit message, followed by
sed -ie 's/progress = progress | /progress |=/g' $(git grep -l 'progress = prog')
ninja -C ~/mesa/build clang-format
cd ~/mesa/src/compiler/nir && clang-format -i *.c
agxfmt
@@
identifier prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
-return prog;
+return nir_progress(prog, impl, metadata);
@@
expression prog_expr, impl, metadata;
@@
-if (prog_expr) {
-nir_metadata_preserve(impl, metadata);
-return true;
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-return false;
-}
+bool progress = prog_expr;
+return nir_progress(progress, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-nir_metadata_preserve(impl, prog ? (metadata) : nir_metadata_all);
-return prog;
+return nir_progress(prog, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-nir_metadata_preserve(impl, prog ? (metadata) : nir_metadata_all);
+nir_progress(prog, impl, metadata);
@@
expression impl, metadata;
@@
-nir_metadata_preserve(impl, metadata);
-return true;
+return nir_progress(true, impl, metadata);
@@
expression impl;
@@
-nir_metadata_preserve(impl, nir_metadata_all);
-return false;
+return nir_no_progress(impl);
@@
identifier other_prog, prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
-other_prog |= prog;
+other_prog = other_prog | nir_progress(prog, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+nir_progress(prog, impl, metadata);
@@
identifier other_prog, prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-other_prog = true;
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+other_prog = other_prog | nir_progress(prog, impl, metadata);
@@
expression prog_expr, impl, metadata;
identifier prog;
@@
-if (prog_expr) {
-nir_metadata_preserve(impl, metadata);
-prog = true;
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+bool impl_progress = prog_expr;
+prog = prog | nir_progress(impl_progress, impl, metadata);
@@
identifier other_prog, prog;
expression impl, metadata;
@@
-if (prog) {
-other_prog = true;
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+other_prog = other_prog | nir_progress(prog, impl, metadata);
@@
expression prog_expr, impl, metadata;
identifier prog;
@@
-if (prog_expr) {
-prog = true;
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+bool impl_progress = prog_expr;
+prog = prog | nir_progress(impl_progress, impl, metadata);
@@
expression prog_expr, impl, metadata;
@@
-if (prog_expr) {
-nir_metadata_preserve(impl, metadata);
-} else {
-nir_metadata_preserve(impl, nir_metadata_all);
-}
+bool impl_progress = prog_expr;
+nir_progress(impl_progress, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-nir_metadata_preserve(impl, metadata);
-prog = true;
+prog = nir_progress(true, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-}
-return prog;
+return nir_progress(prog, impl, metadata);
@@
identifier prog;
expression impl, metadata;
@@
-if (prog) {
-nir_metadata_preserve(impl, metadata);
-}
+nir_progress(prog, impl, metadata);
@@
expression impl;
@@
-nir_metadata_preserve(impl, nir_metadata_all);
+nir_no_progress(impl);
@@
expression impl, metadata;
@@
-nir_metadata_preserve(impl, metadata);
+nir_progress(true, impl, metadata);
squashme! sed -ie 's/progress = progress | /progress |=/g' $(git grep -l 'progress = prog')
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33722 >
2025-02-26 15:19:53 +00:00
Daniel Schürmann
86fd673ade
nir: require nir_metadata_divergence if needed
...
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30814 >
2025-02-13 10:08:43 +00:00
Daniel Schürmann
8d1abd4996
treewide: use nir_src_is_divergent() rather than checking the divergence of the SSA
...
Without LCSSA, divergence between src and def might differ.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30787 >
2024-10-24 10:06:17 +00:00
Daniel Schürmann
c8348139fd
nir: change signature of nir_src_is_divergent()
...
Now, it takes nir_src * instead of nir_src.
Also move the implementation to nir_divergence_analysis.c.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30787 >
2024-10-24 10:06:17 +00:00
Daniel Schürmann
ce0a3fe645
nir/opt_uniform_atomics: don't preserve divergence information
...
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30787 >
2024-10-24 10:06:17 +00:00
Mary Guillemard
84d57e1fb1
nir: Move atomic_op_to_alu to common code
...
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27776 >
2024-10-21 21:47:44 +00:00
Alyssa Rosenzweig
d99c2ef059
nir/opt_uniform_atomics: add fs atomics predicated? flag
...
on agx (and mali), we predicate atomics on "if (!helper)", so doing so again in
this pass is redundant. and would cause a problem since we'd then have to lower
the "is helper inv?" flag late. so just skip the extra lowering code.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30488 >
2024-08-06 11:48:17 -04:00
Rhys Perry
810808b778
nir/opt_uniform_atomics: require block index metadata
...
is_atomic_already_optimized() uses this.
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30518 >
2024-08-06 15:04:21 +00:00
Georg Lehmann
690f880d18
nir/opt_uniform_atomics: handle inverse_ballot when detecting single lane ifs
...
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29502 >
2024-06-04 15:40:57 +00:00
Marek Olšák
1a791c1303
nir: add nir_atomic_op_ordered_add_gfx12_amd
...
for streamout
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Acked-By: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28889 >
2024-04-30 17:17:25 +00:00
Georg Lehmann
bce9bba90d
nir: add nir_scalar intrinsic helpers
...
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24656 >
2023-09-02 00:26:31 +00:00
Faith Ekstrand
b5d6b7c402
nir: Drop most uses if nir_instr_rewrite_src()
...
Generated by the following semantic patch:
@@
expression I, S, D;
@@
-nir_instr_rewrite_src(I, S, nir_src_for_ssa(D));
+nir_src_rewrite(S, D);
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24729 >
2023-08-18 01:00:15 +00:00
Faith Ekstrand
4695bebc79
nir: Drop nir_dest
...
Instead, we replace every use of it with nir_def. Most of this commit
was generated by sed:
sed -i -e 's/dest.ssa/def/g' src/**/*.h src/**/*.c src/**/*.cpp
A few manual fixups were required in lima and the nir_legacy code.
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24674 >
2023-08-14 21:22:53 +00:00
Faith Ekstrand
ed9affa02f
nir: Drop most instances of nir_ssa_dest_init()
...
Generated using the following two semantic patches:
@@
expression I, J, NC, BS;
@@
-nir_ssa_dest_init(I, &J->dest, NC, BS);
+nir_def_init(I, &J->dest.ssa, NC, BS);
@@
expression I, J, NC, BS;
@@
-nir_ssa_dest_init(I, &J->dest.dest, NC, BS);
+nir_def_init(I, &J->dest.dest.ssa, NC, BS);
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24658 >
2023-08-13 17:12:52 +00:00
Alyssa Rosenzweig
09d31922de
nir: Drop "SSA" from NIR language
...
Everything is SSA now.
sed -e 's/nir_ssa_def/nir_def/g' \
-e 's/nir_ssa_undef/nir_undef/g' \
-e 's/nir_ssa_scalar/nir_scalar/g' \
-e 's/nir_src_rewrite_ssa/nir_src_rewrite/g' \
-e 's/nir_gather_ssa_types/nir_gather_types/g' \
-i $(git grep -l nir | grep -v relnotes)
git mv src/compiler/nir/nir_gather_ssa_types.c \
src/compiler/nir/nir_gather_types.c
ninja -C build/ clang-format
cd src/compiler/nir && find *.c *.h -type f -exec clang-format -i \{} \;
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Acked-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24585 >
2023-08-12 16:44:41 -04:00
Faith Ekstrand
777d336b1f
nir: clang-format src/compiler/nir/*.[ch]
...
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24382 >
2023-08-12 19:27:28 +00:00
Alyssa Rosenzweig
190b1fdc64
nir: Convert to nir_foreach_function_impl
...
Done by hand at each call site but going very quickly with funny Vim motions and
common regexes. This is a very common idiom in NIR.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23807 >
2023-06-27 22:44:04 +00:00
Alyssa Rosenzweig
815efcdf7e
nir: Use nir_builder_create
...
perl -p0e 's/nir_builder ([^;]*);\s*nir_builder_init\(&\1, /nir_builder \1 = nir_builder_create(/g' -i $(git grep -l nir_builder_init)
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23860 >
2023-06-27 18:13:02 +00:00
Alyssa Rosenzweig
01e9ee79f7
nir: Drop unused name from nir_ssa_dest_init
...
Since 624e799cc3https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23078 >
2023-05-17 23:46:16 +00:00
Alyssa Rosenzweig
af69d0b1d1
nir/opt_uniform_atomics: Drop legacy atomics
...
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/23036 >
2023-05-16 22:36:22 +00:00
Alyssa Rosenzweig
55f7fd6d0d
nir/opt_uniform_atomics: Handle unified atomics
...
This is the one place where using nir_atomic_op instead of nir_op directly is a
little annoying, since we need to translate between the two enums, but it's not
a big deal.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22914 >
2023-05-12 20:39:46 +00:00
Alyssa Rosenzweig
7f6491b76d
nir: Combine if_uses with instruction uses
...
Every nir_ssa_def is part of a chain of uses, implemented with doubly linked
lists. That means each requires 2 * 64-bit = 16 bytes per def, which is
memory intensive. Together they require 32 bytes per def. Not cool.
To cut that memory use in half, we can combine the two linked lists into a
single use list that contains both regular instruction uses and if-uses. To do
this, we augment the nir_src with a boolean "is_if", and reimplement the
abstract if-uses operations on top of that list. That boolean should fit into
the padding already in nir_src so should not actually affect memory use, and in
the future we sneak it into the bottom bit of a pointer.
However, this creates a new inefficiency: now iterating over regular uses
separate from if-uses is (nominally) more expensive. It turns out virtually
every caller of nir_foreach_if_use(_safe) also calls nir_foreach_use(_safe)
immediately before, so we rewrite most of the callers to instead call a new
single `nir_foreach_use_including_if(_safe)` which predicates the logic based on
`src->is_if`. This should mitigate the performance difference.
There's a bit of churn, but this is largely a mechanical set of changes.
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22343 >
2023-04-07 23:48:03 +00:00
Marcin Ślusarz
6e0bcc1c4d
nir: use nir_metadata_none instead of its value
...
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12282 >
2022-09-26 11:13:03 +00:00
Rhys Perry
5c038b3f02
nir: add _amd global access intrinsics
...
These are the same as the normal ones, but they take an unsigned 32-bit
offset in BASE and another unsigned 32-bit offset in the last source.
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14124 >
2022-04-13 16:23:35 +00:00
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
