Commit graph

123 commits

Author SHA1 Message Date
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
Alyssa Rosenzweig
b824ef83ab util/dynarray: infer type in append
Most of the time, we can infer the type to append in
util_dynarray_append using __typeof__, which is standardized in C23 and
support in Jesse's MSMSVCV. This patch drops the type argument most of
the time, making util_dynarray a little more ergonomic to use.

This is done in four steps.

First, rename util_dynarray_append -> util_dynarray_append_typed

    bash -c "find . -type f -exec sed -i -e 's/util_dynarray_append(/util_dynarray_append_typed(/g' \{} \;"

Then, add a new append that infers the type. This is much more ergonomic
for what you want most of the time.

Next, use type-inferred append as much as possible, via Coccinelle
patch (plus manual fixup):

    @@
    expression dynarray, element;
    type type;
    @@

    -util_dynarray_append_typed(dynarray, type, element);
    +util_dynarray_append(dynarray, element);

Finally, hand fixup cases that Coccinelle missed or incorrectly
translated, of which there were several because we can't used the
untyped append with a literal (since the sizeof won't do what you want).

All four steps are squashed to produce a single patch changing every
util_dynarray_append call site in tree to either drop a type parameter
(if possible) or insert a _typed suffix (if we can't infer). As such,
the final patch is best reviewed by hand even though it was
tool-assisted.

No Long Linguine Meals were involved in the making of this patch.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38038>
2025-10-24 18:32:07 +00:00
Ian Romanick
f7939f2fdc nir/range_analysis: Handle bfi and bitfield_select in get_alu_uub
I noticed some things related to this while implementing support for
bitfield_select / BFN in BRW.

shader-db:

Lunar Lake
total instructions in shared programs: 17183140 -> 17183128 (<.01%)
instructions in affected programs: 3830 -> 3818 (-0.31%)
helped: 6 / HURT: 0

total cycles in shared programs: 889936934 -> 889936056 (<.01%)
cycles in affected programs: 253758 -> 252880 (-0.35%)
helped: 4 / HURT: 2

No shader-db changes on any other Intel platform.

fossil-db:

Lunar Lake
Totals:
Instrs: 233285343 -> 233284796 (-0.00%); split: -0.00%, +0.00%
Cycle count: 32756777978 -> 32756399804 (-0.00%); split: -0.00%, +0.00%
Max live registers: 71738646 -> 71738626 (-0.00%)
Non SSA regs after NIR: 67837900 -> 67837902 (+0.00%)

Totals from 177 (0.02% of 790723) affected shaders:
Instrs: 389849 -> 389302 (-0.14%); split: -0.14%, +0.00%
Cycle count: 356341872 -> 355963698 (-0.11%); split: -0.11%, +0.01%
Max live registers: 39364 -> 39344 (-0.05%)
Non SSA regs after NIR: 70453 -> 70455 (+0.00%)

Meteor Lake, DG2, and Ice Lake had similar results. (Meteor Lake shown)
Totals:
Instrs: 264095611 -> 264095358 (-0.00%)
Cycle count: 26555705299 -> 26554303407 (-0.01%); split: -0.01%, +0.00%
Fill count: 613233 -> 613231 (-0.00%)

Totals from 123 (0.01% of 905547) affected shaders:
Instrs: 334830 -> 334577 (-0.08%)
Cycle count: 326531667 -> 325129775 (-0.43%); split: -0.65%, +0.22%
Fill count: 4145 -> 4143 (-0.05%)

Tiger Lake and Skylake had similar results. (Tiger Lake shown)
Totals:
Instrs: 269733849 -> 269733590 (-0.00%)
Cycle count: 25240548036 -> 25241435039 (+0.00%); split: -0.00%, +0.01%

Totals from 123 (0.01% of 903812) affected shaders:
Instrs: 338617 -> 338358 (-0.08%)
Cycle count: 326605644 -> 327492647 (+0.27%); split: -0.13%, +0.40%

Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37186>
2025-10-10 17:25:08 +00:00
Georg Lehmann
714a149396 nir: remove unsigned upper bound config
All config information is now either in nir->info or nir->options.

Reviewed-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/37361>
2025-09-16 09:24:04 +00:00
Georg Lehmann
bb67dae12d nir/uub: remove max_workgroup_size from config
For most hardware, this is the same as max invocations in the workgroup.

Reviewed-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/37361>
2025-09-16 09:24:04 +00:00
Georg Lehmann
f3c08c9d27 nir/uub: use shader_info subgroup size
Reviewed-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/37361>
2025-09-16 09:24:04 +00:00
Georg Lehmann
112e160946 nir/uub: remove vertex input handling
Unused since radv started lowering vertex inputs in NIR.

Reviewed-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/37361>
2025-09-16 09:24:04 +00:00
Rhys Perry
09ab7ff01e nir: add nir_def_num_lsb_zero
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36760>
2025-08-22 15:45:55 +00:00
Georg Lehmann
e24db36f20 nir/uub: handle bit_count
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36874>
2025-08-21 10:36:09 +00:00
Georg Lehmann
aff391bc77 nir/uub: handle more reduction ops
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36874>
2025-08-21 10:36:09 +00:00
Georg Lehmann
773ee60e48 nir/uub: decrease default max subgroup size to 128
128 is the maximum all apis allow.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36874>
2025-08-21 10:36:09 +00:00
Georg Lehmann
a2e48d2ede nir/uub: fix exclusive scans
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>

Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36874>
2025-08-21 10:36:09 +00:00
Georg Lehmann
de3d04dd72 nir/uub: guard against division by 0
Fixes: 8ee5440073 ("nir/uub: improve ishl/imul with constant sources")

Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36805>
2025-08-19 15:49:57 +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
Georg Lehmann
3168ebe2c5 nir/range_analysis: look through vec2
Foz-DB Navi31:
Totals from 11 out of 14 FSR4 shaders:
Instrs: 58987 -> 58298 (-1.17%)
CodeSize: 402844 -> 397836 (-1.24%)
Latency: 209630 -> 209634 (+0.00%); split: -0.66%, +0.66%
InvThroughput: 230240 -> 229152 (-0.47%); split: -0.48%, +0.00%
VClause: 838 -> 826 (-1.43%); split: -1.55%, +0.12%
Copies: 3019 -> 2954 (-2.15%); split: -2.82%, +0.66%
VALU: 50196 -> 49637 (-1.11%)
VOPD: 1950 -> 1916 (-1.74%); split: +0.72%, -2.46%

Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36468>
2025-08-01 20:29:29 +00:00
Georg Lehmann
caf89c97de nir/range_analysis: look through f2f
Foz-DB Navi31:
Totals from 93 (0.12% of 80273) affected shaders:
Instrs: 123927 -> 121073 (-2.30%); split: -2.30%, +0.00%
CodeSize: 670832 -> 653332 (-2.61%); split: -2.61%, +0.00%
Latency: 337678 -> 322803 (-4.41%); split: -4.41%, +0.00%
InvThroughput: 63277 -> 61083 (-3.47%)
VClause: 460 -> 373 (-18.91%)
SClause: 2178 -> 2100 (-3.58%)
Copies: 7637 -> 7744 (+1.40%)
PreSGPRs: 4414 -> 4287 (-2.88%)
PreVGPRs: 4229 -> 4230 (+0.02%)
VALU: 77375 -> 75693 (-2.17%)
SALU: 16497 -> 16383 (-0.69%); split: -0.73%, +0.04%
VMEM: 561 -> 477 (-14.97%)
SMEM: 3197 -> 3113 (-2.63%)

Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36468>
2025-08-01 20:29:28 +00: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
Antonio Ospite
ddf2aa3a4d build: avoid redefining unreachable() which is standard in C23
In the C23 standard unreachable() is now a predefined function-like
macro in <stddef.h>

See https://android.googlesource.com/platform/bionic/+/HEAD/docs/c23.md#is-now-a-predefined-function_like-macro-in

And this causes build errors when building for C23:

-----------------------------------------------------------------------
In file included from ../src/util/log.h:30,
                 from ../src/util/log.c:30:
../src/util/macros.h:123:9: warning: "unreachable" redefined
  123 | #define unreachable(str)    \
      |         ^~~~~~~~~~~
In file included from ../src/util/macros.h:31:
/usr/lib/gcc/x86_64-linux-gnu/14/include/stddef.h:456:9: note: this is the location of the previous definition
  456 | #define unreachable() (__builtin_unreachable ())
      |         ^~~~~~~~~~~
-----------------------------------------------------------------------

So don't redefine it with the same name, but use the name UNREACHABLE()
to also signify it's a macro.

Using a different name also makes sense because the behavior of the
macro was extending the one of __builtin_unreachable() anyway, and it
also had a different signature, accepting one argument, compared to the
standard unreachable() with no arguments.

This change improves the chances of building mesa with the C23 standard,
which for instance is the default in recent AOSP versions.

All the instances of the macro, including the definition, were updated
with the following command line:

  git grep -l '[^_]unreachable(' -- "src/**" | sort | uniq | \
  while read file; \
  do \
    sed -e 's/\([^_]\)unreachable(/\1UNREACHABLE(/g' -i "$file"; \
  done && \
  sed -e 's/#undef unreachable/#undef UNREACHABLE/g' -i src/intel/isl/isl_aux_info.c

Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36437>
2025-07-31 17:49:42 +00:00
Rhys Perry
a9a1da0264 nir/uub: fix 8/16-bit overflow
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Backport-to: 25.1
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/13552
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/13553
Tested-by: @LingMan
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36372>
2025-07-28 08:46:51 +00:00
Rhys Perry
ea0670dfb5 nir: simplify nir_addition_might_overflow
nir_unsigned_upper_bound is good enough that this isn't needed anymore.

No fossil-db changes.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35514>
2025-06-17 13:28:00 +00:00
Rhys Perry
f3b7ac730c nir/uub: improve ior/ixor with constant sources
No fossil-db changes.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35514>
2025-06-17 13:28:00 +00:00
Rhys Perry
ae6ad8977b nir/uub: improve iand with constant sources
fossil-db (navi21):
Totals from 9 (0.01% of 79653) affected shaders:
Instrs: 11878 -> 11868 (-0.08%)
CodeSize: 61572 -> 61508 (-0.10%)
Latency: 44585 -> 44581 (-0.01%); split: -0.02%, +0.01%
InvThroughput: 9697 -> 9660 (-0.38%)
VALU: 8889 -> 8876 (-0.15%)
SALU: 1339 -> 1342 (+0.22%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35514>
2025-06-17 13:27:59 +00:00
Rhys Perry
8ee5440073 nir/uub: improve ishl/imul with constant sources
fossil-db (navi21):
Totals from 1 (0.00% of 79653) affected shaders:
Instrs: 1339 -> 1338 (-0.07%)
CodeSize: 7244 -> 7240 (-0.06%)
Latency: 19827 -> 19822 (-0.03%)
InvThroughput: 9913 -> 9911 (-0.02%)
SALU: 419 -> 418 (-0.24%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35514>
2025-06-17 13:27:59 +00:00
Marek Olšák
069fdc6f71 nir: handle mov and bcsel in nir_def_bits_used
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
e080833478 nir: handle iand/ior opcodes recursively in nir_def_bits_used
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
a78ed8b8e8 nir: handle extract opcodes recursively in nir_def_bits_used
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
e38a0b9a05 nir: handle u2u/i2i recursively in nir_def_bits_used
to get the number of bits actually used by the uses.

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
15369a792a nir: handle mul24 in nir_def_bits_used
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
7e7ef7b8b7 nir: handle bit shifts by constants in nir_def_bits_used
useful for open-coded bitfield extracts that are not using ubfe/ibfe

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Marek Olšák
7d24a9b649 nir: handle ibfe/ubfe in nir_def_bits_used
it will be used by radeonsi

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34489>
2025-05-13 15:38:37 +00:00
Georg Lehmann
a60d61cce8 nir: improve fadd is_a_number analysis by using the range
Foz-DB Navi21:
Totals from 145 (0.18% of 79789) affected shaders:
Instrs: 168553 -> 168391 (-0.10%); split: -0.10%, +0.00%
CodeSize: 926708 -> 926684 (-0.00%)
Latency: 2210456 -> 2210329 (-0.01%); split: -0.01%, +0.00%
InvThroughput: 545992 -> 545768 (-0.04%)
SClause: 3084 -> 3085 (+0.03%)
VALU: 129521 -> 129360 (-0.12%)
SALU: 13085 -> 13084 (-0.01%)

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34125>
2025-04-22 14:23:05 +00:00
Georg Lehmann
a6fd9f488a nir: add is_a_number analysis for ffma
Foz-DB Navi21:
Totals from 508 (0.64% of 79789) affected shaders:
Instrs: 796183 -> 795838 (-0.04%)
CodeSize: 4303420 -> 4303384 (-0.00%); split: -0.00%, +0.00%
Latency: 7806095 -> 7805458 (-0.01%); split: -0.01%, +0.00%
InvThroughput: 1377028 -> 1376824 (-0.01%); split: -0.01%, +0.00%
Copies: 63297 -> 63299 (+0.00%); split: -0.00%, +0.00%
PreVGPRs: 29818 -> 29819 (+0.00%)
VALU: 562067 -> 561885 (-0.03%); split: -0.03%, +0.00%
SALU: 89896 -> 89733 (-0.18%)

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34125>
2025-04-22 14:23:05 +00:00
Georg Lehmann
cb6d035925 nir: add range analysis for ffmaz
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34125>
2025-04-22 14:23:05 +00:00
Georg Lehmann
317d07484e nir: improve fsqrt range analysis
Foz-DB Navi21:
Totals from 3 (0.00% of 79377) affected shaders:
MaxWaves: 88 -> 96 (+9.09%)
Instrs: 1058 -> 951 (-10.11%)
CodeSize: 5964 -> 5368 (-9.99%)
VGPRs: 104 -> 96 (-7.69%)
Latency: 15283 -> 14099 (-7.75%); split: -8.37%, +0.62%
InvThroughput: 4951 -> 4238 (-14.40%)
Copies: 81 -> 76 (-6.17%)
PreVGPRs: 93 -> 84 (-9.68%)
VALU: 820 -> 737 (-10.12%)
SALU: 115 -> 91 (-20.87%)

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33557>
2025-02-18 20:38:57 +00:00
Georg Lehmann
81b4629636 nir: fix frsq range analysis
Foz-DB Navi21:
Totals from 98 (0.12% of 79377) affected shaders:
Instrs: 157311 -> 157675 (+0.23%); split: -0.03%, +0.26%
CodeSize: 844296 -> 846648 (+0.28%); split: -0.00%, +0.28%
Latency: 1275467 -> 1276259 (+0.06%); split: -0.00%, +0.06%
InvThroughput: 266980 -> 267098 (+0.04%); split: -0.03%, +0.07%
Copies: 11094 -> 11093 (-0.01%)
PreVGPRs: 5945 -> 5977 (+0.54%)
VALU: 110585 -> 110953 (+0.33%); split: -0.04%, +0.38%
SALU: 18481 -> 18476 (-0.03%)

Cc: mesa-stable

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33557>
2025-02-18 20:38:56 +00:00
Georg Lehmann
25300ac18a nir: fix range analysis for frcp
Foz-DB Navi21:
Totals from 448 (0.56% of 79377) affected shaders:
Instrs: 669306 -> 669318 (+0.00%); split: -0.00%, +0.00%
CodeSize: 3736580 -> 3738840 (+0.06%); split: -0.00%, +0.06%
Latency: 5860916 -> 5860961 (+0.00%); split: -0.00%, +0.00%
InvThroughput: 1344094 -> 1344135 (+0.00%); split: -0.00%, +0.00%
VClause: 13878 -> 13879 (+0.01%)
Copies: 58538 -> 58532 (-0.01%)
VALU: 479807 -> 479820 (+0.00%); split: -0.00%, +0.00%

Cc: mesa-stable

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33557>
2025-02-18 20:38:56 +00:00
Georg Lehmann
1f3494b886 nir: range analysis for ffract
Foz-DB Navi21:
Totals from 75 (0.09% of 79377) affected shaders:
Instrs: 69239 -> 68383 (-1.24%)
CodeSize: 385088 -> 379532 (-1.44%)
Latency: 427188 -> 421729 (-1.28%); split: -1.28%, +0.00%
InvThroughput: 103086 -> 101926 (-1.13%)
VClause: 785 -> 753 (-4.08%)
SClause: 1624 -> 1598 (-1.60%)
Copies: 5679 -> 5671 (-0.14%); split: -0.72%, +0.58%
PreSGPRs: 3961 -> 3937 (-0.61%)
VALU: 51107 -> 50457 (-1.27%)
SALU: 9034 -> 8950 (-0.93%)
VMEM: 1123 -> 1091 (-2.85%)
SMEM: 2862 -> 2830 (-1.12%)

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/33557>
2025-02-18 20:38:56 +00:00
Georg Lehmann
e8b29abb25 nir: add unsigned upper bound support for fsat
Foz-DB Navi21:
Totals from 89 (0.11% of 79395) affected shaders:
Instrs: 97018 -> 96995 (-0.02%)
CodeSize: 492996 -> 492488 (-0.10%)
Latency: 504649 -> 504555 (-0.02%)
InvThroughput: 121968 -> 121875 (-0.08%)
VALU: 67427 -> 67404 (-0.03%)

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32565>
2024-12-10 20:53:53 +00:00
Georg Lehmann
e78e63e3fe nir: add unsigned upper bound support for f2i32
Foz-DB Navi21:
Totals from 649 (0.82% of 79395) affected shaders:
CodeSize: 2330592 -> 2314112 (-0.71%)
Latency: 2068161 -> 2053370 (-0.72%)
InvThroughput: 346583 -> 329425 (-4.95%)

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32565>
2024-12-10 20:53:53 +00:00
Georg Lehmann
0b366a7ab2 nir/uub: properly limit float support to 32bit
Cc: mesa-stable

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32565>
2024-12-10 20:53:53 +00:00
Ian Romanick
92befad89f nir/range_analysis: Fix errors in fmin and fmax tables
fmin(x, 0.0) must at least be le_zero, and fmax(x, 0.0) be at least be
ge_zero.

shader-db:

All Intel platforms had similar results. (Meteor Lake shown)
total instructions in shared programs: 19733226 -> 19731919 (<.01%)
instructions in affected programs: 196415 -> 195108 (-0.67%)
helped: 615 / HURT: 0

total cycles in shared programs: 916277979 -> 916265288 (<.01%)
cycles in affected programs: 2482535 -> 2469844 (-0.51%)
helped: 346 / HURT: 178

LOST:   2
GAINED: 1

fossil-db:

All Intel platforms had similar results. (Meteor Lake shown)
Totals:
Instrs: 151531355 -> 151519575 (-0.01%); split: -0.01%, +0.00%
Cycle count: 17209372399 -> 17208402120 (-0.01%); split: -0.01%, +0.01%
Max live registers: 32016490 -> 32016514 (+0.00%)

Totals from 4307 (0.68% of 630198) affected shaders:
Instrs: 4179418 -> 4167638 (-0.28%); split: -0.28%, +0.00%
Cycle count: 1063492212 -> 1062521933 (-0.09%); split: -0.24%, +0.15%
Max live registers: 359250 -> 359274 (+0.01%)

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30158>
2024-07-20 00:19:05 +00:00
Jesse Natalie
d9eacb05c9 nir_range_analysis: Use fmin/fmax to fix NAN handling
The following probable MSVC bug clued us in to some probably-unexpected behavior:
https://developercommunity.visualstudio.com/t/Incorrect-SSE-code-for-minmax-with-NaNs/10687862

Change the logic here so that we're always starting with NANs and use
fmin/fmax, which have more-deterministic handling of NANs. If one argument
is NAN, the non-NAN argument is returned. The previous code would've returned
the second argument if one was NAN.

Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29822>
2024-06-21 20:13:46 +00:00
Rhys Perry
ae54cbeb3f nir: remove sad_u8x4
All uses of this can be replaced with msad_4x8.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26907>
2024-01-05 18:55:22 +00:00
Rhys Perry
0477421f7d nir: add msad_4x8
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26907>
2024-01-05 18:55:22 +00:00
Alyssa Rosenzweig
c39896b17b nir: Use getters for nir_src::parent_*
First, we need to give the parent_instr field a unique name to be able to
replace with a helper.  We have parent_instr fields for both nir_src and
nir_def, so let's rename nir_src::parent_instr in preparation for rework.

This was done with a combination of sed and manual fix-ups.

Then we use semantic patches plus manual fixups:

    @@
    expression s;
    @@

    -s->renamed_parent_instr
    +nir_src_parent_instr(s)

    @@
    expression s;
    @@

    -s.renamed_parent_instr
    +nir_src_parent_instr(&s)

    @@
    expression s;
    @@

    -s->parent_if
    +nir_src_parent_if(s)

    @@
    expression s;
    @@

    -s.renamed_parent_if
    +nir_src_parent_if(&s)

    @@
    expression s;
    @@

    -s->is_if
    +nir_src_is_if(s)

    @@
    expression s;
    @@

    -s.is_if
    +nir_src_is_if(&s)

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24671>
2023-10-10 04:58:05 -04: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
b781dd6200 nir s/nir_get_ssa_scalar/nir_get_scalar/
Generated with sed:

    sed -i -e 's/nir_get_ssa_scalar/nir_get_scalar/g' src/**/*.h src/**/*.c src/**/*.cpp

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24703>
2023-08-15 17:44:27 +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
6c1d32581a nir: Drop nir_alu_dest
Instead, we replace it directly with nir_def.  We could replace it with
nir_dest but the next commit gets rid of that so this avoids unnecessary
churn.  Most of this commit was generated by sed:

   sed -i -e 's/dest.dest.ssa/def/g' src/**/*.h src/**/*.c src/**/*.cpp

There were a few manual fixups required in the nir_legacy.c and
nir_from_ssa.c as nir_legacy_reg and nir_parallel_copy_entry both have a
similar pattern.

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
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