Commit graph

5099 commits

Author SHA1 Message Date
Tapani Pälli
8736d1a9a6 intel/compiler: implement macl part of Wa_18035690555
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Fixes: 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/40941>
2026-04-19 11:39:10 +00:00
Matt Turner
d36a578bc0 intel/elk: Remove dead TXL_LZ/TXF_LZ opcodes
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
These opcodes were only emitted for Gen9+ hardware, but elk only targets
Gen8 and below.

Fixes: 05d78994a7 ("intel/elk: Remove Gfx9+ sampler messages and modes")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41009>
2026-04-17 15:13:00 +00:00
Matt Turner
6856ebe00d intel/elk: Remove some dead code
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/41009>
2026-04-17 15:13:00 +00:00
Mauro Rossi
cc44922048 intel/jay: fix static_assert expression
Some checks are pending
macOS-CI / macOS-CI (dri) (push) Waiting to run
macOS-CI / macOS-CI (xlib) (push) Waiting to run
Fixes the following building error:

FAILED: src/intel/compiler/jay/libintel_compiler_jay.a.p/jay_assign_flags.c.o
...
In file included from ../src/intel/compiler/jay/jay_assign_flags.c:6:
../src/intel/compiler/jay/jay_builder.h:184:24: error: static_assert expression is not an integral constant expression
         static_assert(sizeof(uintptr_t) <= sizeof(uint64_t) &&
                       ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1 error generated.

Fixes: e42e3193 ("intel: add Jay")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
1b029f3279 jay: allow cmod on cvt
it's just a MOV

saves an instruction on dEQP-GLES31.functional.compute.basic.image_atomic_op_local_size_8

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
9df62df55e jay: fix bfn cmod
affects dEQP-GLES31.functional.compute.basic.image_atomic_op_local_size_8

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
b9f8467855 jay: fix a bunch of opcode properties
really need a full audit..

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
8afcbbe77d jay: load_simd_width_intel
dEQP-GLES31.functional.shaders.arrays_of_arrays.es31.array_access.dynamic_expression_access_fragment

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
4eb838eb48 jay: split up jay_from_nir.c
Big monolithic file, split it up into the relevant pieces.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
6925d9ee23 jay: move deswizzle hack outside of swsb
this will eventually enable better swsb for the simd32 payload code.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
48a24f3c27 jay: fix instr counts
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Kenneth Graunke
3a1227f36b jay: Clear default group for quad swizzles
Quad swizzles should always execute with NoMask and group 0.

We skipped initializing the group, and so inherited whatever the state
from the previous instruction was.  This led to incorrect behavior if
the previous instruction was SIMD split:

   (16)        mov.u32 g84<2>, g2        |
   (16|M16)    mov.u32 g86<2>, g3        |
   (32|M16&W)  mov.u32 g2, g126.2<4,4,0> | I@1

Oops.  The final quad swizzle shouldn't have had M16 set.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Kenneth Graunke
8bfb139033 jay: Make lower_immediates bail if there are no sources
For example, JAY_OPCODE_PRELOAD has no sources.  Reading src[other]
would read uninitialized data.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Kenneth Graunke
a79931421c jay: Assert that source is not null in jay_copy_strided
Catch bugs earlier.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
e84cba531b jay: fix simd split swsb bugs
this is subtle, but the relevant igc:

  // In case of shooting down of this instruction, we need to add sync to
  // preserve the swsb id sync, so that it's safe to clear the dep
  if (currInst.hasPredication() ||
      (currInst.getExecSize() != dep.getInstruction()->getExecSize()) ||
      (currInst.getChannelOffset() != dep.getInstruction()->getChannelOffset()))
    needSyncForShootDownInst = true;

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
e66bfbd691 jay: fix SEND scoreboarding
by inspection in a glmark2 shader.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
10fbfe1f65 jay: improve vector affinities
Totals:
Instrs: 2807558 -> 2805545 (-0.07%); split: -0.10%, +0.02%
CodeSize: 45051504 -> 45015888 (-0.08%); split: -0.11%, +0.03%

Totals from 631 (23.84% of 2647) affected shaders:
Instrs: 736113 -> 734100 (-0.27%); split: -0.37%, +0.10%
CodeSize: 11666608 -> 11630992 (-0.31%); split: -0.42%, +0.12%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
dd6efffb47 jay: generalize alignment heuristic
Totals:
Instrs: 2807560 -> 2805375 (-0.08%); split: -0.27%, +0.20%
CodeSize: 45051536 -> 45011488 (-0.09%); split: -0.28%, +0.19%

Totals from 1292 (48.81% of 2647) affected shaders:
Instrs: 2154123 -> 2151938 (-0.10%); split: -0.36%, +0.26%
CodeSize: 34658096 -> 34618048 (-0.12%); split: -0.36%, +0.25%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
e54898fdc8 jay: tweak roundrobin
Totals:
Instrs: 2809726 -> 2807560 (-0.08%); split: -0.39%, +0.31%
CodeSize: 45085472 -> 45051536 (-0.08%); split: -0.41%, +0.33%

Totals from 2229 (84.21% of 2647) affected shaders:
Instrs: 2726162 -> 2723996 (-0.08%); split: -0.40%, +0.32%
CodeSize: 43775120 -> 43741184 (-0.08%); split: -0.42%, +0.34%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
8aeeaa4d85 jay: marginally improve send splitting heuristic
Instrs: 2810815 -> 2809726 (-0.04%); split: -0.07%, +0.04%
CodeSize: 45101440 -> 45085472 (-0.04%); split: -0.07%, +0.04%

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
820e3a9403 jay: roundrobin RA
Totals:
Instrs: 2827788 -> 2810815 (-0.60%); split: -1.12%, +0.52%
CodeSize: 45449488 -> 45101440 (-0.77%); split: -1.28%, +0.51%
Number of spill instructions: 1984 -> 1982 (-0.10%)
Number of fill instructions: 2272 -> 2270 (-0.09%)

Totals from 2449 (92.52% of 2647) affected shaders:
Instrs: 2818824 -> 2801851 (-0.60%); split: -1.13%, +0.53%
CodeSize: 45314880 -> 44966832 (-0.77%); split: -1.28%, +0.51%
Number of spill instructions: 1984 -> 1982 (-0.10%)
Number of fill instructions: 2272 -> 2270 (-0.09%)

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
57a00707a1 jay: drop GRF reg stats
doesn't match what brw does, will revisit for Xe3.

Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
4ab67f80e9 jay: rematerialize address regs
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +00:00
Alyssa Rosenzweig
9a9365ff9b jay: fix W-entry calcs
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/40960>
2026-04-14 23:14:07 +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
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
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
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