Commit graph

222480 commits

Author SHA1 Message Date
Rhys Perry
53cd74c284 aco/ra: don't rename phi operands in get_reg_phi()
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Backport-to: *
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41018>
2026-05-13 15:53:22 +00:00
Rhys Perry
638cefc8f2 aco: fix regclasses for spill/reload subdword temporaries
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/41018>
2026-05-13 15:53:21 +00:00
Rhys Perry
46047cf5ca aco/ra: fix fill() with certain subdword cases
If div_ceil(start.byte() + num_bytes, 4) != div_ceil(num_bytes, 4), like
v3b at byte=2.

Also, the non-const operator[] was unused and unsafe.

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/41018>
2026-05-13 15:53:21 +00:00
Rhys Perry
528aabb056 aco/ra: fix compact_relocate_vars path for get_reg_for_operand
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/41018>
2026-05-13 15:53:20 +00:00
Rhys Perry
7c4429d21a aco: rework subdword definition RA validation a bit
Some issues fixed:
- check that byte=0 for p_as_uniform/etc
- validate 5+ byte definitions
- fix v3b pseudo definition validation
- fix validation when the index of the definition is not 0

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/41018>
2026-05-13 15:53:19 +00:00
Rhys Perry
4249daeedd aco: add helpers to get instruction subdword capabilities
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/41018>
2026-05-13 15:53:19 +00:00
Rhys Perry
43aeeb8b88 aco/ra: test the register file in get_reg_specified() when necessary
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Backport-to: *
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41018>
2026-05-13 15:53:17 +00:00
Rhys Perry
2d5478fc3f aco/ra: fix v3b VALU at byte>0
This can happen with v_cndmask_b32, if we were required to take the
sub-dword path in get_reg_simple().

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/41018>
2026-05-13 15:53:17 +00:00
Rhys Perry
62a268eb5f aco/validate: fix some RA validator error messages
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/41018>
2026-05-13 15:53:17 +00:00
Karol Herbst
47bde74e91 ci: update OpenCL 3.1 piglit fails
Piglit does enforce the pre 3.1 behavior of clSetKernelArg

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
7c07bbcdb5 rusticl: update names of types now core in 3.1
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
6007e35857 bin/gen_release_notes: add paragraph on OpenCL support
Reviewed-by: Eric Engestrom <eric@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
a50d2ce8d7 docs/features: add OpenCL 3.1 section
Reviewed-by: Eric Engestrom <eric@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
47dca06b2e rusticl: implement CL 3.1 device features
Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
626ec83855 rusticl: implement CL 3.1 platform features
Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
e5526bbac0 rusticl: start implementing CL 3.1 support
Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
ac5db2d42e rusticl/kernel: return CL_INVALID_WORK_GROUP_SIZE in clEnqueueNDRangeKernel for an explicit 0 workgroup
See https://github.com/KhronosGroup/OpenCL-Docs/pull/1542

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
f69fe4cd55 rusticl/kernel: update error code handling for clSetKernelExecInfo
See https://github.com/KhronosGroup/OpenCL-Docs/pull/1419

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
dc47061e11 rusticl/program: handle CL_INVALID_CONTEXT for clCompileProgram and clLinkProgram
See https://github.com/KhronosGroup/OpenCL-Docs/pull/1453

Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Karol Herbst
9e673e5773 include: update CL headers
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41358>
2026-05-13 15:21:40 +00:00
Mary Guillemard
dd97257209 nir/lower_bit_size: Preserve float controls when lowering alu ops
fp_math_ctrl should be preserved when recreating alu operations.

Signed-off-by: Mary Guillemard <mary@mary.zone>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/work_items/15455
Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41511>
2026-05-13 14:53:01 +00:00
Danylo Piliaiev
9a247643eb tu: Don't disable UBWC for D24S8+USAGE_SAMPLED+customBorderColorWithoutFormat
Apparently, this is a major footgun since it is not uncommon for apps to
enable all the features exposed by a driver. Having UBWC disabled for
D24S8 can result in a major performance loss, and the reason can be hard
for devs to spot. This footgun is already known to have happened a few
times. Furthermore, disabling UBWC depending on a Vulkan feature being
requested broke D24S8 sharing via external memory when only one device
was created with customBorderColorWithoutFormat.

