The immediate addition can easily be handled by nir_opt_offsets, which
will also take any driver limits into account.
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35092>
In ir3, SSBO offsets are in units of the accessed type size so we want
to start using the new offset_shift index.
Even though the shift is implicit for the ir3 intrinsics, we use
nir_intrinsic_copy_const_indices when creating them so we need to make
sure our indices match the ones used by the generic intrinsics.
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35092>
For intrinsics supporting offset_shift, dealing with their offset is a
bit tricky as we cannot simply add a byte offset to it anymore (which is
what most passes want to do). This commit adds some helpers to add byte
offsets (and adjusting offset_shift accordingly) so that individual
passes don't have to worry about this.
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35092>
For load/store intrinsics that take an offset, this specifies the amount
the offset is shifted left to calculate the final offset:
offset = (offset_src + base) << offset_shift
This is useful for backends that have memory operations that use offset
units other than bytes (i.e., where the shift is implicit).
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/35092>
New type was not handled in the switch which lead to hitting following
assert when running tests with pipeline cache:
deqp-vk: ../src/compiler/glsl_types.c:3334: decode_type_from_blob: Assertion `!"Cannot decode type!"' failed.
Fixes: 9e5d7eb88d ("compiler/types: add a bfloat16 type")
Signed-off-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36833>
Usually we can fold most ldc and ldcx into the instruction using it,
however there are a couple of cases where we can't, e.g. when there is an
indirect offset.
Moving the ldc(x) down to the consumer leads to increase value ranges for
uniform registers, but lowering them for normal registers.
Totals:
CodeSize: 914650304 -> 914469536 (-0.02%); split: -0.05%, +0.03%
Number of GPRs: 3879754 -> 3863818 (-0.41%); split: -0.42%, +0.01%
Static cycle count: 1073273107 -> 1073101189 (-0.02%); split: -0.09%, +0.08%
Spills to reg: 67219 -> 67707 (+0.73%); split: -0.10%, +0.83%
Fills from reg: 79733 -> 80456 (+0.91%); split: -0.10%, +1.01%
Max warps/SM: 3666036 -> 3672668 (+0.18%); split: +0.18%, -0.00%
Totals from 24235 (27.66% of 87622) affected shaders:
CodeSize: 444747392 -> 444566624 (-0.04%); split: -0.11%, +0.07%
Number of GPRs: 1360384 -> 1344448 (-1.17%); split: -1.20%, +0.03%
Static cycle count: 806310857 -> 806138939 (-0.02%); split: -0.12%, +0.10%
Spills to reg: 35826 -> 36314 (+1.36%); split: -0.19%, +1.55%
Fills from reg: 31863 -> 32586 (+2.27%); split: -0.26%, +2.53%
Max warps/SM: 911328 -> 917960 (+0.73%); split: +0.74%, -0.01%
Reviewed-by: Mel Henning <mhenning@darkrefraction.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36536>
Now that we lower all load_per_vertex_input to
r600_load_per_vertex_input we can remove some dead code
and also change the intrinsic to use only one source value.
Signed-off-by: Gert Wollny <gert.wollny@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36488>
Prepare for mesh shader support, change all shared code to
use MESA_SHADER_MESH_STAGES.
Use command:
find src/gallium/auxiliary/ -type f -not -path '*/.git/*' -exec sed -i 's/\bMESA_SHADER_STAGES\b/MESA_SHADER_MESH_STAGES/g' {} +
find src/compiler -type f -not -path '*/.git/*' -exec sed -i 's/\bMESA_SHADER_STAGES\b/MESA_SHADER_MESH_STAGES/g' {} +
find src/mesa -type f -not -path '*/.git/*' -exec sed -i 's/\bMESA_SHADER_STAGES\b/MESA_SHADER_MESH_STAGES/g' {} +
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Yonggang Luo <luoyonggang@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36596>
We're going to support mesh shader in GL, the stage order
is not sorted as comment described.
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Yonggang Luo <luoyonggang@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36596>
glsl needs to plumb this from the backend. we should clean up
nir_lower_subgroups to use this later but I don't have time to churn everything
right now.
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <maraeo@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36649>
As float16_t may comes from arm_neon.h:
arm_neon.h:74:16: error: conflicting declaration 'typedef __fp16 float16_t'
74 | typedef __fp16 float16_t;
Also it is std::float16_t from <stdfloat> as part of C++23
So place it under namespace mesa to avoid conflict when upgrading xxhash.h to v0.8.3
Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Acked-by: Eric Engestrom <eric@igalia.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36654>
For getting clc_helpers.cpp can be compiled with gcc/mingw
Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Reviewed-by: Eric Engestrom <eric@igalia.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36597>
It's not used by anyone.
Acked-by: Mike Blumenkrantz <michael.blumenkrantz@gmail.com>
Acked-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Yonggang Luo <luoyonggang@gmail.com>
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/36569>
