Each shader stage has its own "early preamble" flag.
Early preamble is likely an optimization to hide some of latency
when loading UBOs into consts in the preamble.
Early preamble has the following limitations:
- Only shared, a1, and consts regs could be used
(accessing other regs would result in GPU fault);
- No cat5/cat6, only stc/ldc variants are working;
- Values writen to shared regs are not accessible by the rest
of the shader;
- Instructions before shps are also considered to be a part of
early preamble.
Note, for all shaders from d3d11 games blob produced preambles
compatible with early preamble mode.
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15901>
Since these were reverse-engineered, it's become clear that IBO
descriptors are just a subset of texture descriptors, and bindless reads
of readonly images actually use isam on the IBO descriptor, further
confirming that the two are always compatible, even if not all of the
texture fields exist for IBOs. It's pointless to have a separate type
for IBOs, and just leads to things getting out-of-sync unnecessarily
which has already happened. Just remove it and use TEX_CONST insted.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15114>
Inferring from blob's cmdstream the size of shader instruction
cache for:
- a630 is 64
- a650 is 128
- a660 is 128
On a650 and a660 gpu could hang if we exceed the limit. Though
it is not reproducible with computerator or a single amber
test. Also while blob limits the size to 128 - Turnip still
hangs with it but does not hang with the limit of 127.
On a630 there seem to be no hang when limit is exceeded.
Fixes the hang of compute shader in Alien Isolation on a650/a660.
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/14044>
This backend provides very basic a4xx support. It's enough to run
kernels with explicit stg/etc ops, but not with stgb/ldgb type access.
There is no perfcounter support hooked up yet either.
Signed-off-by: Ilia Mirkin <imirkin@alum.mit.edu>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12784>
Move away from using gpu_id as the primary means to identify which
adreno we are running on, as future GPUs (starting with 7c3) stop
providing a gpu_id as a new naming scheme is introduced.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12159>
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>
Move everything into a struct assocated with the pipe_fence_handle, so
that the drm layer can fill in the seqn/fd fences directly.
This will give us a comvenient place to insert a util_queue_fence in the
next commit.
While we're at it, extract the uint32_t fence (previously called
'timestamp' in place, a kgsl legacy) into a struct that encapsulates
both the kernel fence and the userspace fence.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10444>
Most of them were actually unused. The memory type (KMEM vs SMI) only
applied to very old a2xx era devices that had a small/fast stacked
memory (SMI) vs normal memory (KMEM). And the cache flags are ignored
(ie. everything is writecombine), but we can add new cache flags later
when they actually do something.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10444>
On a6xx/a5xx there is such dependency between branchstack bitfield
and the amount of nested ifs, which could be seen with blob:
IFs BRANCHSTACK
0 0
1 1
2 2
3 2
4 3
5 3
6 4
...
59 30
60 31
61 31
62 32
63 32
64 32
Remove open-coded branchstack for a5xx compute along the way.
Fixes tests:
dEQP-VK.spirv_assembly.instruction.compute.float16.opvectorshuffle.344
dEQP-VK.spirv_assembly.instruction.graphics.float16.opvectorshuffle.344_vert
dEQP-VK.spirv_assembly.instruction.graphics.float16.opvectorshuffle.444_geom
dEQP-VK.spirv_assembly.instruction.graphics.float16.opvectorshuffle.244_tessc
dEQP-VK.spirv_assembly.instruction.graphics.float16.opvectorshuffle.344_frag
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9859>
This needs to be part of the compiler because it's the only piece that
we always have access to in all the places ir3_optimize_loop() is
called, and it's only enabled for the whole Vulkan device. Right now
it's just used for constraining vectorization, but the next commit adds
another use.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7573>
When float16 is enabled this will allow to pass a number of
float16 tests.
When A6XX_SP_FLOAT_CNTL_F16_NO_INF is set - all operations which
generate +-infinity generate +-MAX_HALF_FLOAT.
Fixes some tests from:
dEQP-VK.spirv_assembly.instruction.*.float16.*
dEQP-VK.spirv_assembly.instruction.*.float_controls.fp16.*
E.g.:
dEQP-VK.spirv_assembly.instruction.graphics.float16.arithmetic_1.sinh_vert
dEQP-VK.spirv_assembly.instruction.compute.float16.arithmetic_4.length
dEQP-VK.spirv_assembly.instruction.compute.float_controls.fp16.input_args.log_denorm_flush_to_zero_nostorage
dEQP-VK.spirv_assembly.instruction.compute.float_controls.fp16.input_args.log2_denorm_flush_to_zero_nostorage
dEQP-VK.spirv_assembly.instruction.compute.float_controls.fp16.input_args.inv_sqrt_denorm_flush_to_zero_nostorage
Signed-off-by: Danylo Piliaiev <dpiliaiev@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9840>
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>
The registers were actually different per-stage even though we used the
same type, which resulted in a bunch of incorrectly programmed fields
and confusion. Move the stage-specific values to the registers
themselves, which makes things much less confusing and makes it possible
to set "mergedregs" correctly.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9493>
Reduce noise in a6xx.xml by removing LO/HI versions of address registers.
Also fix type="address" registers in register packing (use bit size instead
of checking for "waddress" to use qword)
Signed-off-by: Jonathan Marek <jonathan@marek.ca>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8423>
Move out of ir3_parse_asm() so we can re-use it in disasm test for
round-tripping asm/disasm. We don't want failures to be fatal (yet)
as there are still some things missing from the assembler.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8175>
computerator depends on ir3_parser.h, which is a generated file, but
this dependency is not expressed in the build.
Fixes: 1e8808a4a0 ("freedreno/ir3: refactor out helper to compile shader from asm")
Signed-off-by: Daniel Stone <daniels@collabora.com>
Reviewed-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7870>
Right now if the shader indirects on some large constant array, we see NIR
load_consts (usually from the const file) of its contents into general
registers, then indirection on the GPRs. This often results in register
allocation failures, as it's easy to go beyond the ~256 dwords of
registers per invocation.
By moving the large constants to a UBO, we can load an arbitrary number of
them. They also can be theoretically moved to the constant reg file (~2k
dwords), though you're unlikely to hit this path without an indirect load
on your large constant, and we don't yet let UBO indirect loads get moved
to constant regs.
This possibly won't work out right if we have 16-bit load_constants, but
without other MRs in flight we won't see 16-bit temps to be lowered to
this.
This allows 2 kerbal-space-program shaders to compile that previously
would fail, and fixes the new dEQP-VK and -GLES2 tests I wrote that
dynamically index a 40-element temporary array of float/vec2/vec3/vec4
with constant element initializers.
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/2789
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5810>
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>
This gets us consistent hcN access with our drivers, for experimenting.
We don't know what the other bit does yet, but let's not have to debug
that later.
Reviewed-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6179>
Instead of building the adreno/foo.xml headers from the toplevel, split
out a subdir(). This fits better with how meson likes things to be
structured. But it does require fixing a bit about how gen_header.py
resolves imports, ie. it cannot assume the src file is at the root of
the $RNN_PATH.
This is needed for the next patch, to add support for installing the
register database for use with installed tools.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6154>
It turns out that this clears CP_LOAD_STATE6 packets, including
disabling any pending loads for SS6_INDIRECT/SS6_BINDLESS (these loads
don't actually happen until the draw itself, and I'm not sure if they
happen if the state is unused by the shader) and marking constants and
UBO descriptors loaded with SS6_DIRECT as invalid. It's used very
differently from HLSQ_UPDATE_CNTL on a4xx from whence the name came, and
unlike on a4xx it's not readable, so this probably doesn't line up with
HLSQ_UPDATE_CNTL on a4xx.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5877>
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>
Deduplicate a bit of hand-building of ir3_shader/_variant from
computerator and delay test. This also removes the need for
external things to depend on generated ir3_parser header.
Signed-off-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5508>
Marek Olšák
Danylo Piliaiev
Rob Clark
Connor Abbott
Ilia Mirkin
Caio Marcelo de Oliveira Filho
Jonathan Marek
Daniel Stone
Eric Anholt