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>
Some infinite loop cases were already covered by other
restrictions (e.g. if the loop had a body), but the case with a single
block in the loop body wasn't yet.
This prevents an infinite loop when optimizing the shader in
dEQP-VK.reconvergence.subgroup_uniform_control_flow_ballot.compute.nesting2.3.2
and various others reconvergence tests.
Fixes: 0881e90c09 ("nir: Split ALU instructions in loops that read phis")
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com> [v1]
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11476>
This is now 100% equivalent to the new rt_resume intrinsic.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8637>
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>
Really copying Jason's pass.
Changes:
- Instead of all the intel lowering introduce rt_{execute_callable,trace_ray,resume}
- Add the ability to use scratch intrinsics directly.
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10339>
About half or more of the text here is actually from Connor Abbot. I've
edited it a bit to bring it up-to-date and make a few things more clear.
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11438>
Long ago, the semantics of bcsel were such that it took a single boolean
value and selected between whole vectors. These days, it takes a vector
boolean with the assumption that if you want the old behavior you can
just use a .xxxx swizzle. There currently are no opcodes which use a
output_size of 0 but have a scalar or fixed-vector input. Let's
disallow it for now to force us to think through the semantics again if
this ever comes up as something someone actually wants.
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11438>
For debug on Android, it's useful to be able to print shaders to the
android log interface, since you don't usually have stdout/stderr.
Reviewed-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9262>
limit==0 is the signal for "don't peephole anything but a move that will
be optimized aways." limit > 0 is "up to N alu instructions may be moved
out." nir-to-tgsi uses ~0 as the indicator of "No, we really need to
eliminate all if instructions" on hardware like i915 that doesn't have
control flow.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11329>
We need to make get it updated after we may have nir_instr_remove()d an
instruction, and when we cross blocks. This didn't really matter before
because the only builder usage was idiv, which other users of
lower_int_to_float were probably never hitting.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Adam Jackson <ajax@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11329>
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>
builtin_builder::~builtin_builder() and builtin_builder::release()
are running into race condition. This leads lightsmark to crash at
the end because both calls ralloc_free which mutates the arguments state
This patch fixes lightsmark2008 crash
Fixes: e4da8b9c33 ("mesa/compiler: rework tear down of builtin/types")
Reviewed-by: Charmaine Lee <charmainel@vmware.com>
Reviewed-by: Neha Bhende <bhenden@vmware.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Tested-by: Neha Bhende <bhenden@vmware.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11385>
Use a single set and ensure dominance by checking after a equivalent
instruction is found.
Besides removing the need to copy a set, this also lets us resize the set
at the start of the pass in the next commit.
ministat (CSE only):
Difference at 95.0% confidence
-984.956 +/- 28.8559
-6.90075% +/- 0.190231%
(Student's t, pooled s = 26.9052)
ministat (entire run):
Difference at 95.0% confidence
-1246.1 +/- 257.253
-0.998972% +/- 0.205094%
(Student's t, pooled s = 239.863)
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Co-authored-by: Daniel Schürmann <daniel@schuermann.dev>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6390>
Now that there's a common NIR pass, there's no point in us doing this in
the back-end anymore. In order to use this pass in i965, we do have to
make one tiny change. Gallium runs the pass after assigning input and
output locations and so needs the pass to respect those locations and
num_inputs. i965, however, runs it before any location assignment or
I/O lowering so we don't care. We do, however, need the pass to succeed
with num_inputs == 0 because we set that later.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11313>
In theory you can rerun the info gather pass, but in practice that
doesn't always end well. Be consistent inside this pass and update the
info.
While we're here, change the inputs read to use VERT_BIT_EDGEFLAG.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11313>
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>
These map directly to v_perm_b32 and v_permlane_b32.
Unfortunately there is no corresponding NIR opcode or
intrinsics, and it's too tedious to puzzle these things
together from the existing NIR instructions.
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>
NIR currently doesn't have any intrinsics for a horizontal packed add,
so this one is modeled after AMD's v_sad_u8.
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>
The helpers will be reused for per-primitive variables that are also
arrayed, so use a more general name.
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/11252>
At best, this is an extra instruction for NIR to optimize out. At worst,
depending on pass ordering nir_load_output could sneak into the final
NIR, even on drivers that don't support fbfetch.
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11255>
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>
if we have
if ... {
return;
} else {
// block X
}
// block Y
phi(X: ...)
then nir_lower_returns tries to move block Y into the else body,
except nir_cf_extract doesn't move the phi. As the return is removed
in the then-body the phi suddenly has the wrong number of arguments
(and the phi doesn't dominate its uses anymore).
In this case we know that the phi has to be single arg, so we can just
rewrite the users of the phis and drop them.
Hit this in my RT adventures, not sure if this is actually reachable
right now, as single arg phis tend to be kind of exceptional outside
of CSSA and we typically call nir_lower_returns pretty early.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11207>
Found in some sottr shaders (originally iand(ishr(a, 16), 0xffff))
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
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>
On AGX, the special register for front facing is inverted from its meaning in
APIs. We need to lower load_front_face to inot(load_back_face). Doing this in
the backend is trivial, but then we would miss out on algebraic optimizations
for the inot.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11199>