Fortunately, there is the depthStencilSwizzleOneSupport feature, which
was added after the above hardware deficiency was found and, when false,
forbids the problematic state combination.

To prevent the footgun described above, we now set
depthStencilSwizzleOneSupport to false by default. This allows UBWC to be
enabled for D24S8 in all cases while remaining conformant. We also have
the tu_enable_d24s8_border_color_workaround driconf option, which enables
the previous workaround for apps that don't know about
depthStencilSwizzleOneSupport, which is currently only the ANGLE
translation layer.

One caveat is that we cannot use the fast border color HW feature for
D24S8+USAGE_SAMPLED+VK_FORMAT_UNDEFINED, so a new driconf toggle is
added. enable_fast_border_color_for_undefined_formats is set for DXVK and
vkd3d-proton since they are known not to use border colors with D24S8.
Lacking fast border colors is a much smaller penalty than not having UBWC
for D24S8.

For some context also see: https://gitlab.khronos.org/Tracker/vk-gl-cts/-/issues/4346

This partially reverts 36916949.

Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41514>
2026-05-13 14:03:00 +00:00
Danylo Piliaiev
ea8de0742b tu/a8xx: Fix reading border_color from sampler memory
Fixes: 77e83d1449 ("tu: gen8 sampler support")

Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41514>
2026-05-13 14:03:00 +00:00
Erik Faye-Lund
50fd51bb9c pan/ci: fix gitlab rules after move
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Without this, changes to the shared compiler sources doesn't trigger any
jobs. Whoops.

Fixes: b02bc53261 ("pan: Move util/* to compiler/")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41528>
2026-05-13 13:13:24 +00:00
Erik Faye-Lund
dfc06fd114 pan/ci: add missing gitlab rule
Fixes: f091bdf392 ("pan: Lift pan_get_model into its own lib")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41528>
2026-05-13 13:13:24 +00:00
Erik Faye-Lund
ffadd54940 pan/ci: remove outdated gitlab rule
This directory doesn't exist any more.

Fixes: e55de285cc ("panfrost: Kill panfrost-job.h")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41528>
2026-05-13 13:13:23 +00:00
Erik Faye-Lund
88bd52cfe3 pan/ci: add missing gitlab rules
Fixes: 20970bcd96 ("panfrost: Add base of OpenCL C infrastructure")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41528>
2026-05-13 13:13:23 +00:00
Icenowy Zheng
98d07d5319 pvr: preserve and pass more data for suspending render passes
Suspending render pass jobs have more things than render targets to
preserve, e.g. occlusion query related information, atomic / compute
overlap enablement information etc.

Preserve them too when suspending. When resuming, for boolean
properties, or'ing them; for other preserved things assign them. This is
for ensuring the last resuming fragment job is compatible with all
suspending geometry jobs, as for suspending render passes the fragment
job is omitted.

The situation of the suspending render pass and the resuming render pass
have different query pools is still not supported, and quite difficult
to support.

Backport-to: 26.0
Signed-off-by: Icenowy Zheng <zhengxingda@iscas.ac.cn>
Reviewed-by: Nick Hamilton <nick.hamilton@imgtec.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41002>
2026-05-13 12:57:23 +00:00
Icenowy Zheng
175784d1a8 pvr: add a structure containing data kept for suspended renderpasses
As more things than render targets data need to be kept for suspending
renderpasses, add a structure to sort out them.

