Commit graph

5961 commits

Author SHA1 Message Date
Job Noorman
60413e11c2 ir3: optimize subgroup operations using brcst.active
Follow the blob and optimize subgroup operation using brcst.active and
getlast when supported.

The transformation consists of two parts. First, a NIR transform
replaces subgroup operations with a sequence of new brcst_active_ir3
intrinsics followed by a new [type]_clusters_ir3 intrinsic (where type
can be reduce, inclusive_scan, or exclusive_scan).

The brcst_active_ir3 intrinsic is lowered directly to a brcst.active
instruction. The other intrinsics get lowered to a new macro
(OPC_SCAN_CLUSTERS_MACRO) which later gets emitted as a loop (using
getlast/getone) that iterates all clusters and produces the requested
scan result.

OPC_SCAN_CLUSTERS_MACRO has a number of optional arguments. First, since
the exclusive scan result is not a natural by-product of the loop but
has to be calculated explicitly, its destination is optional. This is
necessary since adding it unconditionally will produce unused
instructions that won't be DCE'd anymore at this point. Second, when
performing 32b MUL_U reductions (that expand to multiple instructions),
an extra scratch register is necessary.

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/6387
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26950>
2024-02-02 19:49:22 +00:00
Konstantin Seurer
c925b6019d radv/rt: Lower ray payloads like hit attribs
Reviewed-by: Friedrich Vock <friedrich.vock@gmx.de>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27051>
2024-02-02 16:36:15 +00:00
Ian Romanick
c8ba2bc2f0 nir: Pack texture LOD and array index to a single 32-bit value
v2: Fix clamped_ai calculation in nir_lower_tex.c. Add
nir_tex_src_combined_lod_and_array_index_intel to
print_tex_instr. Suggested by Sagar.

Reviewed-by: Sagar Ghuge <sagar.ghuge@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27305>
2024-02-02 02:39:10 +00:00
Konstantin Seurer
e3c2dc2324 nir/print: Rename workgroup-size to workgroup_size
Every other field uses _ instead of -.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27318>
2024-01-30 21:19:40 +00:00
Konstantin Seurer
449e44d6d3 nir/print: Don't print shared_size twice
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27318>
2024-01-30 21:19:40 +00:00
Gert Wollny
0ab3b3c641 nir/builder: Fix compilation with gcc-13 when tsan is enabled
../src/compiler/nir/nir_builder.h: In function ‘nir_build_deref_follower’:
../src/compiler/nir/nir_builder.h:1607:1: error: control reaches end of non-void function [-Werror=return-type]
 1607 | }

Fixes: 4a4e175738
    nir: Support deref instructions in lower_var_copies

Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27345>
2024-01-30 20:42:07 +00:00
Gert Wollny
80a1b91601 nir/lower_int64: Fix compilation with gcc-13 and tsan enabled
../src/compiler/nir/nir_lower_int64.c: In function ‘lower_int64_intrinsic’:
../src/compiler/nir/nir_lower_int64.c:1347:1: error: control reaches end of non-void function [-Werror=return-type]
1347 | }

Fixes: bf7a114246
   nir/lower_int64: Add lowering for some 64-bit subgroup ops

Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27345>
2024-01-30 20:42:07 +00:00
Faith Ekstrand
48ebfeba34 nak: Add a source barrier intrinsic
This just inserts a GPU stall until the given source is available.  We
need this in order to properly implement shader clock.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27303>
2024-01-26 16:55:50 +00:00
Friedrich Vock
9f22b95956 nir: Handle casts in nir_opt_copy_prop_vars
Cc: mesa-stable

Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27197>
2024-01-24 12:39:48 +00:00
Friedrich Vock
6c845ed548 nir: Make is_trivial_deref_cast public
Cc: mesa-stable

Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27197>
2024-01-24 12:39:48 +00:00
Rhys Perry
e465ac2561 nir/lower_shader_calls: remove CF before nir_opt_if
Otherwise, opt_if_simplification() can attempt to insert an inot after a
jump.

Fixes RADV compilation of a Cyberpunk 2077 pipeline with
PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27193>
2024-01-23 19:02:03 +00:00
Rhys Perry
015b0d678f nir/lower_non_uniform: set non_uniform=false when lowering is not needed
Fixes RADV compilation of a Doom Eternal pipeline with
PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT, because
nir_opt_non_uniform_access was skipped and later passes don't expect
non-uniform access.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Fixes: b1619109ca ("nir/lower_non_uniform: remove non_uniform flags after lowering")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27192>
2024-01-23 18:09:39 +00:00
Karol Herbst
f2b7c4ce29 nir: rework and fix rotate lowering
No driver supports urol/uror on all bit sizes. Intel gen11+ only for 16
and 32 bit, Nvidia GV100+ only for 32 bit. Etnaviv can support it on 8,
16 and 32 bit.

Also turn the `lower` into a `has` option as only two drivers actually
support `uror` and `urol` at this momemt.

Fixes crashes with CL integer_rotate on iris and nouveau since we emit
urol for `rotate`.

