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>
These are similar to AYUV, but the channel ordering is different... in
such a way that there's no RGBA format that will make the channels line
up right.
v2: Rebase on bc438c91d9 ("nir/lower_tex: ignore texture_index if
tex_instr has deref src")
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9610>
The CAP for packed transform feedback concerns packing of unrelated
variables into the same varying slot. (On Mali, transform feedback is
implemented on a per-slot basis, so different variables need different
slots to be written to different buffers.) However, this requirement is
tangential to the packing of arrays, matrices, and structures inherent
to GLSL. These array-like values need to be packed /within/ their slot,
even though drivers using the CAP (just Panfrost) cannot pack
independent values in the slot. Transform feedback of individual
elements is not independent, after all.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10778>
We need to duplicate the subscripted members even if they happen to be
aligned, since the other elements may be passed into the consumer
shader. Fixes on Panfrost:
dEQP-GLES3.functional.transform_feedback.array_element.interleaved.lines.highp_float
Note: the test did pass on main previously due to an elaborate set of
driver hacks. I don't believe the old behaviour was correct regardless.
Only Panfrost is affected by this change and the next, as every other
driver sets PIPE_CAP_PACKED_STREAM_OUTPUT.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10778>
The flt version could have been added in 56e21647e2, but our
collective understanding of NaN and comparisons was poor in 2015. The
new "is_a_number" predicate makes the others possible.
All of the helped shaders in shader-db are either from Mad Max or Skia.
Some of the Skia shaders just get decimated by this change:
instructions helped: shaders/skia/580-4.shader_test FS SIMD8: 81 -> 29 (-64.20%) (scheduled: top-down)
I looked at a couple of those shaders, and they had sequences like:
vec1 32 ssa_44 = flt32 ssa_32, ssa_32
vec1 32 ssa_45 = b32csel ssa_44, ssa_43, ssa_0
vec1 32 ssa_46 = fge32 ssa_32, ssa_32
vec1 32 ssa_47 = b32csel ssa_46, ssa_0, ssa_45
vec1 32 ssa_48 = iand ssa_46, ssa_44
vec1 32 ssa_49 = b32csel ssa_48, ssa_43, ssa_0
ssa_44 is replaced with False. Then ssa_47 selects between ssa_0 and
ssa_0, so ssa_47 and ssa_46 are eliminated. ssa_48 is (False && don't
care), so ssa_48 and ssa_49 are eliminated. After that, many
calculations now involve constants of zero, so they are optimized down
too. So it continues until there's not much left!
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
All Intel platforms had similar results. (Tiger Lake shown)
total instructions in shared programs: 21072238 -> 21071386 (<.01%)
instructions in affected programs: 33722 -> 32870 (-2.53%)
helped: 146
HURT: 1
helped stats (abs) min: 1 max: 62 x̄: 5.84 x̃: 2
helped stats (rel) min: 0.19% max: 62.35% x̄: 4.09% x̃: 1.07%
HURT stats (abs) min: 1 max: 1 x̄: 1.00 x̃: 1
HURT stats (rel) min: 0.20% max: 0.20% x̄: 0.20% x̃: 0.20%
95% mean confidence interval for instructions value: -7.94 -3.65
95% mean confidence interval for instructions %-change: -5.87% -2.25%
Instructions are helped.
total cycles in shared programs: 856203326 -> 856192238 (<.01%)
cycles in affected programs: 749966 -> 738878 (-1.48%)
helped: 148
HURT: 0
helped stats (abs) min: 1 max: 1226 x̄: 74.92 x̃: 18
helped stats (rel) min: 0.07% max: 49.70% x̄: 2.69% x̃: 0.46%
95% mean confidence interval for cycles value: -104.82 -45.02
95% mean confidence interval for cycles %-change: -4.01% -1.37%
Cycles are helped.
LOST: 4
GAINED: 0
Fossil-db results:
Tiger Lake
Instructions in all programs: 160915223 -> 160898354 (-0.0%)
SENDs in all programs: 6812780 -> 6812780 (+0.0%)
Loops in all programs: 38340 -> 38340 (+0.0%)
Cycles in all programs: 7434144207 -> 7433978462 (-0.0%)
Spills in all programs: 192582 -> 192582 (+0.0%)
Fills in all programs: 304537 -> 304537 (+0.0%)
Ice Lake
Instructions in all programs: 145296298 -> 145279531 (-0.0%)
SENDs in all programs: 6863692 -> 6863692 (+0.0%)
Loops in all programs: 38334 -> 38334 (+0.0%)
Cycles in all programs: 8800257014 -> 8800088384 (-0.0%)
Spills in all programs: 216880 -> 216880 (+0.0%)
Fills in all programs: 334248 -> 334248 (+0.0%)
Skylake
Instructions in all programs: 135891664 -> 135874910 (-0.0%)
SENDs in all programs: 6802946 -> 6802946 (+0.0%)
Loops in all programs: 38331 -> 38331 (+0.0%)
Cycles in all programs: 8444273433 -> 8444130932 (-0.0%)
Spills in all programs: 194839 -> 194839 (+0.0%)
Fills in all programs: 301114 -> 301114 (+0.0%)
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10012>
If the values are known to be numbers, the the replacements are exact.
This is only applied to the patterns with constants. Constants should
always be numbers, and shaders with NaN constants should be handled in a
different way.
No shader-db or fossil-db changes on any Intel platform. The intention
is to make these patterns more future proof.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10012>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
All Haswell and later Intel platforms had similar results. (Tiger Lake shown)
total instructions in shared programs: 21049056 -> 21048939 (<.01%)
instructions in affected programs: 4716 -> 4599 (-2.48%)
helped: 39
HURT: 0
helped stats (abs) min: 1 max: 6 x̄: 3.00 x̃: 3
helped stats (rel) min: 0.99% max: 5.43% x̄: 2.80% x̃: 2.51%
95% mean confidence interval for instructions value: -3.46 -2.54
95% mean confidence interval for instructions %-change: -3.22% -2.38%
Instructions are helped.
total cycles in shared programs: 855141411 -> 855141159 (<.01%)
cycles in affected programs: 54491 -> 54239 (-0.46%)
helped: 28
HURT: 5
helped stats (abs) min: 2 max: 34 x̄: 12.82 x̃: 12
helped stats (rel) min: 0.06% max: 2.73% x̄: 0.94% x̃: 0.75%
HURT stats (abs) min: 2 max: 52 x̄: 21.40 x̃: 6
HURT stats (rel) min: 0.11% max: 2.46% x̄: 0.90% x̃: 0.56%
95% mean confidence interval for cycles value: -13.72 -1.55
95% mean confidence interval for cycles %-change: -1.01% -0.31%
Cycles are helped.
Tiger Lake
Instructions in all programs: 160902191 -> 160899554 (-0.0%)
SENDs in all programs: 6812435 -> 6812435 (+0.0%)
Loops in all programs: 38225 -> 38225 (+0.0%)
Cycles in all programs: 7428581420 -> 7428555881 (-0.0%)
Spills in all programs: 192582 -> 192582 (+0.0%)
Fills in all programs: 304539 -> 304539 (+0.0%)
A lot of fragment shaders in Shadow of the Tomb Raider were helped, and
a bunch of vertex shaders in Octopath Traveler were hurt.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10012>
It seems worth the small amount of damage to give an extra cushion of
not having to debug problems later.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
All Intel platforms had similar results. (Tiger Lake shown)
total instructions in shared programs: 21043197 -> 21043359 (<.01%)
instructions in affected programs: 4409 -> 4571 (3.67%)
helped: 0
HURT: 25
HURT stats (abs) min: 1 max: 16 x̄: 6.48 x̃: 5
HURT stats (rel) min: 0.39% max: 15.38% x̄: 4.59% x̃: 4.40%
95% mean confidence interval for instructions value: 4.37 8.59
95% mean confidence interval for instructions %-change: 2.93% 6.26%
Instructions are HURT.
total cycles in shared programs: 856175986 -> 856176921 (<.01%)
cycles in affected programs: 58908 -> 59843 (1.59%)
helped: 0
HURT: 25
HURT stats (abs) min: 7 max: 70 x̄: 37.40 x̃: 38
HURT stats (rel) min: 0.27% max: 5.63% x̄: 1.87% x̃: 1.39%
95% mean confidence interval for cycles value: 31.11 43.69
95% mean confidence interval for cycles %-change: 1.35% 2.39%
Cycles are HURT.
No fossil-db changes on any Intel platform.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10012>
When most of these patterns were created, we believed, incorrectly, that
fsat(NaN) was NaN. We have since realized that fsat(NaN) is zero.
Originally, this changed the patterns to use is_a_number. This didn't
help any shaders, so it's easier to just drop the optimizations.
This commit crossed paths with 4c3ad4d065 ("nir/algebraic: mark more
optimization with fsat(NaN) as inexact") and bc123c396a
("nir/algebraic: mark some optimizations with fsat(NaN) as inexact").
Given that these don't impact very many shaders, it seems safer to just
remove them.
As discussed in
https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8716, I tried
modifying these patterns to use !(b cmp a). Unfortunately, on Intel
GPUs, the results were much worse than just removing the patterns
altogether.
Some other related patterns will be addressed in later commits.
There are still a number of patterns that use the identity fsat(1-X) ==
1 - fsat(X). If X is NaN, the former is zero while the latter is 1.0.
I haven't evaluted these patterns yet. If changes are needed in these
patterns, it should be a separate commit anyway.
v2: Replace arrow `=>` with `->` in comments because the `=>` looks a
lot like `<=` comparison. Suggested by Rhys.
Fixes: 92b75c126b ("nir/algebraic: Replace checks that a value is between (or not) [0, 1]")
Fixes: a7f0c57673 ("nir/algebraic: Eliminate useless fsat() on operand of comparison w/value in (0, 1)")
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
All Intel hardware had similar results. (Ice Lake shown)
total instructions in shared programs: 20029060 -> 20029670 (<.01%)
instructions in affected programs: 69236 -> 69846 (0.88%)
helped: 0
HURT: 263
HURT stats (abs) min: 1 max: 20 x̄: 2.32 x̃: 1
HURT stats (rel) min: 0.30% max: 11.11% x̄: 1.35% x̃: 0.98%
95% mean confidence interval for instructions value: 1.86 2.78
95% mean confidence interval for instructions %-change: 1.18% 1.52%
Instructions are HURT.
total cycles in shared programs: 979821278 -> 979834425 (<.01%)
cycles in affected programs: 1476848 -> 1489995 (0.89%)
helped: 49
HURT: 204
helped stats (abs) min: 1 max: 812 x̄: 102.31 x̃: 20
helped stats (rel) min: 0.01% max: 21.43% x̄: 2.23% x̃: 0.52%
HURT stats (abs) min: 2 max: 2600 x̄: 89.02 x̃: 16
HURT stats (rel) min: 0.04% max: 27.27% x̄: 1.49% x̃: 0.72%
95% mean confidence interval for cycles value: 13.18 90.75
95% mean confidence interval for cycles %-change: 0.29% 1.25%
Cycles are HURT.
No fossil-db changes.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10012>
Many fragment shaders do a discard using relatively little information
but still put the discard fairly far down in the shader for no good
reason. If the discard is moved higher up, we can possibly avoid doing
some or almost all of the work in the shader. When this lets us skip
texturing operations, it's an especially high win.
One of the biggest offenders here is DXVK. The D3D APIs have different
rules for discards than OpenGL and Vulkan. One effective way (which is
what DXVK uses) to implement DX behavior on top of GL or Vulkan is to
wait until the very end of the shader to discard. This ends up in the
pessimal case where we always do all of the work before discarding.
This pass helps some DXVK shaders significantly.
v2 (Jason Ekstrand):
- Fix a couple of typos (Grazvydas, Ian)
- Use the new nir_instr_move helper
- Find all movable discards before moving anything so we don't
accidentally re-order anything and break dependencies
v3 (Pierre-Eric): remove the call to nir_opt_conditional_discard based
on Daniel Schürmann comment.
v4 (Pierre-Eric):
- handle demote intrinsics and drop derivatives_safe_after_discard
- add early return if discards/demotes aren't used
v5 (Pierre-Eric):
- use pass_flags instead of instr set (Daniel Schürmann)
v6 (Daniel Schürmann):
- cleanup and fix pass_flags handling
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10522>
Removes an instruction from one place and inserts it at another while
working around a weird cursor corner-case.
v2: change return value to bool (Daniel Schürmann)
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com> (v1)
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10522>
It's perfectly legal to declare multiple SSBOs that point to the same
binding/descriptor_set with different access mask. Currently, it will
always get the first one in the list that matches binding/desc_set
regardless of the access mask, but other variables might have different
access mask.
Fix this by being conservative if another variable uses the same
binding/desc_set because we can't get it reliably without adding
a new field to vulkan_resource_index.
This fixes rendering issues in Resident Evil Village with vkd3d-proton.
This bug has been uncovered by ("spirv: Don't remove variables used by
resource indexing intrinsics") because variables are no longer removed
No fossils-db changes.
Cc: 21.1 mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10692>
We don't want to have to deal with vector phis in freedreno, because
vectors are always split/unsplit around vectorized instructions anyways,
and the stated reason for not scalarising them (it hurting coalescing)
won't apply to us because we won't be using nir_from_ssa. Add this
option so that we don't have to do the equivalent thing while
translating from NIR.
Reviewed-by: Rob Clark <robdclark@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10809>
Backends that don't handle IO component precision can pack more varyings
into one slot if the linker ignores the precision. If the IO is vectorized
then this can save IO instructions.
Related: 165a69d2f7
nir: handle mediump varyings in varying compaction helpers
Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10722>
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>
Otherwise the lowering pass might try to lower any other load from
a deref if its data.location value happens to be zero.
Fixes: 418c4c0d7d
compiler/nir: extend lower_fragcoord_wtrans to support VARYING_SLOT_POS
Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10577>
From the ARB_enhanced_layouts spec:
"As with input layout qualifiers, all shaders except compute shaders
allow *location* layout qualifiers on output variable declarations,
output block declarations, and output block member declarations. Of
these, variables and block members (but not blocks) additionally
allow the *component* layout qualifier."
We previously had compile tests in piglit to make sure this was not a
compile error but no execution tests.
Fixes: d99a040bbf ("i965: enable ARB_enhanced_layouts for gen8+")
Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10763>
These intrinsics represent what the hardware can actually do.
Lowering our shaders to use these intrinsics will allow us to
deal with mapping the classic VS, TES, GS (and the future MS)
stages to the hardware capabilities using NIR, which makes our
backend compilers simpler.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10740>
The lowered NIR code of NGG VS shaders uses this intrinsic
when the VS has to export the primitive ID.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10740>
These allow us to generate slightly better code in some cases,
eg. multiplications in ACO.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10740>
These intrinsics will be used when lowering NGG shaders, including
currently supported stages like VS, TES, GS and also by mesh shaders
in the future.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10740>
Fixes issues with upcoming CTS test testing empty structs.
v2: decorate with UNUSED as only used in assert (Timothy)
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10681>
This pass was originally developed for Panfrost, where it passes the
relevant dEQP tests. Upstreaming so it can be extended and then shared
with:
* Asahi, for blending
* Zink, for logic ops
* Lavapipe, for advanced blending
Note that using this with MRT in a fragment shader (as non-panfrost
drivers will) has not yet been tested. Logic ops with integer
framebuffers are probably todo. It's been enough for Panfrost, will
suffice for ES2 on Asahi, and provides an upstream base for kusma's work
on advanced blending, so overall the merge is a net benefit.
v2: Remove bogus assert that the format layout is PLAIN. We need to
render R11G11B10, which Mesa reports as layout OTHER. The code is still
correct.
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com> [v1]
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10601>
These are equivalent to the 32bit opcodes if there are no more efficient
24bit opcodes available, but inputs are guaranteed to already be 24bit,
so the 24bit opcodes can be used instead if they exist and are efficient.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10549>
Move it out of the "cs" sub-struct, since the bit can be used for
other shader stages in the future.
This also removes a subtle issue in spirv_to_nir:
info.cs.shared_memory_explicit_layout was used without checking for
the CS shader stage. It ended up being "harmless" since the effects
also depended on presence of shared variables.
Fixes: 5de6c5973a ("spirv: Implement SPV_KHR_workgroup_memory_explicit_layout")
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10529>
The V3D hardware allows us to pack multiple workgroups together to avoid
wasting execution lanes in shader cores.
For example, if we dispatch 16 workgroups with a local size of 1 element, we
can pack all 16 workgroups in a single 16-wide dispatch where each lane
executes a different workgroup, instead of 16 1-wide dispatches.
When we do this, we don't have a uniform workgroup id any more.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10541>
The mode one was used before 0bc5a829dd ("nir: Remove shared support from
lower_io").
The others were used before 5f7c7c9a7f ("nir: add src and dest types
to all IO loads and stores for mediump").
All conditions now are always true, so drop them.
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10533>
For unsigned comparisons with zero these ops can be eliminated.
v2: Add comparison optimizations with -1 (Rhys Perry)
Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Reviewed-by: Eric Anholt <eric@anholt.net> (v1)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10583>
Used to split up the fsin/fcos lowering for AGX between NIR and the
backend, to permit algebraic optimizations without polluting NIR with
too many hardware details. The backend NIR lowering produces an
fmul/ffma of the input so we can optimize code like sin(2*x).
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10582>
When we originally added patch variable handling to spirv_to_nir, we
were splitting I/O block variables in spirv_to_nir, so we weren't
guaranteed to have a nir_variable early enough in processing.
Since b0c643d8f5 ("spirv: Use NIR per-member splitting"), we've been
using NIR per-member splitting where we have a nir_variable which has
a separate nir_variable_data per member. With this, we can drop
vtn_variable::patch and use the patch boolean on the nir_variable
instead.
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/10469>
This fixes linking failures with new GL45 linkage tests, no
regressions spotted on existing tests.
v2: add spec reference (Samuel)
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10373>
For non-CL, intrinsic access isn't set, because the image type doesn't
have access qualifier. Instead, the access qualifier is set on the variable.
So, add a mode to this pass which can chase back to the variable in addition
to the intrinsic access. Also, update the variable type and the deref chain
types so everything is consistent, that the tex is accessing a sampler. Note
we can't do this for CL, because void-typed samplers don't exist.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10356>
Fixes crash in
dEQP-GLES31.functional.shaders.framebuffer_fetch.basic.last_frag_data
when using this pass.
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10411>
To vectorize to vec8/16 or vec4 (without vec3), we can't incrementally add
components to a load/store. This patch loops vectorization so that two new
vec2/4/8 operations can be combined into a larger operation.
fossil-db (GFX10.3):
Totals from 22 (0.02% of 139391) affected shaders:
SpillVGPRs: 1749 -> 1771 (+1.26%)
CodeSize: 901212 -> 892532 (-0.96%); split: -1.19%, +0.22%
Scratch: 178176 -> 184320 (+3.45%)
Instrs: 159358 -> 158027 (-0.84%); split: -0.99%, +0.16%
Cycles: 37046772 -> 36738544 (-0.83%); split: -1.00%, +0.17%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10384>
This fixes the
dEQP-GLES3.functional.shaders.invariance.{low,medium,high}p.loop_4 tests when
run in a VM with virgl on a host with iris. virgl mangles the GLSL shaders and
emits shader code for the host driver that contains vec4 arrays. As such, the
test did not fail when running directly on the host.
The test also did not fail if the host was using i965. Disabling
PIPE_CAP_GLSL_OPTIMIZE_CONSERVATIVELY in iris was sufficient to work around it,
so I believe that i965 didn't show the problem because after arrays were split
by optimize_split_arrays(), even though the invariant/precise qualifiers were
lost, do_common_optimization() would be called again and thus
propagate_invariance() would propagate the qualifiers to the new variables
produced by optimize_split_arrays().
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10292>
Doing so allow you to easily tell what the pass did using the existing
infrastructure in the OPT macro.
Reviewed-by: Timothy Arceri <tarceri@itsqueeze.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10292>
This was renamed when I was in high school. I remember updating the
Midgard compiler while sitting in AP Physics.
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/10296>
Some hardware doesn't have a way to check if invocation was demoted,
in such case we have to track it ourselves.
OpIsHelperInvocationEXT is specified as:
"An invocation is currently a helper invocation if it was originally
invoked as a helper invocation or if it has been demoted to a helper
invocation by OpDemoteToHelperInvocationEXT."
Therefore we:
- Set gl_IsHelperInvocationEXT = gl_HelperInvocation
- Add "gl_IsHelperInvocationEXT = true" right before each demote
- Add "gl_IsHelperInvocationEXT = gl_IsHelperInvocationEXT || condition"
right before each demote_if
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Reviewed-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9460>
If a has been lowered to float16 here, then we end up trying to
construct a vector of mixed precision, which the validator asserts
about.
So let's make sure we use the same type for all arguments.
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10201>
Rhys Perry
Timur Kristóf
Caio Marcelo de Oliveira Filho
Alyssa Rosenzweig
Bas Nieuwenhuizen
Alyssa Rosenzweig
Hoe Hao Cheng
Mike Blumenkrantz
Timothy Arceri
Ian Romanick
Alyssa Rosenzweig
Jason Ekstrand
Jason Ekstrand
Samuel Pitoiset
Connor Abbott
Gert Wollny
Marcin Ślusarz
Tapani Pälli
Mauro Rossi
Jesse Natalie
Iago Toral Quiroga
Lionel Landwerlin
Matt Turner
Danylo Piliaiev
Erik Faye-Lund