Commit graph

15803 commits

Author SHA1 Message Date
Caio Oliveira
356d88457a anv: Lower any remaining globals when cmat_calls are inlined
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Fixes: 3debca7dc6 ("anv: Enable cooperativeMatrixPerElementOperations")
Reviewed-by: Dave Airlie <airlied@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40938>
2026-04-14 17:00:53 +00:00
Michael Cheng
8ea53c25d4 anv: log aux disable reasons in image init and DRM modifier selection
Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Nanley Chery <nanley.g.chery@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40658>
2026-04-14 16:14:48 +00:00
Michael Cheng
34e27bbf9f anv: log aux disable and aux-skip reasons during image setup
Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Nanley Chery <nanley.g.chery@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40658>
2026-04-14 16:14:48 +00:00
Michael Cheng
44d06d855c anv: log fast depth clear fallback reasons in vkCmdClearAttachments
Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Nanley Chery <nanley.g.chery@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40658>
2026-04-14 16:14:47 +00:00
Michael Cheng
3b7e56482d anv: log fast color clear fallback reasons in vkCmdClearAttachments
Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Nanley Chery <nanley.g.chery@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40658>
2026-04-14 16:14:47 +00:00
Tapani Pälli
745dc66b69 anv: fix Wa_14024015672 interaction in blorp
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This was causing issue within hiz clears, apply workaround only
when doing RT writes (not for WM_HZ ops).

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/work_items/14533
Fixes: 7be8af1dad ("anv: deal with Wa 14024015672 on the blorp path")
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40913>
2026-04-14 03:15:14 +00:00
Michael Cheng
06c9c08c48 intel/ds: report when OA metrics are unavailable
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Promote DBG() failure paths in enumerate_sysfs_metrics() to mesa_logw()
so users see why OA metrics are unavailable without needing INTEL_DEBUG.
Also log in PPS when no OA queries are available after initialization.

Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40898>
2026-04-13 21:31:51 +00:00
Michael Cheng
16c17d6698 intel/ds: report when OA metric access is blocked by kernel policy
When observation_paranoid (xe) or perf_stream_paranoid (i915) prevents
unprivileged access to OA metrics, the existing code silently returns no
OA queries. PPS then fails with just a segfault.

This patch adds INTEL_PERF_FEATURE_OA_BLOCKED_BY_POLICY to
intel_perf_features, set by both KMD backends when the paranoid sysctl
exists but lacks sufficent privilage. PPS checks this flag immediately
after initialising intel_perf and returns an error before  attempting
metric-set selection.

Signed-off-by: Michael Cheng <michael.cheng@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40898>
2026-04-13 21:31:51 +00:00
Lionel Landwerlin
4dfedcca45 elk: don't support frontfacing ternary optimization on != 32bit
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40931>
2026-04-13 20:32:06 +00:00
Lionel Landwerlin
a84c12414c brw: don't support frontfacing ternary optimization on != 32bit
Fix shader compilation on Crimson Desert :

  16    %1995 = b32csel %1992, %1993, %1994

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40931>
2026-04-13 20:32:06 +00:00
Nanley Chery
b50bb53630 intel/blorp: Fix width scaling for YCBCR copies
Fixes: eb8883f3ef ("intel/blorp: Redescribe surfaces for copies")
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/15267
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40930>
2026-04-13 20:03:41 +00:00
Lionel Landwerlin
0927de4631 anv: enable storageInputOutput16
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40821>
2026-04-13 14:01:04 +00:00
Lionel Landwerlin
46d42b63da brw: add support for < 32bit io values
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40821>
2026-04-13 14:01:04 +00:00
Georg Lehmann
5231c924ea intel/peephole_fma: use nir_fp_no_contract instead of exact
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40872>
2026-04-12 17:10:27 +00:00
Collabora's Gfx CI Team
d15e4274f6 Uprev Piglit to 11ce9eb56edb00e6a7702d13168cc827ce5e0cbd
d0a16eee4f...11ce9eb56e

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40427>
2026-04-10 21:21:52 +00:00
Alyssa Rosenzweig
fc00e2c815 anv: wire up jay
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
e42e319313 intel: add Jay
Jay is a new SSA-based compiler for Intel GPUs. This is an early
work-in-progress. It isn't ready to ship, but we'd like to move development in
tree rather than rebasing the world every week. Please don't bother testing yet
- we know the status and we're working on it!

Jay's design is similar to other modern NIR backends, particularly ACO, NAK and
AGX. It is fully SSA, deconstructing phis after RA. We use a Colombet register
allocator similar to NAK, allowing us to handle Intel's complex register
regioning restrictions in a straightforward way. Spilling logical registers is
straightforward with Braun-Hack.

Thanks to the SSA-based design, the entire backend is essentially linear time,
regardless of register pressure, addressing brw's excessive compile time when
especially spilling with brw.

In this current early draft, we support a limited subset of all three APIs on
Xe2. A lot works and a lot doesn't. The core compiler is there (spilling,
scoreboarding, SIMD32, etc should more or less work), but there are details to
fill in for both performance and correctness. We essentially pass conformance on
OpenGL ES 3.0 and OpenCL 3.0, and we're busy iterating on Vulkan.

