Kenneth Graunke 0b99c88337 nir, brw: lower scratch in NIR ... 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
..
brw	nir, brw: lower scratch in NIR	2026-04-09 21:02:16 +00:00
elk	intel/compiler: Use nir_static_workgroup_size helper	2026-04-09 21:02:16 +00:00
brw_device_sha1_gen_c.py	Rename sha1_* and sha_* names to blake3_*	2026-03-23 07:03:28 +00:00
brw_list.h	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_gfx_ver_enum.h	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir.c	intel/compiler: Use nir_split_conversions()	2025-04-07 17:45:21 -05:00
intel_nir.h	brw: add support for separate tessellation shader compilation	2025-09-05 07:46:17 +00:00
intel_nir_blockify_uniform_loads.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_clamp_image_1d_2d_array_sizes.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_clamp_per_vertex_loads.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_lower_non_uniform_barycentric_at_sample.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_lower_non_uniform_resource_intel.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_lower_printf.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_lower_shading_rate_output.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_lower_sparse.c	intel/nir: Replace tg4 with txl/txb/tex when splitting texture residency	2026-03-24 16:06:29 +00:00
intel_nir_opt_peephole_ffma.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_opt_peephole_imul32x16.c	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_nir_tcs_workarounds.c	nir: add and use block predecessor helpers	2026-04-08 15:06:32 +00:00
intel_prim.h	intel/compiler: Use SPDX annotations	2026-01-24 20:37:31 +00:00
intel_shader_enums.h	brw: Drop BRW_VARYING_SLOT_PAD and brw_varying_slot enum	2026-02-16 15:15:35 -08:00
meson.build	brw: Move into a new src/intel/compiler/brw subdirectory	2025-10-09 07:01:47 +00:00