This should be ignored on non-compute stages but AGX changes 3D shaders
to compute without setting the workgroup size and blows up if it claims
variable workgroups. The safest thing is to only set it from
spirv_to_nir for stages that actually have workgroups.
Fixes: 6d9f563960 ("spirv: Assume variable workgroup size unless it's set")
LoLed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38555>
We don't have to enter the lower-IO-to-temps block for TCS at all.
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38470>
This potentially results in better code because we don't add def uses where
undef is allowed.
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38468>
This fixes an issue a bunch of different components were all working
around themselves where sometimes we don't have a workgroup size but
workgroup_size_variable is false. This also fixes asahi_clc, which
didn't have the workaround and was assuming zero (but not variable!)
workgroup sizes everywhere.
LoLed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Acked-by: Mel Henning <mhenning@darkrefraction.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38538>
This commit also checks for and issues errors for the following:
INVALID_OPERATION is generated if the application attempts enable pixel
local storage while the value of SAMPLE_BUFFERS is one.
INVALID_OPERATION is generated if the application attempts to enable pixel
local storage while the current draw framebuffer is a user-defined frame-
buffer object and has an image attached to any color attachment other than
color attachment zero.
INVALID_OPERATION is generated if the application attempts to enable pixel
local storage while the current draw framebuffer is a user-defined frame-
buffer and the draw buffer for any color output other than color
output zero is not NONE.
INVALID_FRAMEBUFFER_OPERATION is generated if the application attempts to
enable pixel local storage while the current draw framebuffer is
incomplete.
INVALID_OPERATION is generated if pixel local storage is disabled and the
application attempts to issue a rendering command while a program object
that accesses pixel local storage is bound.
INVALID_OPERATION is generated if pixel local storage is enabled and the
application attempts to bind a new draw framebuffer, delete the currently
bound draw framebuffer, change color buffer selection via DrawBuffers, or
modify any attachment of the currently bound draw framebuffer including
their underlying storage.
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/37110>
Shaders might declare PLS vars as inout but might just use them as in
or out but not both. This pass detects those cases and adjusts the
variable/deref modes accordingly.
This pass should be called before nir_lower_io_vars_to_temporaries(),
otherwise the copy_derefs will be inserted, turning unused variables
into used ones.
This should ideally be called after DCE to make sure we don't leave
PLS inout variables behind.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37110>
Pixel local storage variables are like fragment shader outputs that
might be read, written or both. Teach nir_lower_io_vars_to_temporaries()
about these variables so they can be lowered along with the regular
fragment outputs.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.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/37110>
Rather than adding another boolean to optionally lower PLS vars, pass
the types we want to lowers through a nir_variable_mode bitmask.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37110>
The pixel local storage load and store instructions keep track of the
format of the pixel local storage variables. This allows drivers to insert
the appropriate conversions on load/store.
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37110>
This adds some new call operations to handle various parts of the
reductions.
cmat_reduce: is the initial toplevel operation from SPIR-V
this is used after lowering for row/col operation on single hw
supported matrix sizes. The spir-v operation is lowered into
multiple of these on flex dimensions, but also can be lowered into
others.
cmat_reduce_finish:
after multiple reduction operations on a flexible dimension matrix,
there is often subsequent operations on the output matrices to
finish the operation.
cmat_reduce_2x2:
this takes 4 input matrices, and 1 dst to do a 2x2 reduction op.
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38389>
With coopmat2 a bunch of functions need a lot of lowering passes
to happen before they can be lowered, so mark them as to be lowered
later.
Drivers needing these should call the nir_remove_non_cmat_call_entrypoints
where they remove entrypoints now, and call the original nir_remove_non_entrypoints
after lowering coopmat2.
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38389>
This adds a new instruction type to handle cooperative matrix calls.
This clones the call instr, drops callee, and adds a single metadata
slot and a call operation (dummy only for now).
(Not NACKed by Alyssa)
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38389>
This calls nir_separate_merged_clip_cull_io in zink, which is better
than having to handle separate clip & cull arrays in all passes.
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38452>
Only needed by zink. This clip/cull distance separation pass is needed
to remove nir_io_separate_clip_cull_distance_arrays, so that all shared
GLSL code only uses merged clip+cull distance outputs.
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38452>
This aborts if a pass would make any progress. It can be used to assert that:
- our minimalist pass invocation loops in drivers are sufficient and don't
leave any unoptimized code in the shader
- our lowering is sufficient and other passes don't add instructions that
would cause lowering having to be repeated
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Acked-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38406>
Based on existing softfloat64 support and Berkeley SoftFloat. This is
targeted at drivers that can't preserve denorms, so operations where
denorm support is irrelevant like conversions to/from integers aren't
handled.
Because the existing mechanism used by Gallium for softfloat64 doesn't
support includes, we unfortunately can't extract common code into a
header. This can be done later if we switch Gallium to using glslang and
spirv-to-nir.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37608>
We have to invert the condition to select the non-NaN source, instead of
just swapping a and b. The equivalent float32 function was failing
dEQP-VK.spirv_assembly.instruction.graphics.float_controls.fp32.input_args.denorm_nmax_nan_preserve_frag
without this fix.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37608>
Either we need to save this pointer or toss it.
==146166==ERROR: AddressSanitizer: heap-use-after-free on address 0x7bfe77013920 at pc 0x7b9e6fd5b978 bp 0x7ffc30ef18e0 sp 0x7ffc30ef18d8
READ of size 4 at 0x7bfe77013920 thread T0
#0 0x7b9e6fd5b977 in get_header ../src/util/ralloc.c:83
#1 0x7b9e6fd5b977 in ralloc_parent ../src/util/ralloc.c:382
#2 0x7b9e6fd5b977 in reralloc_size ../src/util/ralloc.c:198
#3 0x7b9e6fd5b977 in reralloc_array_size ../src/util/ralloc.c:241
#4 0x7b9e705f83c2 in range_minimum_query_table_resize ../src/util/range_minimum_query.c:21
#5 0x7b9e7018af1d in realloc_info ../src/compiler/nir/nir_dominance_lca.c:33
#6 0x7b9e7018af1d in nir_calc_dominance_lca_impl ../src/compiler/nir/nir_dominance_lca.c:126
#7 0x7b9e6ff9815c in nir_metadata_require ../src/compiler/nir/nir_metadata.c:42
#8 0x7b9e6ff998e4 in nir_metadata_require_most ../src/compiler/nir/nir_metadata.c:200
#9 0x7b9e6f8aab4d in st_finalize_nir ../src/mesa/state_tracker/st_glsl_to_nir.cpp:735
#10 0x7b9e6f0afb14 in st_create_common_variant ../src/mesa/state_tracker/st_program.c:858
#11 0x7b9e6f0be2d3 in st_get_common_variant ../src/mesa/state_tracker/st_program.c:973
#12 0x7b9e6f0bf9cf in st_precompile_shader_variant ../src/mesa/state_tracker/st_program.c:1478
#13 0x7b9e6f0bf9cf in st_finalize_program ../src/mesa/state_tracker/st_program.c:1596
#14 0x7b9e6f8b0127 in st_link_glsl_to_nir ../src/mesa/state_tracker/st_glsl_to_nir.cpp:633
#15 0x7b9e6f8b3611 in st_link_shader ../src/mesa/state_tracker/st_glsl_to_nir.cpp:816
#16 0x7b9e6f7bcf51 in link_program ../src/mesa/main/shaderapi.c:1412
#17 0x7b9e6f7bcf51 in link_program_error ../src/mesa/main/shaderapi.c:1474
#18 0x0000004020b0 in main._omp_fn.0 /home/alyssa/shader-db/run.c:872
#19 0x7f9e7893dd65 in GOMP_parallel (/lib64/libgomp.so.1+0xdd65) (BuildId: 9cc501fdca53b5d4ab094f709486781c98573bc9)
#20 0x000000400d6a in main /home/alyssa/shader-db/run.c:689
#21 0x7f9e78011574 in __libc_start_call_main (/lib64/libc.so.6+0x3574) (BuildId: 48c4b9b1efb1df15da8e787f489128bf31893317)
#22 0x7f9e78011627 in __libc_start_main@GLIBC_2.2.5 (/lib64/libc.so.6+0x3627) (BuildId: 48c4b9b1efb1df15da8e787f489128bf31893317)
#23 0x000000401014 in _start (/home/alyssa/shader-db/run+0x401014) (BuildId: a83b8d830cc265be3f54ea3e7a21a0fb5156624b)
0x7bfe77013920 is located 0 bytes inside of 64-byte region [0x7bfe77013920,0x7bfe77013960)
freed by thread T0 here:
#0 0x7f9e782e5beb in free.part.0 (/usr/lib64/libasan.so.8+0xe5beb) (BuildId: cab80046dbc1c97c6e14490acc37d079701f8d9a)
#1 0x7b9e6fd5bc39 in unsafe_free ../src/util/ralloc.c:319
#2 0x7b9e6fd5bc39 in ralloc_free ../src/util/ralloc.c:264
#3 0x7b9e70063d81 in nir_sweep ../src/compiler/nir/nir_sweep.c:219
#4 0x7b9e6f0bf499 in st_finalize_program ../src/mesa/state_tracker/st_program.c:1585
#5 0x7b9e6f8b0127 in st_link_glsl_to_nir ../src/mesa/state_tracker/st_glsl_to_nir.cpp:633
#6 0x7b9e6f8b3611 in st_link_shader ../src/mesa/state_tracker/st_glsl_to_nir.cpp:816
#7 0x7b9e6f7bcf51 in link_program ../src/mesa/main/shaderapi.c:1412
#8 0x7b9e6f7bcf51 in link_program_error ../src/mesa/main/shaderapi.c:1474
#9 0x0000004020b0 in main._omp_fn.0 /home/alyssa/shader-db/run.c:872
previously allocated by thread T0 here:
#0 0x7f9e782e5e4b in realloc.part.0 (/usr/lib64/libasan.so.8+0xe5e4b) (BuildId: cab80046dbc1c97c6e14490acc37d079701f8d9a)
#1 0x7b9e6fd5a883 in resize ../src/util/ralloc.c:167
#2 0x7b9e705f83c2 in range_minimum_query_table_resize ../src/util/range_minimum_query.c:21
#3 0x7b9e7018af1d in realloc_info ../src/compiler/nir/nir_dominance_lca.c:33
#4 0x7b9e7018af1d in nir_calc_dominance_lca_impl ../src/compiler/nir/nir_dominance_lca.c:126
#5 0x7b9e6ff9815c in nir_metadata_require ../src/compiler/nir/nir_metadata.c:42
#6 0x7b9e6ff998e4 in nir_metadata_require_most ../src/compiler/nir/nir_metadata.c:200
#7 0x7b9e6f8b0ede in st_link_glsl_to_nir ../src/mesa/state_tracker/st_glsl_to_nir.cpp:550
#8 0x7b9e6f8b3611 in st_link_shader ../src/mesa/state_tracker/st_glsl_to_nir.cpp:816
#9 0x7b9e6f7bcf51 in link_program ../src/mesa/main/shaderapi.c:1412
#10 0x7b9e6f7bcf51 in link_program_error ../src/mesa/main/shaderapi.c:1474
#11 0x0000004020b0 in main._omp_fn.0 /home/alyssa/shader-db/run.c:872
Fixes: 17876a00af ("nir: Add a faster lowest common ancestor algorithm")
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Mel Henning <mhenning@darkrefraction.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38412>
This is done by grep ALIGN( to align(
docs,*.xml,blake3 is excluded
Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38365>
We add a bunch of new helpers to avoid the need to touch >parent_instr,
including the full set of:
* nir_def_is_*
* nir_def_as_*_or_null
* nir_def_as_* [assumes the right instr type]
* nir_src_is_*
* nir_src_as_*
* nir_scalar_is_*
* nir_scalar_as_*
Plus nir_def_instr() where there's no more suitable helper.
Also an existing helper is renamed to unify all the names, while we're
churning the tree:
* nir_src_as_alu_instr -> nir_src_as_alu
..and then we port the tree to use the helpers as much as possible, using
nir_def_instr() where that does not work.
Acked-by: Marek Olšák <maraeo@gmail.com>
---
To eliminate nir_def::parent_instr we need to churn the tree anyway, so I'm
taking this opportunity to clean up a lot of NIR patterns.
Co-authored-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38313>
gcc has a a false positive here, silenced with the pragmas, use separate commit
for easily revert latter once gcc fixed it.
Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38313>
Previously the type info for nested values was copied from the source
operand, rather than propagating the new type from the destination
operand.
Fixes: 4c363acf94 ("vtn: Allow for OpCopyLogical with different but compatible types")
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38248>
On Mali, we need not only clamp but also convert to float16 on Valhall+.
We could have a separate pass for this but it fits in nicely with the
rest of nir_lower_point_size() so we might as well put it there.
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38379>
The backend has been fully ignoring all writemasks for a long time,
so it really doesn't make sense to have them on our custom intrinsics.
I'm not sure they even make sense for some of the block intrinsics.
Also, the store_ssbo -> store_ssbo_intel pass was not setting writemask
at all, leaving it at the default value of 0 (aka write nothing, if it
had been respected...)
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38343>
