Commit Graph

480 Commits

Author SHA1 Message Date
Jesse Natalie ee905aa3a0 nir/vtn: CL SPIR-V callers should specify address modes
Instead of inferring the address mode from the environment, allows
callers to override to suit their needs.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6330>
2020-08-17 14:36:18 +00:00
Karol Herbst 217def3ac6 spirv: parse unstructured CFG
v2 (Boris Brezillon): handle functions with return values
v3: call structurizer
v4: entire rewrite
v5: fix handling of already visited default branches
v2 (Jason Ekstrand): Stop walking hash tables

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/2401>
2020-08-14 20:35:37 +00:00
Rhys Perry 389c95a889 spirv: set ACCESS_COHERENT for ssbo/global/image atomic load/store
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6063>
2020-08-05 09:45:54 +00:00
Caio Marcelo de Oliveira Filho b98dd70489 spirv: Propagate explicit layout only in types that need it
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5511>
2020-08-04 07:53:37 -07:00
Jason Ekstrand 4816f6f8d8 spirv: Do more complex unwrapping in get_nir_type
The OpenGL flavor of SPIR-V allows for samplers inside structs.  This
means that our simple array-of-array handling isn't sufficient and we
need something substantially more complex for generating NIR types.

Fixes: 14a12b771d "spirv: Rework our handling of images and samplers"
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6065>
2020-08-04 03:05:34 +00:00
Caio Marcelo de Oliveira Filho 12dd5455f4 spirv: Handle most execution modes earlier
For convenience in e68871f6a4 ("spirv: Handle constants and types
before execution modes") we moved all execution mode parsing after the
constants and types, so that those using OpExecutionModeId could be
handled together.

Later in 84781e1f1d ("spirv/nir: keep track of SPV_KHR_float_controls
execution modes") we had to parse certain non-ID execution modes
before handling constants.

Instead of handling just the float controls related execution modes
early, handle all modes that don't need an ID.  This is a more
"natural" split and will allow other type handling to rely on
execution mode in the future.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6062>
2020-07-27 17:10:23 +00:00
Jason Ekstrand 14a12b771d spirv: Rework our handling of images and samplers
Previously, objects of type OpTypeImage or OpTypeSampler were treated as
vtn_pointers and objects of type OpTypeSampledImage were a special-use
vtn_sampled_image struct.  This commit changes that so that all of those
objects are stored in vtn_ssa_values.  Each of images, samplers, and
sampled images, are stored as a scalar or vector nir_ssa_def whose
components are NIR deref values.  We now use vtn_type_get_nir_type to
re-resolve those as-needed into GLSL sampler types for NIR.

This simplification has a number of benefits:

 1. We can git rid of the rest of our special-cases for handling images
    and samplers in function arguments.  Now that they're treated as
    structs at the glsl_type level, the generic paths can handle images
    and samplers.

 2. We can now construct composite values containing images and samplers
    internally.  It's unclear from the SPIR-V spec whether or not this
    is allowed and it's not a pattern that GLSLang currently generates
    thanks to GLSL rules.  However, if we do start seeing SPIR-V that
    contains such composites, we should now be able to handle it.

 3. SPIR-V OpNull and OpUndef instructions can now create samplers,
    images, and sampled images.  The NIR generated won't likely be fully
    valid but, given a NIR pass to do something sensible, it should be a
    thing we can compile.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 9e3213ad30 spirv: Add a helper for getting the NIR type of a vtn_type
There are a few cases, atomic counters being one example, where the type
used by vtn_ssa_value is not the same as the type we want NIR to use in
derefs and variables.  To solve this, we add a helper which converts
between the types for us.  In the next commit, we'll be adding another
major user of this: images and samplers.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 051f8d3d1c spirv: Give atomic counters their own variable mode
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 8a7932b095 spirv: Drop the sampled boolean from vtn_type
It was set but never used.  We always check the glsl_type instead.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand d0be2fed4e spirv: Add better checks for SSA value types
Primarily, we check for two things:

 1. That we only ever add SSA values via vtn_push_ssa_value and
    vtn_copy_value.

 2. That the type of the SSA value matches the SPIR-V destination type.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 47ebb50cab spirv: Hand-roll fewer vtn_ssa_value creations
Previously, we created our vtn_ssa_value in _vtn_variable_load_store
manually as we did the recursive load/store.  Instead, we now create the
SSA value before calling into the recursive function.  This is a tiny
bit less efficient but it removes a case of hand-rolling vtn_ssa_value
creation.  For symmetry, we make _vtn_block_load_store assume the value
is already created.  Finally, we remove a trivial hand-rolled case in
vtn_composite_extract.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 085ade4125 spirv: Simplify vtn_ssa_value creation
For three different functions which create vtn_ssa_values, we had three
completely different implementations.  This unifies them all to roughly
the same algorithm.  While we're at it, we take advantage of the
nir_build_imm helper to avoid some extra code in vtn_const_ssa_value.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand dbb4a24413 spirv: Refactor vtn_push_ssa
We rename it to vtn_push_ssa_value, move it to spirv_to_nir, and remove
the unnecessary type parameter.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 37ab323480 spirv: Add a vtn_get_nir_ssa helper
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand d8eb6f2499 spirv: Add a vtn_push_nir_ssa helper
This makes it easy to write a simple NIR SSA value

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand aaf1f34215 spirv: Rename push_value_pointer to push_pointer
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand ac1e6d5a46 spirv: Add a helpers for getting types of values
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:43:21 -05:00
Jason Ekstrand 62c53ad20b spirv: Fix indentation in vtn_handle_ptr
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:41:54 -05:00
Jason Ekstrand af81486a8c spirv: Simplify our handling of NonUniform
The original implementation of SPV_EXT_descriptor_indexing was extremely
paranoid about the NonUniform qualifier, trying to fetch it from every
possible location and propagate it through access chains etc.  However,
the Vulkan spec is quite nice to us on this and has very strict rules
for where the NonUniform decoration has to be placed.  For image and
texture operations, we can search for the decoration on the spot when we
process the image or texture op.  For pointers, we continue putting it
on the pointer but we don't bother trying to do anything silly like
propagate it through casts.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5278>
2020-07-23 22:41:54 -05:00
Jason Ekstrand 84086b620e spirv: Add support for SPV_EXT_shader_atomic_float
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5992>
2020-07-21 05:01:34 +00:00
Rhys Perry fac813dc61 spirv: don't split memory barriers
If the SPIR-V had a shared+image memory barrier, we would emit two NIR
barriers: a shared barrier and an image barrier.

Unlike a single barrier, two barriers allows transformations such as:

intrinsic image_deref_store (ssa_27, ssa_33, ssa_34, ssa_32, ssa_25) (1)
intrinsic memory_barrier_shared () ()
intrinsic memory_barrier_image () ()
intrinsic store_shared (ssa_35, ssa_24) (0, 1, 4, 0)
->
intrinsic memory_barrier_shared () ()
intrinsic store_shared (ssa_35, ssa_24) (0, 1, 4, 0)
intrinsic image_deref_store (ssa_27, ssa_33, ssa_34, ssa_32, ssa_25) (1)
intrinsic memory_barrier_image () ()

This commit fixes two dEQP-VK.memory_model.* CTS tests with ACO.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5951>
2020-07-20 12:05:16 +00:00
Jason Ekstrand 351b5137d7 spirv: Allow block-decorated struct types for constants
Whenever a struct type is decorated Block or BufferBlock we turn that
into a GLSL_TYPE_INTERFACE.  Since these decorations can end up random
places, we should allow them for constants.

Closes: #3252
Fixes: 9d0ae777dd "spirv: Use interface type for block and buffer..."
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5855>
2020-07-12 00:02:45 +00:00
Rhys Perry c344c083fc spirv: set variables to restrict by default
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5207>
2020-06-24 10:52:27 +00:00
Rob Clark f94ba1555d spirv: drop some dead code
This case is never hit, we don't have a nir intrinsic for this spirv
opcode.  And when we do, I'm not sure if it would be vectorized or not.
So best just to drop this case.

Signed-off-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5505>
2020-06-18 03:40:54 +00:00
Rob Clark f43a2cd1d9 spirv: atomic_counter_read_deref is not vectorized
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/3141
Signed-off-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5505>
2020-06-18 03:40:54 +00:00
Rob Clark 5b5b45ebf6 spriv: don't set num_components for non-vectorised intrinsics
Signed-off-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5371>
2020-06-16 02:48:18 +00:00
Samuel Pitoiset 9b6a8d1742 spirv: fix using OpSampledImage with OpUndef instead of OpType{Image,Sampler}
This seems valid per the SPIR-V spec to use OpSampledImage with
OpUndef instead of OpTypeImage or OpTypeSampler. When the image
operand is undefined, SPIRV->NIR emits an undef instruction that
can be removed later by the compiler.

This fixes shader compilation crashes with Red Dead Redemption II.

Cc: mesa-stable@lists.freedesktop.org>
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5230>
2020-06-15 17:02:53 +02:00
Caio Marcelo de Oliveira Filho e5bb4b1ee8 spirv: Memory semantics is optional for OpControlBarrier
Fixes: 3ed2123d77 ("spirv: Use scoped barriers for SpvOpControlBarrier")
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5365>
2020-06-08 15:49:24 +00:00
Boris Brezillon 3ed2123d77 spirv: Use scoped barriers for SpvOpControlBarrier
If use_scoped_barrier is set to true, we don't have to split the control
and memory barriers.

Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4900>
2020-06-03 07:39:52 +00:00
Boris Brezillon 345b5847b4 nir: Replace the scoped_memory barrier by a scoped_barrier
SPIRV OpControlBarrier can have both a memory and a control barrier
which some hardware can handle with a single instruction. Let's
turn the scoped_memory_barrier into a scoped barrier which can embed
both barrier types. Note that control-only or memory-only barriers can
be supported through this new intrinsic by passing NIR_SCOPE_NONE to the
unused barrier type.

Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Suggested-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4900>
2020-06-03 07:39:52 +00:00
Boris Brezillon 94438a64bf spirv: Split the vtn_emit_scoped_memory_barrier() logic
We are about to add support for scoped control+memory barriers. Let's
move the convert from SPIRV to NIR enums logic in helpers so we can
easily re-use them.

Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4900>
2020-06-03 07:39:52 +00:00
Timothy Arceri 04dbf709ed nir: add callback to nir_remove_dead_variables()
This allows us to do API specific checks before removing variable
without filling nir_remove_dead_variables() with API specific code.

In the following patches we will use this to support the removal
of dead uniforms in GLSL.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4797>
2020-06-03 02:22:23 +00:00
Samuel Pitoiset 10c4a7cf59 spirv,radv,anv: implement no-op VK_GOOGLE_user_type
This extension only allows HLSL shader compilers to optionally embed
unambiguous type information which can be safely ignored by the driver.

This fixes a crash with the recent Vulkan backend of Path Of Exile
(it uses the extension without checking if it's supported).

Cc: <mesa-stable@lists.freedesktop.org>
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Tested-by: Edmondo Tommasina <edmondo.tommasina@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5237>
2020-05-28 17:30:24 +02:00
Samuel Pitoiset 41dc3ce449 spirv: add support for bias/lod with OpImageGather
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5147>
2020-05-25 08:51:10 +02:00
Samuel Pitoiset dd39bf52b0 spirv: add SpvCapabilityImageGatherBiasLodAMD
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5147>
2020-05-25 08:51:10 +02:00
Samuel Pitoiset 37c88c670f spirv: add ReadClockKHR support with device scope
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5117>
2020-05-24 20:37:50 +02:00
Samuel Pitoiset 844d561c58 spirv: handle OpCopyObject correctly with any types
This implements OpCopyObject as a blind copy and propagates the
access mask properly even if the source object type isn't a SSA
value.

This fixes some recent dEQP-VK.descriptor_indexing.* failures
since CTS changed and now apply nonUniformEXT after constructing
a combined image/sampler.

Original patch is from Jason Ekstrand.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4909>
2020-05-15 19:18:53 +00:00
Jason Ekstrand f4addfdde3 spirv: Use nir_const_value for spec constants
When we originally wrote spirv_to_nir we didn't have a good scalar value
union to handily use so we rolled our own thing for spec constants.  Now
that we have nir_const_value, we can use that and simplify a bunch of
the spec constant logic.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Acked-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4675>
2020-04-24 09:23:59 +00:00
Danylo Piliaiev 66229aa169 spirv: Expand workaround for OpControlBarrier on old GLSLang
In SPIRV of compute shader in Aztec Ruins benchmark there is:

OpControlBarrier %uint_1 %uint_1 %uint_0
// ControlBarrier(Device, Device, rdcspv::MemorySemantics(0));

which is an incorrect translation of glsl barrier().

GLSLang, prior to c3f1cdfa, emitted the OpControlBarrier with
Device instead of Workgroup for execution scope.

2365520c covers similar case but isn't applied when execution_scope
is SpvScopeDevice.

Cc: <mesa-stable@lists.freedesktop.org>
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/2742
Signed-off-by: Danylo Piliaiev <danylo.piliaiev@globallogic.com>
Tested-by: Rafael Antognolli <rafael.antognolli@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4660>
2020-04-22 08:46:12 +00:00
Jason Ekstrand f5deed138a spirv,nir: Move the SPIR-V vector insert code to NIR
This also makes spirv_to_nir a bit simpler because the new
nir_vector_insert helper automatically handles a constant component
selector like nir_vector_extract does.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4495>
2020-04-17 19:21:44 +00:00
Jason Ekstrand feca439697 spirv: Call nir_builder directly for vector_extract
The nir_builder helper already handles checking if the component
selector is an immediate and returns an undef in the OOB case.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4495>
2020-04-17 19:21:44 +00:00
Jason Ekstrand 4b160c6776 spirv: Error if OpCompositeInsert/Extract has OOB indices
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4495>
2020-04-17 19:21:44 +00:00
Jason Ekstrand c478f8ad6c spirv,nir: Add a better vector_insert
The old one in spirv_to_nir was besel'ing the whole vector for every
component.  If we think about this as a vector operation, we can do it
way more efficiently.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4495>
2020-04-17 19:21:44 +00:00
Jason Ekstrand 380bf556bf spirv: Handle OOB vector extract operations
We use vtn_vector_extract to handle vector component level derefs.  This
makes us gracefully handle the case where your vector component is OOB
and give you an undef.  The SPIR-V working group is still working out
whether or not this is technically legal but it's very little code for
us to handle it so we may as well.

Cc: mesa-stable@lists.freedesktop.org
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4495>
2020-04-17 19:21:44 +00:00
Jason Ekstrand d94e464a9f spirv: Make vtn_function a vtn_cf_node
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3820>
2020-04-03 20:54:00 +00:00
Jason Ekstrand 68f325b256 Revert "spirv: Implement OpCopyObject and OpCopyLogical as blind copies"
This reverts commit 7a53e67816.
2020-04-01 12:40:34 -05:00
Jason Ekstrand 7a53e67816 spirv: Implement OpCopyObject and OpCopyLogical as blind copies
Because the types etc. are required to logically match, we can just
copy-propagate the guts of the vtn_value.  This was causing issues with
some new CTS tests that are doing an OpCopyObject of a sampler which is
a special-cased type in spirv_to_nir.  Of course, this is only a partial
solution.  Ideally, we've got a bit of work to do to make all the
composite stuff able to handle all types including images, sampler, and
combined image/samplers but this gets some CTS tests passing.

Cc: mesa-stable@lists.freedesktop.org
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Acked-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4375>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4375>
2020-03-31 17:55:30 +00:00
Rhys Perry 8b361df9cf spirv: fix memory_barrier_tcs_patch emission
Shouldn't affect any driver, since all currently implement
memory_barrier_tcs_patch as a no-op. It also looks like optimizations are
fine

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4003>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4003>
2020-03-03 11:49:40 +00:00
Rhys Perry 6d839addf9 spirv: improve creation of memory_barrier
It shouldn't check for atomic counters or return in case we also need to
create a TCS output barrier.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4003>
2020-03-03 11:49:40 +00:00
Caio Marcelo de Oliveira Filho 956e4b2d37 nir, intel: Move use_scoped_memory_barrier to nir_options
This option will be used later by GLSL, so move to a common struct.

Because nir_options is filled in the compiler instead of the Vulkan
driver, fix that up.  GLSL will ignore that for now.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3913>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3913>
2020-02-24 19:12:11 +00:00
Eric Anholt d37c6ebd3c spirv_to_nir: Reuse glsl_sampler_dim_coordinate_components().
We just needed to move the SUBPASS_MS case in, and the rest of the cases
match up.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3728>
2020-02-24 18:25:02 +00:00
Caio Marcelo de Oliveira Filho f58b384fbe spirv: Be consistent when checking for Shader/Kernel
Use == and != instead of the ordered comparisons.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3911>
2020-02-21 13:09:44 -08:00
Arcady Goldmints-Orlov 5f3cbbd958 spirv: Remove outdated SPIR-V decoration warnings
spirv_to_nir warns if it encounters XFB decorations and errors if
it encounters a Stream decoration with value other than 0, despite
the fact that these decorations are in fact handled correctly.

Fixes dEQP-VK.transform_feedback.simple.query_1_*
Fixes: cd4a14be06 "spirv: Handle XFB variable decorations"

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3910>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3910>
2020-02-21 20:34:03 +00:00
Arcady Goldmints-Orlov e9f83185a2 Rename nir_lower_constant_initializers to nir_lower_variable_initalizers
This is naming is more clear as nir_variables can be initializes not
just with a nir_constant but with a pointer to another nir_variable.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3047>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3047>
2020-02-12 15:41:49 +00:00
Eric Anholt 8d07d66180 glsl,nir: Switch the enum representing shader image formats to PIPE_FORMAT.
This means you can directly use format utils on it without having to have
your own GL enum to number-of-components switch statement (or whatever) in
your vulkan backend.

Thanks to imirkin for fixing up the nouveau driver (and a couple of core
details).

This fixes the computed qualifiers for EXT_shader_image_load_store's
non-integer sizeNxM qualifiers, which we don't have tests for.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com> (v3d)
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3355>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3355>
2020-02-05 10:31:14 -08:00
Samuel Pitoiset 531a26d5aa spirv: implement SPV_AMD_shader_explicit_vertex_parameter
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset 5c053cc6ec spirv: add support for SpvDecorationExplicitInterpAMD
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset 76a34f5d3f spirv: add support for SpvOpFragment{Mask}FetchAMD operations
nir_tex_src_ms_index is re-used for the fragment index with
nir_texop_fragment_fetch to avoid introducing a new texture source type.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3304>
2020-01-23 10:48:02 +00:00
Samuel Pitoiset dea29b3818 spirv: add SpvCapabilityFragmentMaskAMD
This new capability is for SPV_AMD_shader_fragment_mask.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3304>
2020-01-23 10:48:02 +00:00
Ian Romanick 4fcddb55f2 spirv: Add support for IntegerFunctions2INTEL capability
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick aa56934e2a spirv: Silence a bunch of unused parameter warnings
The change to get_uniform_nir_atomic_op make it look like the other
get_*_nir_atomic_op functions.  The rest just add UNUSED or ASSERTED
to parameters required for some of the interfaces.

src/compiler/spirv/spirv_to_nir.c: In function ‘struct_member_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:673:47: warning: unused parameter ‘val’ [-Wunused-parameter]
                             struct vtn_value *val, int member,
                                               ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘struct_member_matrix_stride_cb’:
src/compiler/spirv/spirv_to_nir.c:778:50: warning: unused parameter ‘val’ [-Wunused-parameter]
                                struct vtn_value *val, int member,
                                                  ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘type_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:805:61: warning: unused parameter ‘ctx’ [-Wunused-parameter]
                     const struct vtn_decoration *dec, void *ctx)
                                                             ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘spec_constant_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:1359:70: warning: unused parameter ‘v’ [-Wunused-parameter]
 spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
                                                                      ^
src/compiler/spirv/spirv_to_nir.c: In function ‘handle_workgroup_size_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:1407:43: warning: unused parameter ‘data’ [-Wunused-parameter]
                                     void *data)
                                           ^~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_function_call’:
src/compiler/spirv/spirv_to_nir.c:1806:55: warning: unused parameter ‘opcode’ [-Wunused-parameter]
 vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
                                                       ^~~~~~
src/compiler/spirv/spirv_to_nir.c:1807:54: warning: unused parameter ‘count’ [-Wunused-parameter]
                          const uint32_t *w, unsigned count)
                                                      ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘get_uniform_nir_atomic_op’:
