If the barrier tries to apply to memory that we can't express, just
don't apply the memory portion of the barrier. Similarly, if it tries
to apply a global memory barrier at invocation level, upgrade it to
thread-group.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11670>
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>
This change moves the SRVs associated with read-only SSBOs to be emitted
before any other UAV. We do this because the validator expects resources
to be emitted in a specific order, as noted by `emit_module`.
Previously, we emitted SSBOs as SRVs (read-only) or UAVs (read-write)
after other UAVs.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10514>
Note that it's no longer sufficient to check for >=1 sampler/image
in a potential array, because unbounded arrays have 0 of them.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10298>
For GL, we've previously mostly ignored the binding property for sampler variables
during the shader compilation step. For CL, our image bindings were always 0-based as well.
Now, for Vulkan, we are going to be getting explicit bindings and need to emit DXIL that
respects those bindings. Since Vulkan can also have both split and combined images and samplers,
we now need to be smarter about recognizing when NIR is trying to use a "sampler" as *both* an
image and sampler (in deref mode, the same variable will be deref'd as both image and sampler).
That "being smarter" bit comes next, but first, let's prep GL for building correct root
signatures and binding the resources correctly.
Reviewed-By: Bill Kristiansen <billkris@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10298>
Creating resource handles upfront works well while we have fixed-size resource
counts, but once we start talking about bindless, having arrays or even sets
of handles becomes prohibitive. It also precludes dynamic indexing for textures.
Instead, rely on the load_vulkan_descriptor instruction for UBO/SSBO, and undo
nir_lower_samplers so we continue to have deref chains for image/sampler accesses.
Then, emit handles at the end of a deref chain - the chain should only have
array offsets, so once we get to a type that's not an array anymore, we can
emit the handle.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10288>
Previously UBOs only supported static indices, and SSBOs only
supported dynamic indices. UBO support for descriptors was added
as an alternative to static indices, but the logic for detecting
descriptors to SSBOs couldn't just differentiate on constants vs not.
Add a helper which can differentiate cleanly across the board and
handle pre-created handles from descriptors, or static/dynamic raw
indices.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10149>
Instead of doing all of the handle logic in the descriptor load, split
it so that the resource index is actually computed during resource_index
processing, and it's converted to a handle during the load_descriptor.
At the same time, add SSBO handling and dynamic indexing handling.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10149>
The resources need to be emitted in a particular order, so CBVs
have to be emitted first and can't be emitted as we iterate through
instructions.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10149>
MakeDouble is pretty straightforward, but SplitDouble is interesting
since it returns a unique 2-element struct.
Reviewed-by: Enrico Galli <enrico.galli@intel.com>
Reviewed-by: Michael Tang <tangm@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10063>
Some nir lowerings might need to know if txs is supported by
the backend.
Signed-off-by: Christian Gmeiner <christian.gmeiner@gmail.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8898>
It is currently a bitset on top of a uint64_t but there are already
more than 64 values. Change to use BITSET to cover all the
SYSTEM_VALUE_MAX bits.
Cc: mesa-stable
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Acked-by: Jesse Natalie <jenatali@microsoft.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Acked-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Acked-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8585>
This patch also replaces lower_negate with lower_ineg / lower_fneg.
The fneg semantics have been clarified as of Version 1.5, Revision 1
of the SPIR-V specification, which means that the previous lowering
to fsub is not a viable solution anymore, and is replaced with
lowering to fmul(x, -1.0).
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6597>
In DXIL, the FMA instruction only supports 64-bit operations. However,
back when we implemented support for this, there were only a single
switch for lowering all ffma instructions, so we couldn't easily use it.
But now that there's separate flags to lower ffma on 16, 32 and 64 bit,
we can lower 16 and 32 bit ffmas, and leave 64 bit ffmas alone.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8349>
When I originally added the FFMA opcode here, I added the FMAD opcode
instead of the FMA opcode. The reason for this is that it works on
32-bit values as well, so that seemed like a better fit.
But that's not correct, as the FMA opcode isn't a fused operation, so
let's correct the opcode.
This isn't currently in use, because we currently lower away all ffma
opcodes on the NIR level, but that's about to change.
While we're at it, let's also update the opcode name to match the DXIL
documentation.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8349>
Most of these are adding 'static', for functions that are local
to a translation unit but weren't declared static.
There's one instance of a missing include for bringing the prototype
into the translation unit, one function missing a return type (default-int),
and one which added inline to avoid it being considered unused in some sources.
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7780>
This wasn't implemented yet, because we hadn't encountered it yet. But
now it seems we can trigger this, thanks to the nv_copy_depth_to_color
piglit tests.
This makes the test go from crash to fail, which isn't perfect, but it's
better than nothing.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7855>
This adds a standalone library which can convert through the pipeline of
OpenCL C -> SPIR -> SPIR-V -> NIR -> DXIL. It can add in the libclc
implementations of various library functions in the NIR phase, and
also massages the NIR to shift it more towards graphics-style compute.
This is leveraged by the out-of-tree OpenCLOn12 runtime
(https://github.com/microsoft/OpenCLOn12).
This is the combination of a lot of commits from our development branch,
containing code by several authors.
Co-authored-by: Boris Brezillon <boris.brezillon@collabora.com>
Co-authored-by: Daniel Stone <daniels@collabora.com>
Co-authored-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7565>
This fixes a regression that happened after rebasing on master, where we
end up not writing all components of the clip-distance array, which the
DXIL validation code in the D3D12 runtime treats as an error.
To ensure we don't end up overwriting a previous wrire, enable
nir_shader_compiler_options::lower_all_io_to_temps as well.
Reviewed-By: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7477>
Here's the code to emit DXIL code from NIR. It's big and bulky as-is,
and it needs to be split up a bit.
This is the combination of a lot of commits from our development branch,
containing code by several authors.
Co-authored-by: Bill Kristiansen <billkris@microsoft.com>
Co-authored-by: Boris Brezillon <boris.brezillon@collabora.com>
Co-authored-by: Daniel Stone <daniels@collabora.com>
Co-authored-by: Gert Wollny <gert.wollny@collabora.com>
Co-authored-by: Jesse Natalie <jenatali@microsoft.com>
Co-authored-by: Louis-Francis Ratté-Boulianne <lfrb@collabora.com>
Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7477>
Michael Tang
Enrico Galli
Jesse Natalie
Caio Marcelo de Oliveira Filho
Rhys Perry
Bas Nieuwenhuizen
Christian Gmeiner
Daniel Schürmann
Erik Faye-Lund
Adam Jackson