Commit graph

11086 commits

Author SHA1 Message Date
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
Iago Toral Quiroga
ac11e00b15 nir/serialize: make alu src deserialization consistent for unused swizzles
Currently we have 3 paths for ALU serialization/deserialization in NIR:

1. If the ALU is qualified as packed_src_ssa_16bit (identity swizzle)
  - The write will store an object index only.
  - The read will only load the swizzles actually used, the rest are 0.
2. If the ALU is not qualified as packed_src_ssa_16bit, we have two cases:
  2.1 Up to vec4:
    - The write stores all 4 swizzle components.
    - The read loads all 4 swizzle components.
  2.2 vec8/16
    - The write stores only swizzle components used, the rest are 0.
    - The read loads only swizzle components used, the rest are 0.

This inconsistency in how these paths encode/decode unsused swizzle components
can cause issues in some scenarios where a backend compiler may receive
functionally equivalent NIR shaders from Mesa that won't produce the same sha1,
leading to unnecessary cache misses.

This patch makes path 2.1 always encode and decode unused swizzle components
as 0, making it consistent with the other paths.

This fixes issues where sometimes backends need to compile a shader twice
before it is effectively retrieved from the disk cache. This has been
observed at least with V3d and Panfrost.

The problem occurs when an ALU src with unused swizzle components is serialized
in the Mesa frontend using path 1, but when it later hits the backend it is
serialized using path 2.1. The backend uses the sha1 of the serialized NIR for
the cache key. On the second execution the Mesa frontend has a cache hit and
when it deserializes the alu src, it sets its unused components to 0 but that
will cause the backend to have a cache miss since that NIR doesn't match the one
it cached on the first execution.

By always making unused swizzle components decode and encode consistently to 0
in all paths we ensure the issue never happens and that NIR variants that only
differ in swizzle components that are not used lead to cache hits.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37218>
2025-09-16 07:33:34 +02:00
Georg Lehmann
76df6cd321 nir: remove has_ddx_intrinsics option
I thought this was already gone but apparently not quite.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37362>
2025-09-15 10:16:27 +00:00
Yonggang Luo
bebd167d74 glsl: Fixes warning: deprecated directive: ‘%pure-parser’, ‘%error-verbose’
../../src/compiler/glsl/glcpp/glcpp-parse.y:179.1-12: warning: deprecated directive: ‘%pure-parser’, use ‘%define api.pure’ [-Wdeprecated]
  179 | %pure-parser
      | ^~~~~~~~~~~~
      | %define api.pure
../../src/compiler/glsl/glcpp/glcpp-parse.y:180.1-14: warning: deprecated directive: ‘%error-verbose’, use ‘%define parse.error verbose’ [-Wdeprecated]
  180 | %error-verbose
      | ^~~~~~~~~~~~~~
      | %define parse.error verbose

Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Reviewed-by: Eric Engestrom <eric@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37289>
2025-09-13 08:23:07 +00:00
Christian Gmeiner
5977c0026d nak: Move dataflow to compiler crate
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: Christian Gmeiner <cgmeiner@igalia.com>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37348>
2025-09-13 07:30:50 +00:00
Georg Lehmann
2ac5641473 shader_info: remove gl_subgroup_size enum
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37258>
2025-09-12 21:05:17 +00:00
Georg Lehmann
04d3b3bde5 mesa,glsl,spirv: set new subgroup size info
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37258>
2025-09-12 21:05:17 +00:00
Georg Lehmann
ce91b0be08 nir: define new subgroup size info
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37258>
2025-09-12 21:05:17 +00:00
Georg Lehmann
95c2a65662 nir: remove unused shader_info param in nir_create_shader
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37258>
2025-09-12 21:05:17 +00:00
Konstantin Seurer
2a4b1ea69b nir/opt_ray_queries: Cleanup and return if functions is not singular
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37283>
2025-09-12 19:11:02 +00:00
Mel Henning
611c5ebce5 compiler/rust: impl AsSlice for Box
Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Reviewed-by: Seán de Búrca <leftmostcat@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37315>
2025-09-12 10:33:32 +00:00
Samuel Pitoiset
9fbf29fc6d spirv: add missing non-uniform access for SSBO atomics
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This fixes new VKCTS coverage
dEQP-VK.descriptor_indexing.non_uniform_atomics.

