The comment claims this was to unroll loops, but nir_opt_loop doesn't do that.
Whatever issue the AGX code was originally working around, it doesn't apply now
(I confirmed we produce similar code with or without the pass). In the meantime,
Panfrost and PowerVR cargo-culted the same broken logic. Drop it all.
Closes: #14732
Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39588>
Mali hardware handles a bunch of OOB conditions by using addresses with
the top bit set. When the top bit is set, any load/store from a shader
will treat the address as OOB and read zero and discard writes. We can
use this to implement ro_sink_address_poly.
Signed-off-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38547>
There is nothing IR about it. It's really the compiler interface file,
so it should all go in pan_compiler.h.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38753>
While we're at it, rename pan_lower_* to pan_nir_lower_* for consistency
with everything else. The intention is that eventually this will be a
private header and drivers will stop calling these passes themselves.
However, that is a long way off so for now we'll just move them to the
sensibly named place. Notably, pan_preprocess/optimize/postprocess_nir
stay in pan_compiler.h because those are intended to be external APIs
indefinitely.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38753>
The only thing that only pulls in bifrost is our CL compiler for
pre-builts. Everything that uses pan_shader.h pulls it in because
that's header-only and calls into both back-ends anyway. Some day
we might want to figure out how to properly dead-code the midgard
compiler for Vulkan but today is not that day.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38753>
src/panfrost/util is all compiler stuff, we just put it in util/ so it
could be used by both midgard and bifrost. We also rename all the files
to have a pan_nir prefix. We'll rename the functions later.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38753>
It was this way for a while but then someone decided that we should try
to hide midgard and get it off the main path. However, we still try to
present a unified interface to the GL driver. The combination of these
two things means that those common compiler interfaces have to live
outside of panfrost/compiler which makes everything more awkward than
it needs to be. We either need to move it back into src/gallium and
make abstracting across the two the GL driver's problem (which breaks
other tooling) or need to put both under src/panfrost/compiler. The
midgard compiler can just as easily bitrot in panfrost/compiler/midgard.
The first step is moving the bifrost compiler.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Acked-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38753>
If we're using the singleton, we need to add to it.
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38638>
This fixes an issue a bunch of different components were all working
around themselves where sometimes we don't have a workgroup size but
workgroup_size_variable is false. This also fixes asahi_clc, which
didn't have the workaround and was assuming zero (but not variable!)
workgroup sizes everywhere.
LoLed-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Acked-by: Mel Henning <mhenning@darkrefraction.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38538>
Texel buffers are currently described by a TextureDescriptor,which leads
to restrictive limits on size and alignment.
These limits can be avoided by using a BufferDescriptor instead.
This requires first embedding a ConversionDescriptor into some of the
currently empty space of the BufferDescriptor, and modifying the
compiler so that instead of outputting TEX_FETCH, it will:
1. Load the ConversionDescriptor with LD_PKA
2. Get the buffer address with LEA_BUF[_IMM]
3. Use LD_CVT to get the value
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37007>
Every caller of pan_shader_lower_texture() immediatly called
pan_shader_postprocess() and every caller of pan_shader_postprocess()
lowered textures except blend shaders and those don't texture anyway.
Reviewed-by: Olivia Lee <olivia.lee@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38265>
nir_lower_compute_system_values will attempt to lower
load_workgroup_size unless workgroup_size_variable is set. For precomp
shaders, the workgroup size is set statically for each entrypoint by
nir_precompiled_build_variant. Because we call
lower_compute_system_values early, it sets the workgroup size to zero.
Temporarily setting workgroup_size_variable while we are still
processing all the entrypoints together inhibits this.
Signed-off-by: Olivia Lee <olivia.lee@collabora.com>
Fixes: 20970bcd96 ("panfrost: Add base of OpenCL C infrastructure")
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37799>
lseek can return a negative value on error here. While it's not likely
to happen, let's add some error-checking here to prevent bad behavior if
we're unlucky.
CID: 1648299
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36724>
We now have lower_texture_early and lower_texture.
lower_texture_early handle nir_lower_tex and (in the future) could handle
anything that is backend specific that need to happen before nir_lower_io.
lower_texture handles actual lowering of backend specific things that
must happen after nir_lower_tex and nir_lower_io.
This allows us to finally not run nir_lower_tex two times in panvk.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Olivia Lee <olivia.lee@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36776>
As we are going to move texture and IO lowering, this split preprocess
functions in two, one handling preprocess the other postprocess.
The split is done right before lower_io and has no functional change for
now.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Olivia Lee <olivia.lee@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36776>
Rather than splitting the GPU ID in two, let the GPU ID users do that
when they need.
We also rework the model detection to use a mask so we can customize
the mask if the version major/minor fields are required to differentiate
two GPUs with the same arch major/minor and product major.
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Reviewed-by: Lars-Ivar Hesselberg Simonsen <lars-ivar.simonsen@arm.com>
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35155>
Similar to nir_lower_alu_width(), the callback can return the
desired number of components for a phi, or 0 for no lowering.
The previous behavior of nir_lower_phis_to_scalar() with lower_all=true
can be elicited via nir_lower_all_phis_to_scalar() while the previous
behavior with lower_all=false now corresponds to nir_lower_phis_to_scalar()
with NULL callback.
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Mel Henning <mhenning@darkrefraction.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35783>
We were assigning the scratch size of the whole lib... that was causing
very big TLS usage on draw indirect.
With this TLS usage is way lower now at the cost of running more pass
for every variants.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35639>
At some point CL C shader switched to a format hash approach and we
forgot to properly migrate here.
This should fix all printf again.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35639>
We've been inconsistenly using panfrost_ and pan_ as a prefix for
the common helpers/structs. Let's finally make a clear cut by
prefixing everything that lives in src/panfrost/* with pan_.
No functional changes here, this is just renaming and reformatting
changes.
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Eric R. Smith <eric.smith@collabora.com>
Reviewed-by: Mary Guillemard <mary.guillemard@collabora.com>
Acked-by: Daniel Stone <daniels@collabora.com>
Reviewed-by: Ryan Mckeever <ryan.mckeever@collabora.com>
Reviewed-by: Olivia Lee <olivia.lee@collabora.com>
Reviewed-by: Lars-Ivar Hesselberg Simonsen <lars-ivar.simonsen@arm.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34767>
only the GL driver actually wants this, neither panvk nor internal shaders do.
Cc'd as a prereq to the next patch
Cc: mesa-stable
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/34193>
Those are quite useful for debugging and having sanity checks in place.
It is also quite tidious to get ride of all asserts in every headers we
would ever want to use, lets just accept those now.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32720>
This allows compiling CL shaders into a single SPIR-V library per arch,
NIR call bindings for each functions and precompilled binaries for each
entrypoints.
We are only going to support Bifrost and Valhall for this.
Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/32720>