v2: always lower 64 bit

Fixes: fe0965afa6 ("spirv: Don't use libclc for rotate")
Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by (Intel and nir): Ian Romanick <ian.d.romanick@intel.com>

Reviewed-by: David Heidelberg <david.heidelberg@collabora.com>
Acked-by: Yonggang Luo <luoyonggang@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27090>
2024-01-22 10:27:44 +00:00
Georg Lehmann
d641750573 nir: add lowering for boolean shuffle
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27116>
2024-01-19 20:13:34 +00:00
Georg Lehmann
1cb5bf7009 nir: add ballot_relaxed and as_uniform intrinsics
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27116>
2024-01-19 20:13:33 +00:00
Alyssa Rosenzweig
3a72fc1cb7 nir/passthrough_gs: plug leak
freeing the nir shader should free the xfb info too. found with valgrind
leakcheck.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Reviewed-by: Antonino Maniscalco <antonino.maniscalco@collabora.com>
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27093>
2024-01-19 09:10:29 +00:00
Faith Ekstrand
82fe981e35 nir,spirv: Add support for SPV_NV_shader_sm_builtins
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27154>
2024-01-18 20:20:06 +00:00
Karol Herbst
36012af17f nir/printf: remove treat_doubles_as_floats
It is broken and clang uses fp32 for float constants if the fp64 extension
isn't enabled anyway. SPIRVs can't use fp64 constants with printf unless
they enable the Float64 cap, which also requires cl_khr_fp64 to be
supported.

So just remove it and rely on clang handling -cl-single-precision-constant
correctly, which at the moment doesn't seem to be the case, but we can
think about that once we plan to support cl_khr_fp64.

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26541>
2024-01-18 13:16:13 +01:00
Lionel Landwerlin
a18ea091af nir/comparison_pre_tests: update expectations
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Sviatoslav Peleshko <sviatoslav.peleshko@globallogic.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27087>
2024-01-17 16:01:12 +02:00
Lionel Landwerlin
873fe637e2 nir/alu_srcs_negative_equal: bail earlier if possible
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Sviatoslav Peleshko <sviatoslav.peleshko@globallogic.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27087>
2024-01-17 16:00:30 +02:00
Ian Romanick
4740ee8d67 nir: Minor clean up in nir_alu_srcs_negative_equal
Eliminate some cruft left after a8013644a1.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27087>
2024-01-17 16:00:30 +02:00
Sviatoslav Peleshko
6b0bfdfa9e nir: Use alu source components count in nir_alu_srcs_negative_equal
When we use source from ALU instruction directly, the default swizzle array
should be populated with the same amount of components as the src has.

Otherwise, if we use nir_ssa_alu_instr_src_components, it can return
the destination components count that is lower than component index
actually used in that source. This can lead to false equality
between 0 (uninitialized) and 0 (.x) in swizzle comparison below.

Fixes: c6ee46a7 ("nir: Add nir_alu_srcs_negative_equal")
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/8704
Signed-off-by: Sviatoslav Peleshko <sviatoslav.peleshko@globallogic.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22655>
2024-01-17 08:05:30 +00:00
Alyssa Rosenzweig
8fd18c4f20 nir/lower_flatshade: fix metadata
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
fcae4b469f nir/lower_io_arrays_to_elements: return prog
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
70fd20d2bc nir/lower_passthrough_edgeflags: return progress
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
460d2ca4f3 nir/lower_point_size_mov: return prog
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
8b7d765e59 nir/lower_alpha_test: rewrite with intrinsics_pass
returns progress now

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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
086cbe5da2 nir/lower_bitmap: return prog
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
4833e42721 nir: return prog from drawpixels
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
6fa32b5b83 nir/lower_clip_cull_distance_arrays: return prog
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
a36812d9b4 nir/lower_io_to_temporaries: return prog
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
caffc3abca nir/lower_blend: return progress
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/26976>
2024-01-12 01:13:02 +00:00
Alyssa Rosenzweig
29bd0a8ffa nir/lower_ssbo: rewrite
This pass was a mess. Rewrite it as modern NIR, fixing the metadata issues in
the process.

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/26976>
2024-01-12 01:13:02 +00:00
Yonggang Luo
0b9c96562b treewide: Use util_is_power_of_two_nonzero{64|_uintptr} when needed
Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26909>
2024-01-11 16:45:57 +00:00
Matt Turner
4ed0957ce7 nir/tests: Reenable tests that failed on big-endian
These tests were disabled due to the bug fixed in the previous commit.

Reviewed-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/26964>
2024-01-10 21:47:30 +00:00
Matt Turner
5997cf7587 nir: Fix cast
We were wrongly telling `nir_const_value_as_uint()` that `iter` had
`bit_size` bits, but in one case it is explicitly i64. This works on
little endian platforms, but caused the nir_loop_unroll_test.fadd{,_rev}
tests to fail on big endian platforms.