Found this while implementing a new extension.

Cc: mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37295>
2025-09-12 07:40:50 +00:00
David Rosca
4b54277d2e Remove VDPAU
VDPAU only supports X11 and GL interop. There is no Wayland or Vulkan
interop support. The API has limitations that makes it impossible to
correctly decode certain streams.
Application support is also very limited, and VAAPI is always a better
choice over VDPAU.

Acked-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Acked-by: Adam Jackson <ajax@redhat.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36632>
2025-09-10 12:33:57 +00:00
Timothy Arceri
11a434f3df glsl: remove now unused NumUniformRemapTable
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36997>
2025-09-10 05:11:47 +00:00
Timothy Arceri
e052254066 glsl: make use of u_range_remap for uniform remapping
This will allow ubo buffers to have arrays containing millions of
elements without excessive memory use on a remap table. Before this
change using the max sized array on radeonsi would result in 1.3GB
of memory being used for a remap table in a single shader.

There is also a small functional change here, previously if the
shader used more than GL_MAX_UNIFORM_BLOCK_SIZE mesa would ignore
and allow this as the original ARB_uniform_buffer_object spec
stated:

   "If the amount of storage required for a uniform block exceeds
   this limit, a program may fail to link."

However in OpenGL 4.3 the text was clarified and the "may" was
removed so with this change we enforce the max limit.

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9953
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36997>
2025-09-10 05:11:47 +00:00
Georg Lehmann
08b58c3fac nir/lower_subgroups: remove lower_fp64 option
This was incorrect (it also lowered int64 reductions/scans), and the only
user can just use the general callback to precisely only lower what it wants.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37164>
2025-09-09 11:09:22 +00:00
Georg Lehmann
687510495f nir: remove subgroup size related nir_shader_compiler_options members
This was added with the goal to eventually replace the per
pass subgroup/ballot size options, but that won't work because
some backends don't have a fixed subgroup size across the compilation
process.

It was also mostly added to hack around mesa state tracker behavior,
and we have a better solution there now.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37164>
2025-09-09 11:09:22 +00:00
Georg Lehmann
9bc14a0047 nir/lower_subgroup: optimize reduce/scans with unknown subgroup size
We skip iterations with ifs.
These can be optimized aways after the subgroup size is known.
Every driver should do that because applications depend on it anyway.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37164>
2025-09-09 11:09:21 +00:00
Rhys Perry
c59a85d406 nir/load_store_vectorize: remove offset check in try_vectorize_shared2
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This doesn't seem useful anymore.

fossil-db (gfx1201):
Totals from 111 (0.14% of 79839) affected shaders:
Instrs: 152356 -> 151883 (-0.31%); split: -0.35%, +0.04%
CodeSize: 808484 -> 805584 (-0.36%); split: -0.39%, +0.04%
VGPRs: 7880 -> 7844 (-0.46%); split: -0.91%, +0.46%
Latency: 4121366 -> 4120648 (-0.02%); split: -0.04%, +0.02%
InvThroughput: 814622 -> 815362 (+0.09%); split: -0.02%, +0.11%
VClause: 3066 -> 3065 (-0.03%); split: -0.10%, +0.07%
SClause: 2594 -> 2593 (-0.04%)
Copies: 9412 -> 9447 (+0.37%); split: -0.47%, +0.84%
PreSGPRs: 4012 -> 4026 (+0.35%)
PreVGPRs: 4025 -> 4070 (+1.12%); split: -0.22%, +1.34%
VALU: 80457 -> 81039 (+0.72%); split: -0.08%, +0.80%
SALU: 16542 -> 16528 (-0.08%); split: -0.10%, +0.02%
VOPD: 39 -> 44 (+12.82%)

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/36370>
2025-09-09 10:11:52 +00:00
Rhys Perry
0f364aded3 nir/opt_offsets: improve shared2 optimization
Combine additions too, instead of just constant offsets.