Likewise, additional hardware support will come down the line. There's nothing
fundamentally Xe2-specific here. I just have a Lunarlake laptop on my desk, Ken
has a Battlemage card, and we had to pick _something_ as the first target.

Co-authored-by: Kenneth Graunke <kenneth@whitecape.org>
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Kenneth Graunke
7468261d3d intel/nir: Make intel_nir_lower_sparse work for either brw or jay
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
5c4cae2661 brw: lower ifind_msb for Jay
I'm not sure brw wants the common lowering but Jay definitely does.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
3afd572d8f brw: add Jay-specific SIMD selection rule
In the future this might even do something clever.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
f33454e35a brw: disable nir_opt_uniform_atomics for Jay
While Jay supports subgroups, efficient reductions are TODO so it's probably
better not to run this pass yet.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
4778fc1ab7 brw: disable hw generate local ID for jay
Jay will need more work to handle these payloads properly especially in SIMD32.
For now just disable the optimization for Jay for correctness.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
e8b9f26561 brw: subgroup lowering for jay
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Kenneth Graunke
2780a327fa intel: add INTEL_JAY environment variable
Add a new environment variable controlling which shader stages use the
experimental compiler.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Kenneth Graunke
09089fdd13 nir: Add nir_texop_sparse_residency[_txf]_intel operations
These lowered versions map to what Jay can deal with. The hardware is more
flexible but we're not due to data model restrictions. We choose to lower to get
us off the ground, we can revisit later.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40835>
2026-04-10 18:21:21 +00:00
Alyssa Rosenzweig
181611786c brw: round up block components
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
this ensures we don't see vec5 @load_ssbo_uniform_block_intel which
requires special backend handling, instead rounding up in NIR to vec8
which the LSC can do. affects
dEQP-GLES31.functional.shaders.builtin_functions.integer.bitfieldextract.ivec3_lowp_compute.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40877>
2026-04-10 09:16:43 +00:00
Alyssa Rosenzweig
70e246d7bc brw: chop up unaligned access
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40877>
2026-04-10 09:16:43 +00:00
Alyssa Rosenzweig
c1e1cc9b01 brw: lower mem access sizes even for UBOs
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40877>
2026-04-10 09:16:43 +00:00
Alyssa Rosenzweig
9d82888383 brw: lower 16-bit mulh
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Acked-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40877>
2026-04-10 09:16:42 +00:00
Tapani Pälli
c5bfa688b4 intel/dev: update mesa_defs.json from workaround database
This updates 14024997852 with BMG and brings in media WA
16021867713.

Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Acked-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40881>
2026-04-10 09:50:41 +03:00
Kenneth Graunke
0b99c88337 nir, brw: lower scratch in NIR
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This will let us share a common scratch swizzling between brw and jay.

Changes by Ken:
- Use an immediate SIMD width when known so we don't need to re-lower
- Switch to load_simd_width_intel because it may not match
  info->api_subgroup_size on Vulkan without VK_EXT_subgroup_size_control
- Stop using DWord Scattered Write messages for scratch.  These take an
  offset in DWords, and our offsets are now always in bytes.  This also
  means that we no longer create MEMORY_OPCODE_* IR with inconsistent
  units of either bytes or dwords.  Yikes.  We use byte scattered
  messages now.

fossil-db stats on Battlemage:

   Instrs: 500477504 -> 500450056 (-0.01%); split: -0.01%, +0.00%
   CodeSize: 7807432368 -> 7806786192 (-0.01%); split: -0.01%, +0.00%
   Cycle count: 62404008370 -> 62398437734 (-0.01%); split: -0.01%, +0.00%
   Fill count: 546690 -> 546695 (+0.00%); split: -0.00%, +0.00%
   Max live registers: 141257956 -> 141258100 (+0.00%); split: -0.00%, +0.00%
   Non SSA regs after NIR: 72350283 -> 72336544 (-0.02%)

   Totals from 99 (0.01% of 1581969) affected shaders:
   Instrs: 366593 -> 339145 (-7.49%); split: -7.58%, +0.09%
   CodeSize: 6425936 -> 5779760 (-10.06%); split: -10.06%, +0.00%
   Cycle count: 2412009876 -> 2406439240 (-0.23%); split: -0.26%, +0.03%
   Fill count: 19675 -> 19680 (+0.03%); split: -0.02%, +0.04%
   Max live registers: 17600 -> 17744 (+0.82%); split: -0.09%, +0.91%
   Non SSA regs after NIR: 37894 -> 24155 (-36.26%)

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Signed-off-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Alyssa Rosenzweig
140616d26a brw: scalarize even 64-bit scratch access
No, I don't know how this worked before, thanks for asking.

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/40843>
2026-04-09 21:02:16 +00:00
Alyssa Rosenzweig
15b11635a2 brw: Move intel_nir_opt_peephole_imul32x16 later in compilation
(Split by Ken out of a patch authored by Alyssa.)

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Kenneth Graunke
e5598166b0 brw: Have brw_nir_apply_key call brw_nir_lower_simd for all stages
brw_nir_apply_key typically knows the dispatch width (it's fixed for
geometry stages, and we clone the NIR for compute and mesh shaders).
For compute/mesh, this was the very next thing called.  For the others,
if we know the width, there's no reason not to lower it.