Backport-to: 26.0
Signed-off-by: Icenowy Zheng <zhengxingda@iscas.ac.cn>
Reviewed-by: Nick Hamilton <nick.hamilton@imgtec.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41002>
2026-05-13 12:57:23 +00:00
Icenowy Zheng
7e56e4a030 pvr: stop to derive rt datasets based on geometry_terminate
As we're going to kick frag for suspending rendering passes to mitigate
frag job inconsistency between suspending rendering passes and resuming
render passes, deriving render target datasets based on
geometry_terminate property will be incorrect.

Stop to use geometry_terminate to decide whether to remember render
target datasets, instead use is_suspend directly.

In addition, is_resume is now also used instead of checking whether
suspended render taget datasets is available. This will help when either
the suspending render pass or the resuming render pass have multiple
graphics sub_cmds.

Backport-to: 26.0
Signed-off-by: Icenowy Zheng <zhengxingda@iscas.ac.cn>
Reviewed-by: Nick Hamilton <nick.hamilton@imgtec.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41002>
2026-05-13 12:57:23 +00:00
Icenowy Zheng
d7f79b34c7 pvr: copy sub_cmd flags except owned when executing subcmds out of pass
When executing a secondary command buffer outside a renderpass, the
sub_cmds of that secondary command buffer is simply copied into the
primary command buffer. However, the 4 flags outside the type-specific
structures are not copied. Although owned flag is intentionally set to
false, the other 3 flags should be preserved.

Copy these 3 flags when executing sub_cmds of a secondary command buffer
outside renderpasses.

Backport-to: 26.0
Signed-off-by: Icenowy Zheng <zhengxingda@iscas.ac.cn>
Reviewed-by: Nick Hamilton <nick.hamilton@imgtec.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41002>
2026-05-13 12:57:23 +00:00
Icenowy Zheng
351eaffdc5 pvr: fix handling of invalid attachment info in pvr_init_fs_outputs_mrt
The attachments field of the render pass state could be
MESA_VK_RP_ATTACHMENT_INFO_INVALID, which indicates no attachment
information is valid. If such situation really happens when initializing
the fragment state of a pipeline, this means neither a render pass nor a
VkPipelineRenderingCreateInfo structure is available -- in this case,
the specificiation for that structure says colorAttachmentCount is
considered as 0, so the loop iterating color attachments should just not
happen.

Skip iterating color attachments if the render pass has a attachments
field with value MESA_VK_RP_ATTACHMENT_INFO.

This fixes some regression on the Vulkan CTS testcase
dEQP-VK.pipeline.monolithic.misc.no_rendering introduced by !40870, in
which MESA_VK_RP_ATTACHMENT_INFO instead of 0 is set as the value of the
attachments field of the render pass state, if neither a render pass nor
the VkPipelineRenderingCreateInfo structure is available.

Fixes: 1950b6c1a7 ("vulkan: mark RP attachments as invalid when no rendering create info")
Signed-off-by: Icenowy Zheng <zhengxingda@iscas.ac.cn>
Acked-by: Frank Binns <frank.binns@imgtec.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41032>
2026-05-13 12:40:33 +00:00
Dmitry Osipenko
85cb633871 intel/virtio: Preserve errno properly when handling ioctl
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Avoid changing errno when ioctl succeeds.

Fixes: b06d759a93 ("intel: Add virtio-gpu native context")
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/work_items/15446
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: Dmitry Osipenko <dmitry.osipenko@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41512>
2026-05-13 11:34:11 +00:00
Timothy Arceri
0429a73f47 mesa: fix typo in validation string
We are processing the ff fragment shader not the vertex shader.

Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41524>
2026-05-13 11:01:07 +00:00
hwandy
c96e73aa93 Revert "intel/decoder: make libvulkan_intel to depend on stub decoder when buildtyle=release."
This reverts commit 2ee6b4d96e.

The previous change avoids 0.25MB (1%) size change on the driver binary file,
but blocks the runtime enablement for some intel tools which is critical
to our optimization tasks.

It's not a good tradeoff based on the new need of the tool in runtime,
so revert this change.

