Use the new nir_var_mem_image mode for images that are not known to be
used with a sampler (i.e. storage images).
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4743>
Now that all drivers have been patched to convert sysvals to input
varyings when they have too, we can safely declare PointCoord as a sysval
too.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13017>
Now that all spirv_to_nir() users take care of converting sysvals to
varyings, we can unconditionally declare FragCoord as a sysval.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13017>
If an OpVariable's initializer is undef, there is no need to
initialize the variable.
v2: Comment the code (Caio)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13030>
The LLVM-SPIRV translator creates variables with initializers, but
most of those are actually undef initializers. We can just skip
composites that are entirely made of undefs, but for partially undefs,
we will still zero initialize.
v2: Rename wa_llvm_spirv_undef_initializer to wa_llvm_spirv_ignore_workgroup_initializer (Caio)
Limit workaround to OpenCL (Caio)
Make workaround clearer (Caio)
v3: Only apply workaround on workgroup storage (Caio)
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13030>
This allows the variables decorated with RelaxedPrecision to have the
proper precision. It is worth to note that the decorator can be
applied on other cases, but those would be handled on the future.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7614>
spirv_to_nir now requires NIR variables to be created for everything.
Fixes: 10b3eecd36 "spirv: Don't remove variables used by resource..."
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8637>
Move it out the "cs" sub-struct, since the bit will be used for other
shader stages in the future.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225>
Be consistent with other usages in Vulkan and SPIR-V, and the recently
added workgroup_size field.
Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
VS outputs are "per vertex" but not the kind of I/O we want to match
with this helper. Change to a name that covers the "arrayness"
required by the type.
Name inspired by the GLSL spec definition of arrayed I/O.
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10493>
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>
When tyding up this section in 1e5b09f42f ("spirv: Tidy some repeated
if checks by using a switch statement.") the break got lost. It is
not a real problem because the next case just break, but better to
have it explicitly here instead of a FALLTHROUGH.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9440>
Newer versions of SPIR-V require that all the global variables used by
the entry point are declared (in contrast to only I/O in previous
versions), so there's no need to remove dead variables or keep track
of the indirectly used variables.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8456>
OpEntryPoint declares the list of variables in Input and Output
storage classes that are used. Use that information to skip creating
other variables from such storage classes that are unused by the entry
point.
After that change, is not necessary to use remove dead variables for
those types of variables; and because of that is also not necessary to
lower initalizers for output variables.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8456>
Fail when parsing Initializers used in Variables with Storage Classes
that doesn't support it.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8820>
This will be used to implement
VK_KHR_zero_initialize_workgroup_memory.
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8708>
Not only these are recalculated in nir_shader_gather_info, but
currently they are also counting all the images / textures in the
module instead of in the shader (entrypoint).
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8786>
In Vulkan, for some variable modes, the generated NIR will have derefs
pointing to resource index intrinsics instead of the variable. This
was letting nir_remove_dead_variables pass remove those variables,
which would lose information relevant for later passes after
spirv2nir.
Add a set to keep track of such variables and prevent them to be
removed when producing the NIR output.
Issue reported by Rhys.
Fixes: c4c9c780b1 ("spirv: Remove more dead variables")
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8706>
Uses same NIR intrinsic as glsl_to_nir. Make it an option so it is
easy later to move Vulkan drivers incrementally to use it.
Fixes piglit test spec/arb_gl_spirv/execution/ssbo/unsized-array-length.
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/3691
Fixes: 15e43907 ("iris: Enable ARB_gl_spirv and ARB_spirv_extensions")
Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8136>
Unless the non uniform work group extension is supported, this
just aliases workgroupsize, so just do that for now.
Fixes:
CL CTS basic enqueued_local_size
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7642>
The SPV_KHR_ray_tracing extension adds 6 new storage classes which is a
bit on the ridiculous side. In order to avoid adding that many variable
modes to NIR, we make a few simplifying assumptions:
1. CallableData and RayPayload data actually lives on the stack
somewhere, presumably in the caller's stack. We assume that these
are no different from global variables and use nir_var_shader_temp
for them. We still need a separate storage class for the incoming
variants but only so we can figure out which one the incoming one
is and lower it to something useful.
2. There's no difference between incoming CallableData and RayPaolad
data. We can use a single storage class for both.
3. ShaderRecordBuffer data is just a global memory access. This lets
us avoid NIR variables entirely and just fetch the pointer via the
shader_record_ptr system value and it's accessed using a 64-bit
global memory pointer.
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6479>
Missing in this commit are NIR intrinsics for the ObjectToWorld and
WorldToObject built-ins. Those are matrices and so they take a bit more
work and justify a separate commit. For now, we add the enums and leave
the SYSTEM_VALUE <-> nir_intrinsic conversion commented out.
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6479>
For now, we assume its a 64-bit global pointer.
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6479>
We already fail in these same cases in vk_desc_type_for_mode. These
additional assertions are just extra code to update.
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6479>
Most of this is fairly straightforward; we just set all the modes on any
derefs which are generic. The one tricky bit is OpGenericCastToPtrExplicit.
Instead of adding NIR intrinsics to do the cast, we add NIR intrinsics
to do a storage class check and then bcsel based on that.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6332>
We rename it to "modes" to make it clear that it may contain more than
one mode and adjust all the uses of nir_deref_instr::modes to attempt to
handle multiple modes.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6332>
