The full form for ldg.a/stg.a offset is:
g[reg_address + reg_offset << (imm_shift + 2) + imm_offset << 2]
where imm_shift is in [0, 3] and imm_offset is in [0, 3]
a6xx blob was found to produce a bit simplier offset calculations
for TES/TCS shaders in GTA V:
[c002000a_03c14215] ldg.a.f32 r2.z, g[r1.y+((r2.z+1)<<2)], 3;
[c0020004_01c14609] ldg.a.f32 r1.x, g[r1.y+((r1.x+3)<<2)], 1;
Our new syntax:
stg.a.u32 g[r2.x+(r1.x+1)<<2], r5.x, 1
stg.a.u32 g[r2.x+r1.x<<4+3<<2], r5.x, 1
ldg.a.f32 r1.w, g[r1.y+(r1.w+1)<<2], 3
ldg.a.f32 r1.w, g[r1.y+r1.w<<5+2<<2], 3
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11431>
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>
immediates_count and immediates_size are supposed to have the same
units, but it was only incrementing immediates_count by 1. While we're
here, also fix the case where constants are specified out-of-order.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10291>
We want to use the local_size when available to calculate the threadsize
in ir3, and we need it to work with e.g. computerator where we don't
have a nir shader. Add a local_size field and use that in computerator
instead of of a separate structure that's inaccessable to core ir3.
Also set a dummy local_size in the tests to avoid a divide-by-zero.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9498>
Before, decoding showed that there is an error:
sam.base0 (f32)(xyzw)r0.x, r0.z, a1.x ; no field 'HAS_SAMP',
WARNING: unexpected bits[0:7] in #cat5-samp-s2en-bindless-a1: 0x1 vs 0x0
After:
sam.base0 (f32)(xyzw)r0.x, r0.z, s#1, a1.x
Fixes textures on the ground in TauCeti Vulkan Technology Benchmark
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9628>
By keeping track of nonuniform access from nir and storing it to ir3.
Signed-off-by: Hyunjun Ko <zzoon@igalia.com>
Reviewed-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Reviewed-by: Eric Anholt <eric@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9125>
This fixes the assembly for many scenarios where you want to use shader
replacement.
Note: unfortunately this leaks the identifier string created while
lexing, but I couldn't find a way to avoid leaking it except for
bringing in ralloc or something (which would be way more complicated).
The only other place doing something similar in mesa is the glsl parser,
which is using ralloc (actually a linear context).
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9463>
Trying to specify a floating-point value in a @const line would result
in it getting interpreted as a FLUT value and failing parsing. Fix this
by making the various FLUT tokens include the surrounding parentheses.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9463>
The offset can also be a register, in which case we need to shuffle
around the src order. Add a few more test vectors to cover each
permutation (no offset, immed offset, gpr offset).
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Note that this shows up a slight encoding difference compared to test
vector extracted from blob deqp runs. We think these should be dontcare
bits. For now, add a note and replace the encoded value in the disasm
test.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Add a `.b`.. for the atomic instructions it should be `atomic.b.op` but
for now put the `.b` at the end to simplify life for the existing disasm
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Well, really just resinfo.. dealing with the different ldib/stib syntax
for a6xx+ vs earlier seems a bit too painful to deal with. But resinfo
at least gives us some encoding test coverage of this group of instrs.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Currently ir3 (incl emit_cat5()) expects the samp/tex src register to be
first.. which requires some fixup for the parser to match.
TODO we might want to revisit the src reg order when adding new instr
packing/encoding. For now, lets just make the parser match the rest of
ir3.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
There was some src2 vs src3 confusion, but since the syntax is like:
ldl.f32 rDst, l[rBase+off], ncomp
it makes more sense to call the offset src2 and ncomp src3, than the
way we had it. This is also easier to deal with for the ir3 assembly
parser.
Also, src_offset was only ever used by the assembly parser, and was
handled incorrectly in emit_cat6(), resulting that cat6 load instrs
would not work properly in (for ex) computerator. Since we are
cleaning things up, drop src_offset and make the asm parser work in
the same way as the nir->ir3 frontend.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Convert everything remaining over to the version which takes # of
register (src + dst) and drop the ir3_instr_create2() version.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
Update the IR and packer to handle the additional cat0 fields, in
prep for adding support in the assembler (in prep for adding round
trip parsing/packing test coverage).
We don't actually use these yet from the ir3 compiler, but at least
this is one less thing to worry about when we start trying to use
them.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
It seems the src_offset and dst_offset are unused for these, and the
offset is expected to be an immediate register. Also we forgot to add a
dummy dst for the store instructions.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7386>
Rework setup_{input,output} to be called during emit_intrinsic, in a way
which allows struct/array/matrix type varyings to work.
This allows turnip to pass dEQP-VK.glsl.linkage.varying.struct.*
Signed-off-by: Jonathan Marek <jonathan@marek.ca>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6181>
I got tripped up again with the index vs count vs size fields and I'd
rather we didn't store the redundant info. Settle on immediates_count as
"how many dwords of immediates we have"
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5990>
Make it an rzalloc'd ptr instead of embedded struct, so it can serve as
the mem ctx for immediates. This gets rid of needing to explicitly free
the immediates, so one less thing to deal with when moving const_state.
(Also, after we move const_state to the shader variant, we won't need
one for binning pass variants)
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5508>
We are going to want to move this back to the variant, and come up with
a different strategy for binning/nonbinning to share the same constant
layout, in order to implement shader-cache support. (Since then we
can have a mix of dynamically compiled variants and cache hits, so there
is no good place to serialize the const-state.)
To reduce the churn as we re-arrange things, move direct access to the
const-state to a helper fxn. This patch is the boring churny part.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5508>