fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-18 20:18:06 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2
219538 commits 66 branches 935 tags 952 MiB
Commit graph

133 commits

Author SHA1 Message Date
Rhys Perry
e43caba5f4 nir/range_analysis: use sparse array for float analysis 
This seems to be faster.

ministat (nir_analyze_fp_range):
Difference at 95.0% confidence
        -592900 +/- 2302.24
        -27.6432% +/- 0.0998961%
        (Student's t, pooled s = 2719.05)

ministat (overall):
Difference at 95.0% confidence
        -76.8333 +/- 27.2345
        -0.632558% +/- 0.223407%
        (Student's t, pooled s = 46.867)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40190>
 2026-03-05 11:26:25 +00:00
Rhys Perry
aecbb2a903 nir/range_analysis: use function pointers for lookup 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40190>
 2026-03-05 11:26:25 +00:00
Rhys Perry
2731c34891 nir/range_analysis: use SSA index for hash table keys 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40190>
 2026-03-05 11:26:25 +00:00
Rhys Perry
5e376e3ed2 nir: add nir_fp_analysis_state 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40190>
 2026-03-05 11:26:25 +00:00
Rhys Perry
c0079e09ca nir/range_analysis: set deleted key 
If (uintptr_t)&deleted_key is small enough, inserting entries into the
hash table might not work correctly.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Backport-to: 26.0
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40190>
 2026-03-05 11:26:25 +00:00
Georg Lehmann
bca5aab2be nir: let nir_analyze_fp_range take a nir_def 
This is midly worse for vector constants, but so much simpler.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39756>
 2026-02-16 18:08:53 +00:00
Georg Lehmann
474af815ff nir: rename nir_analyze_range because it's float only 
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39756>
 2026-02-16 18:08:53 +00:00
Georg Lehmann
f2a59fdea6 nir: remove non float nir_analyse_range support 
This was always unused/unfinished.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39756>
 2026-02-16 18:08:53 +00:00
Marek Olšák
0a9bdcac79 ac: lower load_workgroup_ids for ACO in NIR 
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Acked-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39638>
 2026-02-13 15:33:19 +00:00
Emma Anholt
10ba7675c8 nir/uub: Use an optional max_samples from drivers for sample counts. 
This triggers some unrolling in Fallout 4, GTAV, and Rocky Planet in my
shader-db.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38585>
 2025-12-11 14:26:11 +00:00
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
Details
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Details 
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
Details
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Details 
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