spirv_to_nir sometimes wraps derefs in vec2 or mov instructions as part of
its texture handling. These get in the way of
nir_rematerialize_derefs_in_use_blocks_impl. Running copy propagation
should get rid of the extra move instructions and get us back to intact
deref chains for everything except variable pointer use-cases.
fossil-db (Sienna Cichlid):
Totals from 6 (0.00% of 134572) affected shaders:
CodeSize: 92656 -> 93088 (+0.47%)
Instrs: 17060 -> 17138 (+0.46%)
Latency: 224408 -> 227539 (+1.40%)
InvThroughput: 37402 -> 37924 (+1.40%)
VClause: 408 -> 402 (-1.47%)
Copies: 1065 -> 1107 (+3.94%)
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/5668
Fixes: 14a12b771d ("spirv: Rework our handling of images and samplers")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13924>
This fixes CLs 1.2 1Darray and 2Darray images.
Fixes: 589d918a4f
("spirv: Add 0.5 to integer coordinates for OpImageSampleExplicitLod")
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Signed-off-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13401>
Instead of using gsamplerND types for sampled images, use the new
gtextureND types for sampled images and reserve gsamplerND for combined
image+samplers. Combined image+sampler bindings still get a gsamplerND
type.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13389>
Use the new nir_var_mem_image mode for images that are not known to be
used with a sampler (i.e. storage images).
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4743>
Now that all drivers have been patched to convert sysvals to input
varyings when they have too, we can safely declare PointCoord as a sysval
too.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13017>
Now that all spirv_to_nir() users take care of converting sysvals to
varyings, we can unconditionally declare FragCoord as a sysval.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13017>
No shader-db or fossil-db changes on any Intel platform.
v2: Keep the flt <-> fge switcharoo local to the SpvOpFUnordLessThan,
etc. handling. Add a comment explaining why the suboptimal
SpvOpFUnordEqual implementation is used here. Suggested by Caio.
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12320>
If an OpVariable's initializer is undef, there is no need to
initialize the variable.
v2: Comment the code (Caio)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13030>
The LLVM-SPIRV translator creates variables with initializers, but
most of those are actually undef initializers. We can just skip
composites that are entirely made of undefs, but for partially undefs,
we will still zero initialize.
v2: Rename wa_llvm_spirv_undef_initializer to wa_llvm_spirv_ignore_workgroup_initializer (Caio)
Limit workaround to OpenCL (Caio)
Make workaround clearer (Caio)
v3: Only apply workaround on workgroup storage (Caio)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13030>
Since only Anv uses the value, I'm only enabling this on anv.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Fixes: 518693c3ec ("spirv: Handle the SubgroupSize execution mode")
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13034>
v2 (Ivan): Add missing capability enum handling.
v3 (idr): Properly handle cases where dest_size != 32.
v4 (idr): Rewrite most of the error checking to use vtn_fail_if. Use
nir_ssa_def with vtn_push_nir_ssa instead of vtn_ssa_value with
vtn_push_ssa_value. All suggested by Jason. Massive rewrite of the
handling of packed 4x8 saturating opcodes. Based on some observations
made by Jason.
v5 (idr): Remove some debugging cruft accidentally added in v4. Noticed
by Jason.
v6: Emit packed versions of vectored instructions when possible.
Suggested by Jason.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12142>
This shouldn't do anything but will make testing a later patch easier.
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8056>
This adds support for SpvOpAtomicFlag operations.
This is just a simple implementation that lowers
Clear to Store 0
and
TestAndSet to Cas (0, -1)
There are likely platforms/hw that will want to
lower this all the way through NIR and into their
backend, but this will do for now.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12229>
This allows the variables decorated with RelaxedPrecision to have the
proper precision. It is worth to note that the decorator can be
applied on other cases, but those would be handled on the future.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7614>
It's intel-specific, used to get at MSAA compression information.
Reviewed-by: Emma Anholt <emma@anholt.net>
Reviewed-by: Adam Jackson <ajax@redhat.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11775>
There's no SPIR-V Capability associated, so check in the Execution
Mode. For now, don't keep track of whether a shader uses uniform
control flow in the shader_info, we can add that when/if a driver
actually need that information.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11476>
spirv_to_nir now requires NIR variables to be created for everything.
Fixes: 10b3eecd36 "spirv: Don't remove variables used by resource..."
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8637>
During dEQP runs for radv, I see a lot of warnings like,
ERROR - dEQP error: SPIR-V WARNING:
ERROR - dEQP error: In file ../src/compiler/spirv/spirv_to_nir.c:1073
ERROR - dEQP error: Decoration not allowed on struct members: SpvDecorationRestrict
ERROR - dEQP error: 408 bytes into the SPIR-V binary
This fails jobs on Gitlab, due to,
Job's log exceeded limit of 4194304 bytes.
Job execution will continue but no more output will be collected.
Since it doesn't seem feasible right now to fix the many shaders in
the VK-CTS triggering this warning, add an environment toggle that
allows test runners to only see the level of commentary they want.
v2 from Martin:
- Add my SoB
v3 from Martin:
- fix the indentation (suggested by Eric)
- put the declarations at the top of the function
v4 from Martin:
- make vtn_default_log_level() static (Marcin)
- cache the default level in vtn_log (Marcin)
- move vtn_log_level_strings inside vtn_default_log_level()
- Fix the build issue on MSC
Signed-off-by: Martin Peres <martin.peres@mupuf.org>
Reviewed-by: Eric Engestrom <eric@engestrom.ch>
Acked-by: Andres Gomez <agomez@igalia.com>
Acked-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Acked-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11491>
When an OpBranchConditional that had two equal branches was parsed, we
were treating it as a regular OpBranch. However this doesn't work
well when there's an associated OpSelectionMerge. We ended up
skipping marking the merge block as such, and depending on what was
inside the construct we would end up trying to process the block
twice.
Fix this by keeping the vtn_if around, but when emitting NIR identify
the two equal branch case.
Fixes: 9c2a11430e ("spirv: Rewrite CFG construction")
Closes: #3786, #4580
Reviewed-by: Yevhenii Kolesnikov <yevhenii.kolesnikov@globallogic.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9297>
The v_mbcnt instructions can take an extra source that they add to
the result. This is not exposed in SPIR-V but we now expose it in NIR.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072>
Move it out the "cs" sub-struct, since the bit will be used for other
shader stages in the future.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225>
Be consistent with other usages in Vulkan and SPIR-V, and the recently
added workgroup_size field.
Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
VS outputs are "per vertex" but not the kind of I/O we want to match
with this helper. Change to a name that covers the "arrayness"
required by the type.
Name inspired by the GLSL spec definition of arrayed I/O.
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10493>