Scratch lowering will start using load_simd_width_intel soon, so we
need it to work in all stages.

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Kenneth Graunke
765d74eebe brw: Set nir->info.{min,max}_subgroup_size in brw_nir_apply_key
This records the actual SIMD width we selected for the shader, in
all cases except fragment shaders, where we don't know it yet.

MR 37258 notes that "Backends can update [these fields] when they make
new decisions about the subgroup size" - which is what we now do.

Note that nir->info.api_subgroup_size may be different than min/max
subgroup size on Vulkan prior to SPV1.6/VK_EXT_subgroup_size_control,
so we do not alter that.

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Kenneth Graunke
d7d2d7aceb brw: Support load_simd_width_intel for fragment shaders
This lets us emit NIR code based on the SIMD size.  For non-fragment
stages, we'll replace it with a constant and optimize, but for FS,
we delay it until the backend.

Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Kenneth Graunke
cac9f670d1 intel/compiler: Use nir_static_workgroup_size helper
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40843>
2026-04-09 21:02:16 +00:00
Tapani Pälli
3ab9145393 intel/compiler: implement dummy mov for Wa_18035690555
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37804>
2026-04-09 07:30:01 +00:00
Tapani Pälli
4bb68d7474 intel/compiler: expose inferred_exec_pipe from scoreboarding
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37804>
2026-04-09 07:30:01 +00:00
Sagar Ghuge
2bf520340d intel/compiler: Remove unused brw_nir_memclear_global helper
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
This is a dead code, we can remvoe it for now.

Signed-off-by: Sagar Ghuge <sagar.ghuge@intel.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenz.ca>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40801>
2026-04-09 05:06:05 +00:00
José Roberto de Souza
1e052f0bb5 intel/brw: Remove unsed functions to get data port message type
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40832>
2026-04-08 17:44:52 +00:00
José Roberto de Souza
667a58ab38 anv: Use helper to get anv_address in emit_simple_shader_dispatch()
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40832>
2026-04-08 17:44:52 +00:00
José Roberto de Souza
a69e02d97c anv: Make use of anv_shader_get_scratch_surf() in genX_cmd_compute.c
genX_cmd_compute.c has 2 places that is had a code very similar to
anv_shader_get_scratch_surf() but we could not make use of this function without
change it parameters.

Now it takes the shader stage and the total_scratch instead of anv_shader because
cmd_buffer_trace_rays() don't have a shader.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40832>
2026-04-08 17:44:52 +00:00
José Roberto de Souza
fd420e80e2 anv: Rename and share get_scratch_surf() with other files
We will need to call get_scratch_surf() from other files, so here removing the
static and adding it to anv_private.h.

No changes in behavior expected here.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40832>
2026-04-08 17:44:51 +00:00
Kenneth Graunke
b391f2d888 anv: Use nir_lower_memory_model
This replaces NIR_MEMORY_MAKE_{AVALIABLE,VISIBLE} with COHERENT.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40833>
2026-04-08 16:07:35 +00:00
Alyssa Rosenzweig
73701c305e brw: wire up MACL
New on Xe2, this instruction enables faster 32x32 integer multiply at the cost
of extra accumulator usage. Add it to the opcode list for future use.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40833>
2026-04-08 16:07:35 +00:00
Rhys Perry
463e3643f2 nir: add and use block predecessor helpers
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40242>
2026-04-08 15:06:32 +00:00
Valentine Burley
61a2e53b7b anv/ci: Add full VKCTS pre-merge job on Raptor Lake
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
We have 11 new RPL-U Brya Chromebooks in the Collabora lab, allowing the
full VKCTS test suite to run pre-merge for the first time without a
fraction.

Signed-off-by: Valentine Burley <valentine.burley@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40782>
2026-04-08 11:11:17 +00:00
Ian Romanick
cfdb3ddb93 brw: brw_reg::nr for an accumulator is not part of the offset
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Without this, reg_offset will return 1024 for acc0. This causes
has_invalid_dst_region to decide that the destination region is invalid
(because 1024 != 0), and the lowering code tries to treat the floating
point accumulators as integers. It's a mess.

v2: Add and use set_gfx_platform. Suggested by Caio.

Fixes: 937373eb25 ("i965/fs: Handle fixed HW GRF subnr in reg_offset().")
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40716>
2026-04-08 00:36:39 +00:00
Ian Romanick
ffdc310bf1 brw/const: Don't allow type changes when accumulators are involved
Integer accumulators and float accumulators do not occupy the same bits,
so the types cannot be arbitrarily changed.

No shader-db or fossil-db changes on any Intel platform.

v2: Use is_accumulator() instead if brw_reg_is_arf(). Add an extra test
to show the desired behavior when an accumulator is not
involved. Suggested by Caio.

Fixes: 64c251bb3a ("intel/fs: Combine constants for SEL instructions too")
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40638>
2026-04-07 23:37:26 +00:00