Bug: https://bugs.gentoo.org/921297
Fixes: 268ad47c11 ("nir/loop_analyze: Handle bit sizes correctly in calculate_iterations")
Reviewed-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/26964>
2024-01-10 21:47:30 +00:00
Alyssa Rosenzweig
8ddd89ffa5 nir,zink: Redefine flat_mask in terms of I/O locations
Robust against separable shaders, and still makes sense for lowered I/O drivers,
whereas just counting FS variables and expecting them to match with the VS is...
questionable.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Signed-off-by: antonino <antonino.maniscalco@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26888>
2024-01-10 14:30:14 +00:00
Alyssa Rosenzweig
97f9f7ab0a asahi: implement point sprites w/o shader key
we can replace varyings with point sprites, we just need to fix up .zw
appropriately. do that with some bcsels, ALU is cheap.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26963>
2024-01-10 08:44:38 -04:00
Caio Oliveira
e0eea5ea4e nir: Disable -Wmisleading-indentation when compiling with GCC
When a file is too large, -Wmisleading-indentantion will give the warning
below, that we can't prevent from a #pragma:

```
src/compiler/nir/nir_opt_algebraic.c: In function ‘nir_opt_algebraic’:
src/compiler/nir/nir_opt_algebraic.c:1469069: note: ‘-Wmisleading-indentation’ is disabled from this point onwards, since column-tracking was disabled due to the size of the code/headers
1469069 |    nir_foreach_function_impl(impl, shader) {
        |
src/compiler/nir/nir_opt_algebraic.c:1469069: note: adding ‘-flarge-source-files’ will allow for more column-tracking support, at the expense of compilation time and memory
```

See https://gcc.gnu.org/bugzilla/show_bug.cgi?id=89549 for details.

Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Reviewed-by: Eric Engestrom <eric@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25315>
2024-01-09 01:40:22 +00:00
Daniel Schürmann
c1ef6037fd nir/gather_info: fix enumeration of wide subgroup intrinsics
nir_intrinsic_ballot_* are no subgroup operations.
nir_intrinsic_rotate was missing.
nir_intrinsic_mbcnt_amd is not a subgroup operation.
nir_intrinsic_writelane_amd only affects a single invocation.

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18249>
2024-01-08 10:01:47 +00:00
Daniel Schürmann
d434a127f9 nir/opt_move_discards_to_top: don't schedule discard/demote across subgroup operations
Fixes: b447f5049b ('nir: Add a discard optimization pass')
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18249>
2024-01-08 10:01:47 +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
e86ab8173b nir/algebraic: optimize vkd3d-proton's MSAD
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
d32daa3fb2 nir/validate: allow bias on nir_texop_lod
AGX seems to support it, and it's very convenient for implementing sampler LOD
bias together with a clamped LOD query.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26861>
2024-01-04 01:51:07 +00:00
Daniel Schürmann
bf43af984a nir/opt_loop_cf: generalize removal of "trivial" continues
So that is also handles break statements and works in
arbitrarily nested control flow.

Totals from 905 (1.18% of 76636) affected shaders: (RADV, GFX11)
Instrs: 605164 -> 605548 (+0.06%); split: -0.01%, +0.08%
CodeSize: 3162036 -> 3163472 (+0.05%); split: -0.01%, +0.06%
Latency: 2045559 -> 1387622 (-32.16%)
InvThroughput: 352344 -> 231676 (-34.25%)
SClause: 16092 -> 16088 (-0.02%); split: -0.04%, +0.02%
Copies: 41286 -> 41297 (+0.03%); split: -0.02%, +0.05%
Branches: 19949 -> 19929 (-0.10%)
PreSGPRs: 33413 -> 33385 (-0.08%)
PreVGPRs: 19177 -> 19135 (-0.22%)

Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24940>
2024-01-03 20:48:05 +00:00
Daniel Schürmann
bdbf873b0f nir: remove redundant passes from nir_opt_if()
These are now covered by nir_opt_loop():
- opt_if_loop_last_continue()
- opt_merge_breaks()
- opt_if_loop_terminator()

Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24940>
2024-01-03 20:48:05 +00:00
Daniel Schürmann
5b1b5cd794 nir: remove nir_opt_trivial_continues()
This pass is superseded by nir_opt_loop()

Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24940>
2024-01-03 20:48:04 +00:00
Daniel Schürmann
9808ef0349 nir/opt_loop: move loop control-flow optimizations into separate pass
This new pass aims to simplify loop control-flow by reducing the number
of break and continue statements. It also supersedes nir_opt_trivial_continues().

For this purpose, it implements 3 optimizations:
- opt_loop_terminator(), as previously
- opt_loop_merge_break_continue(), similar to opt_merge_breaks() incl. continues
- opt_loop_last_block(), a generalization of opt_if_loop_last_continue()

Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24940>
2024-01-03 20:48:04 +00:00
Christian Gmeiner
0158075b22 nir/opt_peephole_select: handle speculative ubo loads
Some platforms may be able to speculate ubo loads safely.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8299>
2024-01-03 20:02:25 +00:00