fossil-db (gfx1201):
Totals from 97 (0.12% of 79839) affected shaders:
Instrs: 145269 -> 144886 (-0.26%); split: -0.27%, +0.01%
CodeSize: 762184 -> 759556 (-0.34%); split: -0.36%, +0.01%
VGPRs: 5812 -> 5764 (-0.83%)
Latency: 4050681 -> 4050528 (-0.00%); split: -0.01%, +0.00%
InvThroughput: 617458 -> 617181 (-0.04%); split: -0.05%, +0.00%
Copies: 8719 -> 8672 (-0.54%); split: -0.70%, +0.16%
PreVGPRs: 3558 -> 3543 (-0.42%); split: -0.59%, +0.17%
VALU: 77793 -> 77462 (-0.43%); split: -0.44%, +0.01%
SALU: 17028 -> 17009 (-0.11%)

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/36370>
2025-09-09 10:11:51 +00:00
Rhys Perry
c10e495182 nir/opt_offsets: fix progress determination with offsets that add to zero
If the offset is iadd(iadd(iadd(a, 1), b), -1), try_extract_const_addition
will create a dead iadd(a, b) and claim that it didn't modify the shader.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Marek Olšák <maraeo@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36370>
2025-09-09 10:11:50 +00:00
Rhys Perry
9aad852af8 nir/opt_offsets: report progress if NUW is set
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Reviewed-by: Marek Olšák <maraeo@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36370>
2025-09-09 10:11:50 +00:00
Mel Henning
eba08245a8 treewide: Spell indices correctly
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
LOLed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36184>
2025-09-08 23:03:13 +00:00
Mel Henning
17876a00af nir: Add a faster lowest common ancestor algorithm
On a fossil from the blender 4.5.0 vulkan backend, this improves compile
times in nak by about 17%. Compile time of other shaders improves by a
more modest 1.2%.

No stat changes on shader-db.

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36184>
2025-09-08 23:03:13 +00:00
Mel Henning
cd06366ca2 nir/phi_builder: Adjust valid_metadata assert
so we can add more metadata to nir_metadata_control_flow.

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36184>
2025-09-08 23:03:13 +00:00
Mel Henning
ee8d448241 nir: Don't require nir_metadata_control_flow
We're about to add to nir_metadata_control_flow, and we don't want
passes to require the new metadata.

Via coccinelle:

@@
expression e1;
@@
- nir_metadata_require(e1, nir_metadata_control_flow)
+ nir_metadata_require(e1, nir_metadata_block_index | nir_metadata_dominance)

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36184>
2025-09-08 23:03:13 +00:00
Daniel Schürmann
a53190a426 nir/load_store_vectorize: hoist base addr instead of subtracting
Totals from 3130 (3.92% of 79839) affected shaders: (Navi48)
Instrs: 2634316 -> 2633652 (-0.03%); split: -0.06%, +0.04%
CodeSize: 13999784 -> 13996888 (-0.02%); split: -0.05%, +0.03%
SpillSGPRs: 1771 -> 1778 (+0.40%)
Latency: 27233464 -> 27230934 (-0.01%); split: -0.02%, +0.01%
InvThroughput: 4234587 -> 4234550 (-0.00%); split: -0.00%, +0.00%
VClause: 54684 -> 54689 (+0.01%)
SClause: 62743 -> 62912 (+0.27%); split: -0.08%, +0.35%
Copies: 162594 -> 163729 (+0.70%); split: -0.22%, +0.91%
PreSGPRs: 146909 -> 146914 (+0.00%); split: -0.01%, +0.01%
VALU: 1558771 -> 1558778 (+0.00%)
SALU: 337715 -> 338168 (+0.13%); split: -0.30%, +0.44%

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37163>
2025-09-08 09:56:04 +00:00
Rhys Perry
cfba417316 nir/load_store_vectorize: optimize accesses with u2u64(ishl.nuw(iadd))
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/37163>
2025-09-08 09:56:04 +00:00
Rhys Perry
4bc4322150 nir/load_store_vectorize: call nir_def_num_lsb_zero less
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/37163>
2025-09-08 09:56:03 +00:00
Rhys Perry
491b7e851f nir/load_store_vectorize: refactor entry key creation
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/37163>
2025-09-08 09:56:03 +00:00
Rhys Perry
8888c2471b nir/load_store_vectorize: refactor offset parsing
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/37163>
2025-09-08 09:56:03 +00:00
Daniel Schürmann
acb47d2c78 nir/load_store_vectorize: also parse offsets through u2u64 if additions don't wrap around
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37163>
2025-09-08 09:56:03 +00:00
Timothy Arceri
8b1d48cf0b nir: move nir_lower_point_size_mov() to st
Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37037>
2025-09-07 23:13:24 +00:00
Timothy Arceri
450419c3f4 nir: move nir_lower_alpha_test() to the st
This is gl specific and a following fix will add more gl specific
params so here we move it to the st to avoid filling nir.h with
more junk.

Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37037>
2025-09-07 23:13:23 +00:00
Timothy Arceri
8417f4a8eb nir: move nir_lower_drawpixels() to the state tracker
This is gl specific and a following fix will add more gl specific
params so here we move it to the st to avoid filling nir.h with
more junk.

Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37037>
2025-09-07 23:13:22 +00:00
Daniel Schürmann
c78f1d516c nir/algebraic: add pattern for (a << #b) * #c => a * (#c << #b)
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Totals from 2545 (3.19% of 79839) affected shaders: (Navi48)

Instrs: 6371003 -> 6364130 (-0.11%); split: -0.12%, +0.01%
CodeSize: 33827548 -> 33812244 (-0.05%); split: -0.06%, +0.01%
Latency: 47451755 -> 47430108 (-0.05%); split: -0.05%, +0.00%
InvThroughput: 10442450 -> 10437159 (-0.05%); split: -0.05%, +0.00%
SClause: 159829 -> 159874 (+0.03%); split: -0.01%, +0.04%
Copies: 500725 -> 500721 (-0.00%); split: -0.01%, +0.01%
PreSGPRs: 110482 -> 110478 (-0.00%); split: -0.00%, +0.00%
PreVGPRs: 147289 -> 147287 (-0.00%); split: -0.00%, +0.00%
VALU: 3456135 -> 3454241 (-0.05%); split: -0.06%, +0.01%
SALU: 925982 -> 923616 (-0.26%)
VOPD: 1243 -> 1212 (-2.49%)
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37173>
2025-09-06 10:18:42 +00:00
Rhys Perry
efe536dbe9 vtn: use vtn_has_decoration more
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/37175>
2025-09-05 15:58:03 +00:00
Christoph Pillmayer
f81f3c85e2 nir/opt_algebraic: Convert a + b + a to b + 2a
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This allows fusing into one FMA later.

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37113>
2025-09-05 11:39:51 +00:00
Lionel Landwerlin
7cbabcad36 compiler: add stage_is_graphics() helper
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34872>
2025-09-05 07:46:17 +00:00
Lionel Landwerlin
afea98593e nir: add a new intrinsic for load dynamic tessellation config
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34872>
2025-09-05 07:46:15 +00:00
Rob Clark
d5a8233598 nir/lower-amul: Comment fix
Signed-off-by: Rob Clark <rob.clark@oss.qualcomm.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37063>
2025-09-04 15:21:38 +00:00
Rob Clark
55d77749ed nir/lower-amul: Fix crash with unused SSBO
Since https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12175
we should be able to rely on driver_location for both UBOs and SSBOs.

Signed-off-by: Rob Clark <rob.clark@oss.qualcomm.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37063>
2025-09-04 15:21:38 +00:00
Georg Lehmann
796f0847a6 nir/lower_subgroups: recursively lower ballot scans
This should be better for backends that have le/lt mask intrinsics.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:04:00 +00:00
Georg Lehmann
2725eaf9a2 nir/lower_subgroups: change filter to intrinsic callback
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:04:00 +00:00
Georg Lehmann
d14897b2f7 nir/lower_subgroups: don't use get_max_subgroup_size for lowering boolean rotates
The lowering won't work with an unknown subgroup size, and we correctly
assert that at the top of the function.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:03:59 +00:00
Georg Lehmann
516c766c71 spirv: ensure ballot find_lsb/find_msb/bit_count have 32bit result
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:03:58 +00:00
Georg Lehmann
f8633511be nir: make ballot find_lsb/msb/bit_count 32bit only
The lowering is 32bit only too.

Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:03:58 +00:00
Georg Lehmann
276fce4f13 spirv: handle ballot bit_extract separately
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37178>
2025-09-04 14:03:58 +00:00