Test: meson setup builddir -Dallow-fallback-for=libdrm -D build-tests=true -Dbuildtype=release --reconfigure && ninja -C builddir && cd builddir && meson test

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: hwandy <hwandy@google.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41525>
2026-05-13 10:21:08 +00:00
Lorenzo Rossi
32ca5f7515 panfrost: Plumb VS varying_layout in FS
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This allows us to use LD_VAR_BUF instead of LD_VAR when the shaders are
linked together.

Signed-off-by: Lorenzo Rossi <lorenzo.rossi@collabora.com>
Reviewed-by: Christoph Pillmayer <christoph.pillmayer@arm.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40761>
2026-05-13 09:15:58 +00:00
Lorenzo Rossi
12b9a23cb5 panfrost: Split default key creation in helper function
This will make it easier to create new default keys in other places

Signed-off-by: Lorenzo Rossi <lorenzo.rossi@collabora.com>
Reviewed-by: Christoph Pillmayer <christoph.pillmayer@arm.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40761>
2026-05-13 09:15:58 +00:00
Lorenzo Rossi
1f80394c67 pan/compiler/lower_fs_inputs: Do not trust slot->alu_type
pan_varying_layout contains both layout and format, in lower_fs_inputs
though the layout is referring to the VS layout and the format might
differ from what the FS layout expects.  We cannot use the VS format as
FS format otherwise we risk interpolating an integer.

Fixes: 66bee415ad ("pan/compiler: Split lower_varyings_io into fs_inputs and vs_outputs")
Signed-off-by: Lorenzo Rossi <lorenzo.rossi@collabora.com>
Reviewed-by: Christoph Pillmayer <christoph.pillmayer@arm.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40761>
2026-05-13 09:15:58 +00:00
Faith Ekstrand
e3d5d13d6d pan/bi: Use LOD_MODE_EXPLICIT for the 2nd half of textureGrad() on Bifrost
textureGrad() has to be split into two halves on Mali: Computing the
gradient/LOD and doing the actual texture operation.  On Valhal, we do
this with LOD_MODE_GRDESC but on Bifrost, we use LOD_MOD_EXPLICIT.  When
converting to NIR, I missed this.

Fixes: 05a066c921 ("pan/nir: Add bifrost support to pan_nir_lower_tex()")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41513>
2026-05-13 08:10:23 +00:00
David Rosca
804ff19a5a radeonsi/uvd_enc: Skip extra padding bytes in output bitstream
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Trailing zeroes should be harmless, but it seems to cause issues with
latest ffmpeg (which looks like an ffmpeg bug).
The extra bytes are useless, so we can just skip them like we already
do on VCN to workaround it.

Cc: mesa-stable
Reviewed-by: Benjamin Cheng <benjamin.cheng@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41485>
2026-05-13 07:12:53 +00:00
Samuel Pitoiset
aee1043227 radv/meta: adjust an assertion for HTILE expand on SDMA with compute fallback
Because SDMA doesn't support MSAA, it's possible to get there because
RADV fallback to compute queue in this case.

Some tests only pass because RDNA2 and older don't support image
stores with depth/stencil and MSAA.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41492>
2026-05-13 06:10:02 +00:00
Alyssa Rosenzweig
db95df3da4 jay/opt_propagate: propagate undefs
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
allows deleting piles of moves & pressure.

simd16 results:

   Totals:
   Instrs: 2759547 -> 2753358 (-0.22%); split: -0.29%, +0.06%
   CodeSize: 41141280 -> 41071072 (-0.17%); split: -0.23%, +0.06%

   Totals from 332 (12.54% of 2647) affected shaders:
   Instrs: 648080 -> 641891 (-0.95%); split: -1.23%, +0.28%
   CodeSize: 9782272 -> 9712064 (-0.72%); split: -0.97%, +0.25%