src/compiler/spirv/spirv_to_nir.c:2548:47: warning: unused parameter ‘b’ [-Wunused-parameter]
 get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
                                               ^
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_atomics’:
src/compiler/spirv/spirv_to_nir.c:2633:48: warning: unused parameter ‘count’ [-Wunused-parameter]
                    const uint32_t *w, unsigned count)
                                                ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_barrier’:
src/compiler/spirv/spirv_to_nir.c:3197:48: warning: unused parameter ‘count’ [-Wunused-parameter]
                    const uint32_t *w, unsigned count)
                                                ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_execution_mode’:
src/compiler/spirv/spirv_to_nir.c:3618:68: warning: unused parameter ‘data’ [-Wunused-parameter]
                           const struct vtn_decoration *mode, void *data)
                                                                    ^~~~

Acked-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick 44471a76e9 nir/spirv: Translate SPIR-V to NIR for new INTEL_shader_integer_functions2 opcodes
v2: Rebase on 272e927d0e ("nir/spirv: initial handling of OpenCL.std
extension opcodes")

v3: Add missing SpvOpUCountTrailingZerosINTEL case to switch in
vtn_handle_body_instruction. Remove stray semicolon in
vtn_nir_alu_op_for_spirv_opcode. Use umin instead of umax for
SpvOpUCountTrailingZerosINTEL "lowering" in vtn_handle_alu.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Caio Marcelo de Oliveira Filho d8440a3d2f spirv: Handle PhysicalStorageBuffer in memory barriers
PhysicalStorageBuffer is lowered to nir_var_mem_global, and
SPIR-V 1.5rev1 in section "3.25. Memory Semantics <id>" says

    UniformMemory

    Apply the memory-ordering constraints to StorageBuffer,
    PhysicalStorageBuffer, or Uniform Storage Class memory.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3322>
2020-01-14 14:42:12 -08:00
Caio Marcelo de Oliveira Filho 1ec0d4fdff spirv: Drop EXT for PhysicalStorageBuffer symbols
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3322>
2020-01-14 14:42:12 -08:00
Jason Ekstrand e40b11bbcb nir: Rename nir_intrinsic_barrier to control_barrier
This is a more explicit name now that we don't want it to be doing any
memory barrier stuff for us.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3307>
2020-01-13 17:23:47 +00:00
Jason Ekstrand a4125b4d26 spirv: Add output memory semantics to OpControlBarrier in TCS
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3307>
2020-01-13 17:23:47 +00:00
Jason Ekstrand 2365520c9d spirv: Add a workaround for OpControlBarrier on old GLSLang
As per the Vulkan memory model, the proper translation of GLSL barrier()
is an OpControlBarrier with a scope of Workgroup and semantics of
Acquire, Release, and WorkgroupMemory.  Older versions of GLSLang gave
an OpControlBarrier with semantics of None so we need to patch it up on
those versions.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3307>
2020-01-13 17:23:47 +00:00
Samuel Pitoiset 1b808d208f spirv,nir: add new lod parameter to image_{load,store} intrinsics
SPV_AMD_shader_image_load_store_lod allows to use a lod parameter
with OpImageRead, OpImageWrite and OpImageSparseRead.

According to the specification, this parameter should be a 32-bit
integer. It is initialized to 0 when no lod parameter is found
during SPIR-V->NIR translation.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2020-01-09 07:58:33 +01:00
Samuel Pitoiset 37bfd854c7 spirv: add SpvCapabilityImageReadWriteLodAMD
New SPIR-V capability for SPV_AMD_shader_image_load_store_lod.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2020-01-09 07:58:33 +01:00
Rob Clark a8ec4082a4 nir+vtn: vec8+vec16 support
This introduces new vec8 and vec16 instructions (which are the only
instructions taking more than 4 sources), in order to construct 8 and 16
component vectors.

In order to avoid fixing up the non-autogenerated nir_build_alu() sites
and making them pass 16 src args for the benefit of the two instructions
that take more than 4 srcs (ie vec8 and vec16), nir_build_alu() is has
nir_build_alu_tail() split out and re-used by nir_build_alu2() (which is
used for the > 4 src args case).

v2 (Karol Herbst):
  use nir_build_alu2 for vec8 and vec16
  use python's array multiplication syntax
  add nir_op_vec helper
  simplify nir_vec
  nir_build_alu_tail -> nir_builder_alu_instr_finish_and_insert
  use nir_build_alu for opcodes with <= 4 sources
v3 (Karol Herbst):
  fix nir_serialize
v4 (Dave Airlie):
  fix serialization of glsl_type
  handle vec8/16 in lowering of bools
v5 (Karol Herbst):
  fix load store vectorizer

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Dave Airlie <airlied@redhat.com>
2019-12-21 11:00:17 +00:00
Caio Marcelo de Oliveira Filho 4fbc99c124 spirv: Implement SPV_KHR_non_semantic_info
Do nothing for OpExtInst from extended instruction sets that name
start with "NonSemantic.".

Since they can be used within the "preamble" to annotate global
decorations, also don't stop iterating when one of them is found.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3154>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3154>
2019-12-19 22:49:39 -08:00
Rob Clark ddb9701a3c spirv: add OpLifetime*
These are just hints so we can ignore them.

Signed-off-by: Rob Clark <robdclark@gmail.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Dave Airlie <airlied@redhat.com>
Signed-off-by: Karol Herbst <kherbst@redhat.com>
2019-12-11 23:54:39 +00:00
Karol Herbst 2402232c90 spirv: handle UniformConstant for OpenCL kernels
The caller is responsible for setting up the ubo_addr_format value as
contrary to shared and global, it's not controlled by the spirv.

Right now clovers implementation of CL constant memory uses a 24/8 bit format
to encode the buffer index and offset, but that code is dead as all backends
treat constants as global memory to workaround annoying issues within OpenCL.

Maybe that will change, maybe not. But just in case somebody wants to look at
it, add a toggle for this inside vtn.

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Dave Airlie <airlied@redhat.com>
2019-12-11 23:54:39 +00:00
Dave Airlie b62a925ad1 spirv: don't store 0 to cs.ptr_size for non kernel stages.
cs is a union so storing this there is wrong.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-11-22 04:37:25 +10:00
Jason Ekstrand 9e440b8d0b spirv: Sort out the mess that is sampled image
This commit makes two major changes.  First, we add a second case to
OpLoad for sampled images which constructs a vtn_sampled_image and
stashes that rather than stashing a pointer to the combined image
sampler like we do for bare samplers and images.  This should be more in
line with how SPIR-V is intended to work and hopefully doesn't cause any
weird problems.  The second is a rework of vtn_handle_texture to assume
that everything has an image but not everything has a sampler.  We also
add a vtn_fail_if for the case where a texture instructions require a
sampler but none is provided.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-11-09 15:29:01 +00:00
Jason Ekstrand 4f9688e571 spirv: Remove the type from sampled_image
We have types on all vtn_values at this point so there's no reason to
carry the redundant type information.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-11-09 15:29:01 +00:00
Caio Marcelo de Oliveira Filho 087ecd9ca5 spirv: Don't leak GS initialization to other stages
The stage specific fields of shader_info are in an union.  We've
likely been lucky that this value was either overwritten or ignored by
other stages.  The recent change in shader_info layout in commit
84a1a2578d ("compiler: pack shader_info from 160 bytes to 96 bytes")
made this issue visible.

Fixes: cf2257069c ("nir/spirv: Set a default number of invocations for geometry shaders")
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
2019-11-08 16:28:21 -08:00
Caio Marcelo de Oliveira Filho 9c3c206e71 spirv: Don't fail if multiple ordering semantics bits are set
Vulkan requires that only one bit for the ordering is set, but old
versions of GLSLang just set all the bits.  This was fixed as part of
c51287d744
but we can still find older versions (or shaders compiled with it)
around.

So instead of failing, emit a warning and fallback to the effective
result of any combination of multiple bits: AcquireRelease.

Closes: https://gitlab.freedesktop.org/mesa/mesa/issues/2018
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-10-29 14:53:46 -07:00
Caio Marcelo de Oliveira Filho d31f415ba0 spirv: Add helper to find args of Image Operands
Avoid keeping track of the idx and all possible image operands for
each operation.  Note for convenience we split up the handling of
ImageOperandsOffsetMask and ImageOperandsConstOffsetMask.

Suggested by Jason Ekstrand.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho c7d8fe2f0d spirv: Check that only one offset is defined as Image Operand
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho 129c85c28b spirv: Handle MakeTexelAvailable/Visible
Set the memory semantics and scope for later emitting the barrier.
Note the barrier emission code already exist in vtn_handle_image for
the Image atomics.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho c649e64edc spirv: Add option to emit scoped memory barriers
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho c022043102 spirv: Add SpvMemoryModelVulkan and related capabilities
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho 1bb191a0d1 spirv: Emit memory barriers for atomic operations
Add a helper to split the memory semantics into before and after the
operation, and use that result to emit memory barriers.

v2: Be more explicit about which bits we are keeping around when
    splitting memory semantics into a before and after.  For now
    we are ignoring Volatile.  (Jason)

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Caio Marcelo de Oliveira Filho d6992f996b spirv: Parse memory semantics for atomic operations
Including the right storage memory semantic based on the storage class
of the operation.  These will be used later to emit memory barriers.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-10-24 11:39:56 -07:00
Eduardo Lima Mitev f1d4fadf1b nir: Add new texop nir_texop_tex_prefetch
This is like nir_texop_tex, but signals that the sampling coordinates
are immutable during the shader stage, in a way that allows the HW
that supports pre-dispatching sampling operations to pre-fetch
the result prior to scheduling the shader stage.

This is introduced to support the feature in Freedreno. Adreno HW
from a4xx supports it.

A NIR pass introduced later in this series will detect sampling
operations that are eligible for pre-dispatch, and replace
nir_texop_tex by this new op, to tell the backend to enable
pre-fetch.

Signed-off-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
2019-10-18 21:11:54 +00:00
Jonathan Marek 39d7cb36ff spirv: set correct dest_type for texture query ops
Signed-off-by: Jonathan Marek <jonathan@marek.ca>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-10-15 08:42:22 -04:00
Caio Marcelo de Oliveira Filho f20cea0162 spirv: Implement SPV_KHR_shader_clock
We only have the subgroup variant in NIR (equivalent to clockARB), so
only support that for now.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
2019-10-07 09:12:12 -07:00
Caio Marcelo de Oliveira Filho d38e0a6326 spirv: Add missing break for capability handling
New added cases "stole" the previous break.

Fixes: 420ad0a1a3 ("spirv: check support for SPV_KHR_float_controls capabilities")
Reviewed-by: Eric Engestrom <eric.engestrom@intel.com>
2019-09-18 15:49:14 -07:00
Samuel Iglesias Gonsálvez f7d73db353 nir: add support for flushing to zero denorm constants
v2:
- Refactor conditions and shared function (Connor).
- Move code to nir_eval_const_opcode() (Connor).
- Don't flush to zero on fquantize2f16
  From Vulkan spec, VK_KHR_shader_float_controls section:

  "3) Do denorm and rounding mode controls apply to OpSpecConstantOp?

  RESOLVED: Yes, except when the opcode is OpQuantizeToF16."

v3:
- Fix bit size (Connor).
- Fix execution mode on nir_loop_analize (Connor).

v4:
- Adapt after API changes to nir_eval_const_opcode (Andres).

v5:
- Simplify constant_denorm_flush_to_zero (Caio).

v6:
- Adapt after API changes and to use the new constant
  constructors (Andres).
- Replace MAYBE_UNUSED with UNUSED as the first is going
  away (Andres).

v7:
- Adapt to newly added calls (Andres).
- Simplified the auxiliary to flush denorms to zero (Caio).
- Updated to renamed supported capabilities member (Andres).

Signed-off-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Signed-off-by: Andres Gomez <agomez@igalia.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com> [v4]
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-09-17 23:39:18 +03:00
Samuel Iglesias Gonsálvez 84781e1f1d spirv/nir: keep track of SPV_KHR_float_controls execution modes
v2:
- Add support for rounding modes for each floating point bit size.

v3:
- Commit e68871f6a4 ("spirv: Handle constants and types before
  execution modes") changed when the execution modes are handled,
  which affects the result of the floating point constants when the
  rounding mode is set in the execution mode. Moved the handling of
  the rounding modes before we handle the constants.

v4:
- Rename vtn_decoration "literals" to "operands" (Andres).
- Simplify execution mode parsing util function (Caio).
- Extend the comment about the timing of the handling of the rounding
  modes (Caio).

v5:
- Correct extension name (Caio).
- Rename shader info member (Andres).
- Rename float controls enum (Andres).

Signed-off-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Signed-off-by: Andres Gomez <agomez@igalia.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com> [v3]
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-09-17 23:39:18 +03:00
Samuel Iglesias Gonsálvez 420ad0a1a3 spirv: check support for SPV_KHR_float_controls capabilities
v2:
- Correct extension name (Caio).
- Rename supported capabilities member (Andres).

Signed-off-by: Samuel Iglesias Gonsálvez <siglesias@igalia.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com> [v1]
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-09-17 23:39:18 +03:00
Caio Marcelo de Oliveira Filho 9cf1bfcdd7 spirv: Handle ShaderLayer and ShaderViewportIndex capabilities
SPIR-V 1.5 incorported the SPV_EXT_shader_viewport_index_layer but
splitting into the two capabilities above.  Just handle them as we
support the extension already.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-09-16 19:18:01 -07:00
Jason Ekstrand 951cf94521 nir: Add explicit signs to image min/max intrinsics
This better matches all the other atomic intrinsics such as those for
SSBOs and shared variables where the sign is part of the intrinsic
opcode.  Both generators (GLSL and SPIR-V) know the sign from the type
of the image variable or handle.  In SPIR-V, signed min/max are separate
opcodes from unsigned.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Eric Anholt <eric@anholt.net>
2019-08-21 17:19:55 +00:00
Caio Marcelo de Oliveira Filho 5ed4e31c08 spirv: Drop lower_workgroup_access_to_offsets
Intel drivers are not using this anymore, and turnip still don't have
Compute Shaders, so won't make a difference to stop using this option.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Rob Clark <robdclark@chromium.org>
2019-08-10 22:15:35 -07:00
Lionel Landwerlin 0fb61dfdeb spirv: propagate access qualifiers through ssa & pointer
Not only variables can be flagged as NonUniformEXT but also
expressions. We're currently ignoring it in an expression such as :

   imageLoad(data[nonuniformEXT(rIndex)], 0)

The associated SPIRV :

   OpDecorate %69 NonUniformEXT
   ...
   %69 = OpLoad %61 %68

This changes propagates access qualifiers through ssa & pointers so
that when it hits a OpLoad/OpStore style instructions, qualifiers are
not forgotten.

Fixes failure the following tests :

   dEQP-VK.descriptor_indexing.*

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Fixes: 8ed583fe52 ("spirv: Handle the NonUniformEXT decoration")
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-07-26 14:09:55 +00:00
Lionel Landwerlin 86b53770e1 spirv: wrap push ssa/pointer values
This refactor allows for common code to apply decoration on all
ssa/pointer values. In particular this will allow to propagage access
qualifiers.

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Suggested-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-07-26 14:09:55 +00:00
Daniel Schürmann e352b4d650 spirv: Fix order of barriers in SpvOpControlBarrier
Semantically, the memory barrier has to come first to wait
for the completion of pending memory requests.
Afterwards, the workgroups can be synchronized.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-07-19 10:37:37 -07:00
Jason Ekstrand 6394680f6b spirv: Add a warning for ArrayStride on arrays of blocks
It's disallowed according to the SPIR-V spec or at least I think that's
what the spec says.  It's in a section explicitly about explicit layout
of things in the StorageBuffer, Uniform, and PushConstant storage
classes so it's not 100% clear that it applies with other storage
classes.  However, it seems like it should apply in general and
violating it can trigger (fairly harmless) asserts in NIR.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-07-16 17:02:08 -05:00
Caio Marcelo de Oliveira Filho 1210e8caaf spirv: Ignore ArrayStride for storage classes that should not use it
The stride was already overriden when using
lower_workgroup_access_to_offsets, so elaborate a bit the commentary
there.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-07-15 16:18:57 -07:00
Caio Marcelo de Oliveira Filho 026cfa1099 spirv: Fix stride calculation when lowering Workgroup to offsets
Use alignment to calculate the stride associated with the pointer
types.  That stride is used when the pointers are casted to arrays.

Note that size alone is not sufficient, e.g. struct { vec2 a; vec1 b;
} will have element an element size of 12 bytes, but the stride needs
to be 16 bytes to respect the 8 byte alignment.

Fixes: 050eb6389a "spirv: Ignore ArrayStride in OpPtrAccessChain for Workgroup"
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-07-15 16:18:46 -07:00
Caio Marcelo de Oliveira Filho 1a83c9a619 spirv: Implement SPV_EXT_demote_to_helper_invocation
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-07-08 08:57:25 -07:00
Caio Marcelo de Oliveira Filho 050eb6389a spirv: Ignore ArrayStride in OpPtrAccessChain for Workgroup
From OpPtrAccessChain description in the SPIR-V spec (1.4 rev 1):

    For objects in the Uniform, StorageBuffer, or PushConstant storage
    classes, the element’s address or location is calculated using a
    stride, which will be the Base-type’s Array Stride when the Base
    type is decorated with ArrayStride. For all other objects, the
    implementation will calculate the element’s address or location.

For non-CL shaders the driver should layout the Workgroup storage
class, so override any explicitly set ArrayStride in the shader.  This
currently fixes only the lower_workgroup_access_to_offsets case, which
is used by anv.

Reviewed-by: Juan A. Suarez <jasuarez@igalia.com>
2019-07-02 12:15:01 -07:00
Jason Ekstrand 81e51b412e nir: Make nir_constant a vector rather than a matrix
Most places in NIR, we treat matrices like arrays.  The one annoying
exception to this has been nir_constant where a matrix is a first-class
thing.  This commit changes that so a matrix nir_constant is the same as
an array nir_constant.  This makes matrix nir_constants a tiny bit more
expensive but shrinks all others by 96B.

Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-06-19 21:05:54 +00:00
Jason Ekstrand a54e397152 spirv: Only copy needed components for OpSpecConstantOp
Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-06-19 21:05:54 +00:00
Jason Ekstrand 96bb9c9277 spirv: Use a single path for OpSpecConstantOp of OpVectorShuffle
Now that nir_const_value is a scalar, there's no reason why we need
multiple paths here and it's just extra paths to keep working.  While
we're here, we also add a vtn_fail_if check that component indices are
in-bounds.

Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-06-19 21:05:54 +00:00
Jason Ekstrand 280e5442e5 spirv: Use vtn_constan_uint() for array lengths and gather components
Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-06-19 21:05:54 +00:00
Daniel Schürmann 7a858f274c spirv/nir: add support for AMD_shader_ballot and Groups capability
This commit also renames existing AMD capabilities:
 - gcn_shader -> amd_gcn_shader
 - trinary_minmax -> amd_trinary_minmax

Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
2019-06-13 12:44:23 +00:00
Daniel Schürmann 1b89ebeede nir/spirv: add support for the SubgroupBallotKHR SPIR-V capability
This capability is required for the VK_EXT_shader_subgroup_ballot extension.

Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
2019-06-13 12:44:23 +00:00
Daniel Schürmann de56ebadce nir/spirv: add support for the SubgroupVoteKHR SPIR-V capability
This capability is required for the VK_EXT_shader_subgroup_vote extension.

Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
2019-06-13 12:44:23 +00:00
Jason Ekstrand 5176805471 spirv: Implement SPV_EXT_fragment_shader_interlock
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-06-04 17:30:51 +00:00
Caio Marcelo de Oliveira Filho 61de825e11 spirv: Like Uniform, do nothing for UniformId
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-06-03 17:20:54 -07:00
Caio Marcelo de Oliveira Filho b4eff83180 spirv: Implement SpvOpCopyLogical
This is the same as SpvOpCopyObject but without the type checking,
which is how vtn_composite_copy works, so we just need to hook the
operation.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-06-03 17:20:54 -07:00
Caio Marcelo de Oliveira Filho 81586e9f53 spirv: Generalize OpSelect
SPIR-V 1.4 supports OpSelect over any composite type, and also allows
scalar boolean condition for vector types -- a case which we already
handled to support old GLSLang.

Added a helper function to recursively perform nir_bcsel, that makes
easier to support structs.

v2: Replace asserts() with vtn_fail_if().  (Jason)

v3: Simplify Condition and Result types verifications.  (Jason)

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-06-03 17:20:54 -07:00
Caio Marcelo de Oliveira Filho 17630291e5 spirv: Move OpSelect handling to a function
This will make a later change easier to review.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-06-03 17:20:54 -07:00
Caio Marcelo de Oliveira Filho 1f8546ba2f spirv: Implement OpPtrEqual, OpPtrNotEqual and OpPtrDiff
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-06-03 13:45:09 -07:00
Caio Marcelo de Oliveira Filho e45bf01940 spirv: Change spirv_to_nir() to return a nir_shader
spirv_to_nir() returned the nir_function corresponding to the
entrypoint, as a way to identify it.  There's now a bool is_entrypoint
in nir_function and also a helper function to get the entry_point from
a nir_shader.

The return type reflects better what the function name suggests.  It
also helps drivers avoid the mistake of reusing internal shader
references after running NIR_PASS on it.  When using NIR_TEST_CLONE or
NIR_TEST_SERIALIZE, those would be invalidated right in the first pass
executed.

Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-29 10:34:35 -07:00
Jason Ekstrand ddd08e1888 nir/builder: Remove the use_fmov parameter from nir_swizzle
This flag has caused more confusion than good in most cases.  You can
validly use imov for floats or fmov for integers because, without source
modifiers, neither modify their input in any way.  Using imov for floats
is more reliable so we go that direction.

Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
2019-05-24 08:38:11 -05:00
Caio Marcelo de Oliveira Filho 9b9f7030c6 spirv: Drop GOOGLE suffix from names incorporated to SPIR-V
SPV_GOOGLE_decorate_string and SPV_GOOGLE_hlsl_functionality1 were
incorporated to SPIR-V.  Let's pick the names used by SPIR-V core.

Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-05-21 11:52:41 -07:00
Caio Marcelo de Oliveira Filho 02d140ce9a spirv: Pick the right bitsize when doing SpvUConvert
Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-05-21 11:52:29 -07:00
Caio Marcelo de Oliveira Filho 7e5723d6d7 spirv: Generate proper NULL pointer values
Use the storage class address format information to pick the right
constant values for a NULL pointer.

v2: Don't add a deref_cast to the values.  (Jason)

v3: Update to use vtn_storage_class_to_mode() and
    vtn_mode_to_address_format() explicitly.  (Jason)

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-20 10:53:38 -07:00
Caio Marcelo de Oliveira Filho 83550b7dc4 spirv: Reuse helpers in vtn_handle_type()
And change vtn_storage_class_to_mode() to accept NULL as
interface_type.  In this case, if we have a SpvStorageClassUniform, we
assume it is uses an ubo_addr_format, like the code being replaced by
the helper.

That assumption is a problem, but no different than the previous
code.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-20 10:53:38 -07:00
Caio Marcelo de Oliveira Filho f9336751bc spirv: Add and use vtn_type_without_array() helper
v2: Renamed from vtn_interface_type. (Jason)
    Accept any type not only pointers.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-20 10:53:38 -07:00
Caio Marcelo de Oliveira Filho 8af9de0a38 spirv: Change vtn_null_constant() to use vtn_type
This is a preparation to handle OpConstantNull for pointers, we'll use
the vtn_type to get to the address format and then the appropriate
representation of NULL pointer.

v2: Move rest of body to use vtn_type. (Jason)

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-20 10:53:38 -07:00
Caio Marcelo de Oliveira Filho 31a7476335 spirv, radv, anv: Replace ptr_type with addr_format
Instead of setting the glsl types of the pointers for each resource,
set the nir_address_format, from which we can derive the glsl_type,
and in the future the bit pattern representing a NULL pointer.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-05-20 10:53:38 -07:00
Karol Herbst d11b807da5 nir: Add nir_op_vec helper
with that we can simplify code where nir vectors are created

v2: merge both lines in nir_vec

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-05-04 12:27:51 +02:00
Karol Herbst c91ea6343f vtn: handle bitcast with pointer src/dest
v2: use vtn_push_ssa and vtn_ssa_value

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-05-04 12:27:51 +02:00
Rob Clark a99c360a46 nir: add pass to lower fb reads
Signed-off-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
2019-05-02 11:19:22 -07:00
Juan A. Suarez Romero bbbe00a101 spirv: add missing SPV_EXT_descriptor_indexing capabilities
Add ShaderNonUniformEXT, UniformBufferArrayNonUniformIndexingEXT,
SampledImageArrayNonUniformIndexingEXT,
StorageBufferArrayNonUniformIndexingEXT,
StorageImageArrayNonUniformIndexingEXT,
InputAttachmentArrayNonUniformIndexingEXT,
UniformTexelBufferArrayNonUniformIndexingEXT and
StorageTexelBufferArrayNonUniformIndexingEXT capabilities.

Cc: mesa-stable@lists.freedesktop.org
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-04-30 09:22:45 +02:00
Caio Marcelo de Oliveira Filho 1fb6630636 spirv: Properly handle SpvOpAtomicCompareExchangeWeak
The code was handling the Weak variant in some cases, but missing
others, e.g. the get_deref_nir_atomic_op.  Add all the missing cases
with the same behavior of the non-Weak SpvOpAtomicCompareExchange.

Note that the Weak variant is basically an alias, as SPIR-V 1.3,
Revision 7 says

    "OpAtomicCompareExchangeWeak

    Deprecated (use OpAtomicCompareExchange).

    Has the same semantics as OpAtomicCompareExchange."

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-29 19:02:44 -07:00
Caio Marcelo de Oliveira Filho 7e2684ce01 spirv: Handle SpvOpDecorateId
This operation decorate with an Id instead of a Literal or String.

It is used by HlslCounterBufferGOOGLE (provided by
SPV_GOOGLE_hlsl_functionality1).  Even if we don't do anything with
that decoration, we must be able to parse SPIR-V that uses it.

Fixes: 891886da2f "spirv: Add no-op support for VK_GOOGLE_hlsl_functionality1"
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-23 14:58:01 -07:00
Caio Marcelo de Oliveira Filho 7b66d584a3 spirv: Rename vtn_decoration literals to operands
Decorations (and ExecutionModes) can have not only literals, but also
Ids associated with them.  So rename the field to the more general
name "Operand" used by the spec.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-23 14:58:01 -07:00
Iago Toral Quiroga e6ee07a664 compiler/spirv: move the check for Int8 capability
So it is right after the checks for the other various Int* capabilities.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-18 13:23:03 +02:00
Caio Marcelo de Oliveira Filho a0dae78e72 spirv: Tell which opcode or value is unhandled when failing
v2: When available, include the opcode name too. (Karol)

v3: Use more to_string helpers. (Karol)
    Include the wrong bit_size in those failures.
    Include the capability number in spv_check_supported.
    Provide vtn_fail_with_* macros to avoid noise in the call sites.

v4: Provide macros only for opcode and decoration, which have enough
    usages to justify them. (Jason)

Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-04-16 11:11:10 -07:00
Samuel Pitoiset bbe8febd93 spirv: add SpvCapabilityFloat16 support
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
2019-04-15 10:43:52 +02:00
Jason Ekstrand 893dd34702 spirv: Drop some unneeded bit size switch statements
Now that nir_const_value is a scalar, we don't need the switch on bit
size in order copy components around properly.

Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-04-14 22:25:56 +02:00
Karol Herbst 14531d676b nir: make nir_const_value scalar
v2: remove & operator in a couple of memsets
    add some memsets
v3: fixup lima

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net> (v2)
2019-04-14 22:25:56 +02:00
Karol Herbst 73d883037d spirv: reduce array size in vtn_handle_constant
we already assert above that there are no more than 3 sources, so it
doesn't make sense to use an array of 4 sources

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-14 22:25:56 +02:00
Caio Marcelo de Oliveira Filho bd73531677 spirv: Add support for DerivativeGroup capabilities
As defined in SPV_NV_compute_shader_derivatives. These control how the
invocations are arranged in a CS when doing derivative and related
operations (which are also enabled by the extension).

Since we expect valid SPIR-V, we don't need to do more work at SPIR-V
level to enable the derivative and related operations to be called.

Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-04-08 19:29:33 -07:00
Jason Ekstrand 8ed583fe52 spirv: Handle the NonUniformEXT decoration 2019-03-25 16:12:09 -05:00
Caio Marcelo de Oliveira Filho 9d0ae777dd spirv: Use interface type for block and buffer block
Also handle GLSL_TYPE_INTERFACE the same way we do GLSL_TYPE_STRUCT in
various places.  Motivated by ARB_gl_spirv work, that will take
advantage of the interface types when handling NIR coming from SPIR-V.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-03-23 10:22:39 -07:00
Jason Ekstrand 6e19348ad1 spirv: Drop inline tg4 lowering
Reviewed-by: Karol Herbst <kherbst@redhat.com>
2019-03-21 02:58:41 +00:00
Juan A. Suarez Romero efcf9c9f9f nir: deref only for OpTypePointer
Fixes dEQP-VK.binding_model.buffer_device_address.* and
dEQP-VK.ssbo.phys.layout* Vulkan CTS tests.

v2: set val->type->stride in the section below (Jason)

v3: restore val->type->type to original place (Jason)

Fixes: d0ba326f23 ("nir/spirv: support physical pointers")
CC: Karol Herbst <kherbst@redhat.com>
CC: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-03-20 19:26:32 +00:00
Jason Ekstrand cbfe31ccbe Revert "nir: const `nir_call_instr::callee`"
This reverts commit db57db5317.  When
building IR, nothing is really immutable and, since C has no concept of
constness propagating beyond the first pointer, we have to be vary
careful with how we use it.  To just throw const into a function like
this is a lie.

Instead, we should just drop the unneeded const in spirv_to_nir which
this commit does along with the revert.
2019-03-19 10:19:42 -05:00
Karol Herbst d0ba326f23 nir/spirv: support physical pointers
v2: add load_kernel_input

Signed-off-by: Karol Herbst <kherbst@redhat.com>

squash! nir/spirv: support physical pointers
2019-03-19 04:08:07 +00:00
Karol Herbst c95afe56a8 nir/spirv: handle kernel function parameters
the idea here is to generate an entry point stub function wrapping around the
actual kernel function and turn all parameters into shader inputs with byte
addressing instead of vec4.

This gives us several advantages:
1. calling kernel functions doesn't differ from calling any other function
2. CL inputs match uniforms in most ways and we can just take advantage of most
   of nir_lower_io

v2: move code into a seperate function
v3: verify the entry point got a name
    fix minor typo
v4: make vtn_emit_kernel_entry_point_wrapper take the old entry point as an arg

Signed-off-by: Karol Herbst <kherbst@redhat.com>
2019-03-19 04:08:07 +00:00
Karol Herbst 659f333b3a glsl: add packed for struct types
We need this for OpenCL kernels because we have to apply C rules for alignment
and padding inside structs and for this we also have to know if a struct is
packed or not.

v2: fix for kernel params

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
2019-03-19 04:08:07 +00:00
Jason Ekstrand fe9a6c0f14 nir/builder: Add a vector extract helper
This one's a tiny bit better than what we had in spirv_to_nir because it
emits a binary tree rather than a linear walk.  It also doesn't leave
around unneeded bcsel instructions for a constant index and returns an
undef for constant OOB access.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2019-03-15 23:10:26 -05:00
Xavier Bouchoux c5236fc6e2 nir/spirv: Fix assert when unsampled OpTypeImage has unknown 'Depth'
'dxc' hlsl-to-spirv compiler appears to emit 2 (Unknown) in the depth field,
when the image is not sampled and the value is not needed.

Previously, shaders failed with:

SPIR-V parsing FAILED:
    In file ../src/compiler/spirv/spirv_to_nir.c:1412
    !is_shadow
    632 bytes into the SPIR-V binary

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
2019-03-11 23:28:39 +01:00