simd32 is a loss because of RA being stupid. again, this is obviously the right
thing to do so we're doing it. stats are just a hint.

   Totals:
   Instrs: 4683556 -> 4689193 (+0.12%); split: -0.25%, +0.37%
   CodeSize: 70072256 -> 70171920 (+0.14%); split: -0.23%, +0.38%
   Number of spill instructions: 50320 -> 50316 (-0.01%)
   Number of fill instructions: 51530 -> 51526 (-0.01%)

   Totals from 351 (13.26% of 2647) affected shaders:
   Instrs: 1349954 -> 1355591 (+0.42%); split: -0.86%, +1.28%
   CodeSize: 20484224 -> 20583888 (+0.49%); split: -0.80%, +1.29%
   Number of spill instructions: 21762 -> 21758 (-0.02%)
   Number of fill instructions: 26328 -> 26324 (-0.02%)

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:36 +00:00
Alyssa Rosenzweig
21e527ceec jay/opt_propagate: fix NOT propagation
and add a test for it. oops.

Totals:
Instrs: 4700885 -> 4683707 (-0.37%); split: -1.36%, +1.00%
CodeSize: 70551872 -> 70285088 (-0.38%); split: -1.35%, +0.97%
Number of spill instructions: 50325 -> 50320 (-0.01%)
Number of fill instructions: 51541 -> 51530 (-0.02%)

Totals from 1261 (47.64% of 2647) affected shaders:
Instrs: 3932922 -> 3915744 (-0.44%); split: -1.63%, +1.19%
CodeSize: 59196320 -> 58929536 (-0.45%); split: -1.60%, +1.15%
Number of spill instructions: 47901 -> 47896 (-0.01%)
Number of fill instructions: 48420 -> 48409 (-0.02%)

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:36 +00:00
Alyssa Rosenzweig
5cbf0002c4 jay/register_allocate: tweak roundrobin heuristic
Totals:
Instrs: 4706214 -> 4700132 (-0.13%); split: -1.03%, +0.90%
CodeSize: 70628880 -> 70540336 (-0.13%); split: -1.02%, +0.89%

Totals from 2084 (78.73% of 2647) affected shaders:
Instrs: 4515981 -> 4509899 (-0.13%); split: -1.08%, +0.94%
CodeSize: 67822800 -> 67734256 (-0.13%); split: -1.06%, +0.93%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:35 +00:00
Alyssa Rosenzweig
37e4144693 jay/register_allocate: set num_regs[MEM] properly
this is both a correctness fix (insufficient MEM registers reserved in some
cases) and a performance fix (unnecessary allocations & zeroing in the RA when
we don't spill).

fixes dEQP-VK.dgc.ext.compute.misc.scratch_space

stats are noise but positive i guess.

Totals from 35 (1.32% of 2647) affected shaders:
Instrs: 396770 -> 396690 (-0.02%)
CodeSize: 6040832 -> 6039600 (-0.02%)

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:35 +00:00
Alyssa Rosenzweig
d67e37a24c jay/lower_scoreboard: use sbid syncs to elide regdist deps
Totals from 1522 (57.50% of 2647) affected shaders:
CodeSize: 65268400 -> 65056176 (-0.33%); split: -0.33%, +0.00%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:35 +00:00
Alyssa Rosenzweig
89e33407e4 jay/lower_scoreboard: use CFG for RegDist scoreboarding
this is now properly global.

Totals from 558 (21.08% of 2647) affected shaders:
CodeSize: 42098496 -> 42078256 (-0.05%); split: -0.05%, +0.00%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:35 +00:00
Alyssa Rosenzweig
c2a423b5b5 jay/lower_scoreboard: rename gpr_range -> key
for clarity since UGPRs are here too.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:34 +00:00
Alyssa Rosenzweig
d549fb9c04 jay/lower_scoreboard: compact inst_exec_pipe
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:34 +00:00
Alyssa Rosenzweig
adaae3baf1 jay/lower_scoreboard: control flow is int pipe
according to IGC output.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41510>
2026-05-12 22:46:34 +00:00