fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-22 13:08:09 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2
142964 commits 66 branches 936 tags 949 MiB
Commit graph

931 commits

Author SHA1 Message Date
Alyssa Rosenzweig
f9442e46ff pan/mdg: Garbage collect silly quirk 
As discussed with Jason and Connor, this is probably subtly broken on
Mali T720.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11775>
 2021-07-23 15:53:57 +00:00
Icecream95
c246af0dd8 panfrost: Only upload UBOs when needed 
If all of the used values from a UBO are pushed, it doesn't need to be
uploaded.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11700>
 2021-07-03 13:23:29 +00:00
Icecream95
b8a7355c03 pan/mdg: Create a mask of UBOs that need to be uploaded 
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11700>
 2021-07-03 13:23:29 +00:00
Icecream95
08495a948e pan/mdg: Add 16 bytes of padding to the end of shaders 
Fixes INSTR_INVALID_PC faults when a shader ends on a 16MB boundary.

Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11551>
 2021-06-23 14:42:55 +00:00
Boris Brezillon
20b22efdcb pan/midg: Add a flag to dump internal shaders 
Signed-off-by: Boris Brezillon <boris.brezillon@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11139>
 2021-06-22 14:07:33 +00:00
Jason Ekstrand
0afbfee8da nir,panfrost: Suffix fsat_signed and fclamp_pos with _mali 
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11463>
 2021-06-21 09:03:34 -05:00
Icecream95
38e8d7afe3 pan/mdg: Fix reading a spilt register in the bundle it's written 
Read directly from the instruction getting spilt. Otherwise a fill
will be inserted before the spill writing the value, so the
instruction reading the spilt value gets garbage data.

Use the bundle_id to check if the instructions are in the same bundle.

Insert a move instruction, as the spill needs the value in a LD/ST
register such as AL0, while the ALU instruction reading the value
needs it in a work register such as R0.

Cc: mesa-stable
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/4857
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11212>
 2021-06-17 19:53:14 +00:00
Icecream95
31d26ebf1b pan/mdg: Fill from TLS before spilling non-SSA nodes 
Otherwise the data already written to the node will get overwritten.

Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11212>
 2021-06-17 19:53:14 +00:00
Icecream95
ed9a9a09f2 pan/mdg: Reorder some code in mir_spill_register 
Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11212>
 2021-06-17 19:53:14 +00:00
Icecream95
1490e7e622 pan/mdg: Add a bundle ID to instructions 
So that it is possible to check if two instructions were scheduled
into the same bundle.

Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11212>
 2021-06-17 19:53:14 +00:00
Alyssa Rosenzweig
43cff98dff pan/mdg: Insert moves to load/store registers 
Ensures a valid schedule/regalloc is possible when vectors are used in
funny ways, as occurs in dEQP-GLES31 resulting in a scheduler hang (or
with prior patches, an RA failure).

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
5f37474403 pan/mdg: Assert scheduled instructions are reasonable 
Would've got a scheduler hang.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
fa20037895 pan/mdg: Don't skip unit-based checks in choose_instruction 
If an explicit unit isn't specified, we still should check.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
edbdf4f4e7 pan/mdg: Use more accurate ld/st reg estimates 
And assert that we got them right.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
68846ba4a8 pan/mdg: Lower away gl_VertexID offset 
Technically we can stick the offset in the vertex ID attribute record,
but this is a faster way to get the test passing and Midgard perf?
what's that?

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
26baad41f1 pan/mdg: Wire in PAN_SYSVAL_VERTEX_INSTANCE_OFFSETS 
If we're going to advertise the CAP, better not crash..

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
6936298d75 pan/mdg: Fix incorrect rewrite in Midgard scheduler 
Fixes on Midgard
dEQP-GLES31.functional.shaders.builtin_functions.uniform.findLSBMinusOne.highp_fragment

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
def3d52a15 pan/mdg: Update r1.w comment 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
1369d5e43a pan/mdg: Handle {i,u}{add,sub}_sat 
As SATADD with different modifiers.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
bdb32eec9a pan/mdg: Fix units for SUBSAT 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
22a973601b pan/mdg: Stub memory_barrier{_image} 
Same as we do for Bifrost.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
21d06a41f7 pan/mdg: Make -Wswitch happy 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Alyssa Rosenzweig
34c6d105f6 pan/mdg: Use consistent casing in midgard_print 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11123>
 2021-06-10 18:06:10 +00:00
Rhys Perry
1cbcfb8b38 nir, nir/algebraic: add byte/word insertion instructions 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
 2021-06-08 08:57:42 +00:00
Caio Marcelo de Oliveira Filho
c8a7bd0dc8 nir: Rename WORK_GROUP (and similar) to WORKGROUP 
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>
 2021-06-07 22:34:42 +00:00
Caio Marcelo de Oliveira Filho
a71a780598 nir: Rename nir_intrinsic_load_local_group_size to nir_intrinsic_load_workgroup_size 
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>
 2021-06-07 22:34:42 +00:00
Icecream95
f18c55708a pan/mdg: Try scheduling load/store ops in pairs 
If there are an even number of load/store ops to be scheduled, and
only one load/store op is available for scheduling, try using another
instruction type.

Helps bundle count at the cost of register pressure.

total instructions in shared programs: 333405 -> 333599 (0.06%)
instructions in affected programs: 27576 -> 27770 (0.70%)
helped: 43
HURT: 69
helped stats (abs) min: 1 max: 61 x̄: 5.49 x̃: 1
helped stats (rel) min: 0.18% max: 11.71% x̄: 2.27% x̃: 1.75%
HURT stats (abs)   min: 1 max: 95 x̄: 6.23 x̃: 2
HURT stats (rel)   min: 0.06% max: 32.42% x̄: 2.59% x̃: 1.53%
95% mean confidence interval for instructions value: -0.93 4.40
95% mean confidence interval for instructions %-change: -0.09% 1.53%
Inconclusive result (value mean confidence interval includes 0).

total bundles in shared programs: 155785 -> 152371 (-2.19%)
bundles in affected programs: 83689 -> 80275 (-4.08%)
helped: 2538
HURT: 110
helped stats (abs) min: 1 max: 59 x̄: 1.53 x̃: 1
helped stats (rel) min: 0.14% max: 22.52% x̄: 8.71% x̃: 7.69%
HURT stats (abs)   min: 1 max: 92 x̄: 4.32 x̃: 1
HURT stats (rel)   min: 0.21% max: 55.76% x̄: 4.61% x̃: 2.86%
95% mean confidence interval for bundles value: -1.41 -1.17
95% mean confidence interval for bundles %-change: -8.37% -7.94%
Bundles are helped.

total quadwords in shared programs: 264143 -> 260520 (-1.37%)
quadwords in affected programs: 141705 -> 138082 (-2.56%)
helped: 2560
HURT: 96
helped stats (abs) min: 1 max: 15 x̄: 1.49 x̃: 1
helped stats (rel) min: 0.06% max: 14.29% x̄: 5.62% x̃: 5.00%
HURT stats (abs)   min: 1 max: 11 x̄: 2.02 x̃: 2
HURT stats (rel)   min: 0.12% max: 6.20% x̄: 1.94% x̃: 1.47%
95% mean confidence interval for quadwords value: -1.42 -1.31
95% mean confidence interval for quadwords %-change: -5.50% -5.20%
Quadwords are helped.

total registers in shared programs: 21709 -> 22156 (2.06%)
registers in affected programs: 2684 -> 3131 (16.65%)
helped: 55
HURT: 470
helped stats (abs) min: 1 max: 2 x̄: 1.05 x̃: 1
helped stats (rel) min: 6.67% max: 40.00% x̄: 15.37% x̃: 14.29%
HURT stats (abs)   min: 1 max: 4 x̄: 1.07 x̃: 1
HURT stats (rel)   min: 6.67% max: 100.00% x̄: 31.63% x̃: 25.00%
95% mean confidence interval for registers value: 0.79 0.91
95% mean confidence interval for registers %-change: 24.69% 28.72%
Registers are HURT.

total threads in shared programs: 24450 -> 24360 (-0.37%)
threads in affected programs: 234 -> 144 (-38.46%)
helped: 12
HURT: 63
helped stats (abs) min: 1 max: 2 x̄: 1.50 x̃: 1
helped stats (rel) min: 100.00% max: 100.00% x̄: 100.00% x̃: 100.00%
HURT stats (abs)   min: 1 max: 2 x̄: 1.71 x̃: 2
HURT stats (rel)   min: 50.00% max: 50.00% x̄: 50.00% x̃: 50.00%
95% mean confidence interval for threads value: -1.49 -0.91
95% mean confidence interval for threads %-change: -38.74% -13.26%
Threads are [HURT].

total loops in shared programs: 286 -> 286 (0.00%)
loops in affected programs: 0 -> 0
helped: 0
HURT: 0

total spills in shared programs: 521 -> 593 (13.82%)
spills in affected programs: 260 -> 332 (27.69%)
helped: 8
HURT: 9

total fills in shared programs: 1598 -> 1659 (3.82%)
fills in affected programs: 839 -> 900 (7.27%)
helped: 9

HURT: 10
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5745>
 2021-05-24 20:54:37 +00:00
Alyssa Rosenzweig
b131f96aa8 pan/mdg: Set lower_uniforms_to_ubo 
Rip off the band-aid. No other driver in Mesa has this combination, and
there's a reason for that. It confuses nir_to_tgsi; this commit fixes
GALLIUM_HUD (again) on Midgard.

shader-db stats aren't as bad as they appear at first blush, since the
added cycles are from added if-else branches (and only one side of the
if is taken on Midgard, which does no warping*).

total instructions in shared programs: 97036 -> 98107 (1.10%)
instructions in affected programs: 8297 -> 9368 (12.91%)
helped: 0
HURT: 45
HURT stats (abs)   min: 1 max: 52 x̄: 23.80 x̃: 25
HURT stats (rel)   min: 0.61% max: 61.90% x̄: 16.18% x̃: 14.66%
95% mean confidence interval for instructions value: 20.04 27.56
95% mean confidence interval for instructions %-change: 12.62% 19.74%
Instructions are HURT.

total bundles in shared programs: 45507 -> 46091 (1.28%)
bundles in affected programs: 3138 -> 3722 (18.61%)
helped: 2
HURT: 40
helped stats (abs) min: 1 max: 1 x̄: 1.00 x̃: 1
helped stats (rel) min: 1.54% max: 2.17% x̄: 1.86% x̃: 1.86%
HURT stats (abs)   min: 2 max: 44 x̄: 14.65 x̃: 12
HURT stats (rel)   min: 2.70% max: 50.00% x̄: 23.03% x̃: 21.25%
95% mean confidence interval for bundles value: 10.35 17.46
95% mean confidence interval for bundles %-change: 16.81% 26.88%
Bundles are HURT.

total quadwords in shared programs: 76952 -> 77849 (1.17%)
quadwords in affected programs: 6556 -> 7453 (13.68%)
helped: 2
HURT: 44
helped stats (abs) min: 1 max: 5 x̄: 3.00 x̃: 3
helped stats (rel) min: 1.73% max: 4.55% x̄: 3.14% x̃: 3.14%
HURT stats (abs)   min: 2 max: 58 x̄: 20.52 x̃: 18
HURT stats (rel)   min: 2.11% max: 46.34% x̄: 17.20% x̃: 12.96%
95% mean confidence interval for quadwords value: 15.18 23.82
95% mean confidence interval for quadwords %-change: 12.68% 19.96%
Quadwords are HURT.

total registers in shared programs: 6966 -> 6925 (-0.59%)
registers in affected programs: 347 -> 306 (-11.82%)
helped: 26
HURT: 8
helped stats (abs) min: 1 max: 4 x̄: 2.04 x̃: 2
helped stats (rel) min: 6.67% max: 42.86% x̄: 20.92% x̃: 22.22%
HURT stats (abs)   min: 1 max: 5 x̄: 1.50 x̃: 1
HURT stats (rel)   min: 9.09% max: 50.00% x̄: 17.19% x̃: 11.81%
95% mean confidence interval for registers value: -1.85 -0.56
95% mean confidence interval for registers %-change: -18.97% -4.93%
Registers are helped.

total threads in shared programs: 5040 -> 5050 (0.20%)
threads in affected programs: 13 -> 23 (76.92%)
helped: 10
HURT: 1
helped stats (abs) min: 1 max: 2 x̄: 1.10 x̃: 1
helped stats (rel) min: 100.00% max: 100.00% x̄: 100.00% x̃: 100.00%
HURT stats (abs)   min: 1 max: 1 x̄: 1.00 x̃: 1
HURT stats (rel)   min: 50.00% max: 50.00% x̄: 50.00% x̃: 50.00%
95% mean confidence interval for threads value: 0.44 1.38
95% mean confidence interval for threads %-change: 55.98% 116.75%
Threads are helped.

Fixes: 24d7c413fe ("panfrost: Enable packed uniforms.")
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10866>
 2021-05-18 19:19:01 +00:00
Alyssa Rosenzweig
a89bc59980 pan/mdg: Fix spills to TLS 
LOCAL_STORAGE.zw is workgroup local memory, whereas LOCAL_STORAGE.xy is
thread local memory. Likewise PC_SP.zw is the stack pointer, which is
initialized to (LOCAL_STORAGE.zw + offset) but is modifiable by the
shader. Panfrost doesn't modify the s tack pointer, and the register
allocation logic assumes a zero offset, so let's always spill to thread
local memory = LOCAL_STORAGE.xy, as was intended by Italo's cleanup.

This is visible on any shader that spills. Compute shaders aren't
advertised yet, so WLS will be null, causing a fault like the following
(reproduced on Mali T860 with the glyphy trace):

[15634.148873] panfrost ff9a0000.gpu: Unhandled Page fault in AS0 at VA 0x0000000000000000
          Reason: TODO
          raw fault status: 0x70003C2
          decoded fault status: SLAVE FAULT
          exception type 0xC2: TRANSLATION_FAULT_LEVEL2
          access type 0x3: WRITE
          source id 0x700
[15634.658170] panfrost ff9a0000.gpu: gpu sched timeout, js=0, config=0x3300, status=0x8, head=0x31d4540,
tail=0x31d4540, sched_job=00000000e8101b2e

Fixes: 6a12ea02fe ("pan/mdg: properly encode/decode ldst instructions")
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10866>
 2021-05-18 19:19:01 +00:00
Alyssa Rosenzweig
8d2d711026 pan/mdg: Fix output types for scalar fields 
Already fixed vector, but scalar was missed.

Fixes: 4d9c0a32e7 ("pan/mdg: Use _output_ type for outmod printing")
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10866>
 2021-05-18 19:19:01 +00:00
Alyssa Rosenzweig
6f2ea57e1d pan/mdg: Remove unused midgard_int_alu_op_prefix 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10866>
 2021-05-18 19:19:01 +00:00
Emma Anholt
958f11d537 midgard: Fix type for vertex_builtin_arg() and compute_builtin_arg(). 
It takes an intrinsic, not an ALU op.  Fixes a clang complaint about enum
conversion.

Fixes: 306800d747 ("pan/midgard: Lower gl_VertexID/gl_InstanceID to attributes")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10843>
 2021-05-17 17:58:49 +00:00
Alyssa Rosenzweig
01ef56a7e4 pan/mdg: Use smaller LD_UNIFORM instructions 
If we only read 8 bytes from a UBO, we need to use LD_UNIFORM.64 as
opposed to LD_UNIFORM.128. In addition to probably being faster, this
fixes a buffer overrun manifesting as MMU faults with source ID
0x500/0x600/0x700, visible in WebGL Aquarium.

This is essentially the same patch as 616394cf31 ("pan/mdg: Use
appropriate sizes for global loads/stores"), only this is for UBOs where
that was for SSBOs.

Before enabling PACKED_UNIFORMS, this bug was not visible since we could
guarantee the UBO size was a multiple of 16. We no longer have that
invariant, and in rare cases the last 8 bytes of the last 16-byte slot
of a mapped uniform buffer would overrun the BO and trigger a fault,
even if the result is unused.

Fixes: 24d7c413fe ("panfrost: Enable packed uniforms.")
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10772>
 2021-05-12 20:04:21 +00:00
Tomeu Vizoso
e882b0505e pan/midgard: Don't emit zero padding 
util_dynarray_grow_bytes doesn't cope well with zero grow bytes.

../mesa/src/util/u_dynarray.h:134:8: runtime error: division by zero
    #0 0xffffa2d99288 in util_dynarray_grow_bytes ../mesa/src/util/u_dynarray.h:134
    #1 0xffffa2d99288 in emit_alu_bundle ../mesa/src/panfrost/midgard/midgard_emit.c:350
    #2 0xffffa2d99288 in emit_binary_bundle ../mesa/src/panfrost/midgard/midgard_emit.c:390
    #3 0xffffa2d738a8 in midgard_compile_shader_nir ../mesa/src/panfrost/midgard/midgard_compile.c:2740
    #4 0xffffa2cc8aa8 in panfrost_shader_compile ../mesa/src/gallium/drivers/panfrost/pan_assemble.c:68
    #5 0xffffa2c9918c in panfrost_bind_shader_state ../mesa/src/gallium/drivers/panfrost/pan_context.c:2015
    #6 0xffffa0726d94 in st_update_vp ../mesa/src/mesa/state_tracker/st_atom_shader.c:238
    #7 0xffffa070e028 in st_validate_state ../mesa/src/mesa/state_tracker/st_atom.c:261
    #8 0xffffa06c30bc in prepare_draw ../mesa/src/mesa/state_tracker/st_draw.c:132
    #9 0xffffa06c30bc in st_draw_vbo ../mesa/src/mesa/state_tracker/st_draw.c:184
    #10 0xffffa167bb20 in _mesa_validated_drawrangeelements ../mesa/src/mesa/main/draw.c:816
    #11 0xffffa167c684 in _mesa_DrawElements ../mesa/src/mesa/main/draw.c:970
    #12 0xaaaadfa5b55c in glu::drawFromUserPointers(glu::RenderContext const&, unsigned int, int, glu::VertexArrayBinding const*, glu::PrimitiveList const&, glu::DrawUtilCallback*) (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x3b355c)
    #13 0xaaaadf9ce514 in deqp::gls::ShaderLibraryCase::execute() (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x326514)
    #14 0xaaaadf9cfb38 in deqp::gls::ShaderLibraryCase::iterate() (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x327b38)
    #15 0xaaaadf70cc9c in deqp::gles2::TestCaseWrapper::iterate(tcu::TestCase*) (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x64c9c)
    #16 0xaaaadfb9a47c in tcu::TestSessionExecutor::iterateTestCase(tcu::TestCase*) (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x4f247c)
    #17 0xaaaadfb9b240 in tcu::TestSessionExecutor::iterate() (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x4f3240)
    #18 0xaaaadfb7a564 in tcu::App::iterate() (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x4d2564)
    #19 0xaaaadf70a20c in main (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x6220c)
    #20 0xffffaf1e12a8 in __libc_start_main ../csu/libc-start.c:308
    #21 0xaaaadf70a8dc  (/home/tomeu/deqp-build/modules/gles2/deqp-gles2+0x628dc)

Signed-off-by: Tomeu Vizoso <tomeu.vizoso@collabora.com>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10695>
 2021-05-10 17:02:09 +00:00
Alyssa Rosenzweig
54046d61f8 pan/mdg: Model blend shader interference 
Backport of 4439757db2 ("pan/bi: Use the interference mechanism
to describe blend shader reg use") to Midgard.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9239>
 2021-05-06 23:26:21 +00:00
Alyssa Rosenzweig
1378c67bcf panfrost/blend: Inline blend constants 
If we're going to key them in NIR, we might as well get the benefit of
constant folding them too.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10652>
 2021-05-06 16:40:50 +00:00
Alyssa Rosenzweig
ba39367b96 pan/mdg: Enable nir_opt_{move, sink} 
I felt bad about the last patch regressing Midgard perf, so here's some
moar Midgard perf for you ^^

total instructions in shared programs: 97089 -> 97036 (-0.05%)
instructions in affected programs: 5230 -> 5177 (-1.01%)
helped: 53
HURT: 31
helped stats (abs) min: 1 max: 17 x̄: 4.40 x̃: 6
helped stats (rel) min: 0.61% max: 12.24% x̄: 7.74% x̃: 11.54%
HURT stats (abs)   min: 1 max: 8 x̄: 5.81 x̃: 8
HURT stats (rel)   min: 1.08% max: 13.79% x̄: 9.69% x̃: 11.11%
95% mean confidence interval for instructions value: -1.89 0.63
95% mean confidence interval for instructions %-change: -3.41% 0.80%
Inconclusive result (value mean confidence interval includes 0).

total bundles in shared programs: 45612 -> 45507 (-0.23%)
bundles in affected programs: 17331 -> 17226 (-0.61%)
helped: 139
HURT: 166
helped stats (abs) min: 1 max: 21 x̄: 3.76 x̃: 2
helped stats (rel) min: 0.85% max: 18.37% x̄: 6.38% x̃: 4.55%
HURT stats (abs)   min: 1 max: 10 x̄: 2.51 x̃: 1
HURT stats (rel)   min: 0.79% max: 31.25% x̄: 7.54% x̃: 4.55%
95% mean confidence interval for bundles value: -0.90 0.21
95% mean confidence interval for bundles %-change: 0.05% 2.34%
Inconclusive result (value mean confidence interval includes 0).

total quadwords in shared programs: 77275 -> 76952 (-0.42%)
quadwords in affected programs: 32314 -> 31991 (-1.00%)
helped: 142
HURT: 179
helped stats (abs) min: 1 max: 28 x̄: 4.38 x̃: 2
helped stats (rel) min: 0.34% max: 13.79% x̄: 4.29% x̃: 2.78%
HURT stats (abs)   min: 1 max: 6 x̄: 1.67 x̃: 2
HURT stats (rel)   min: 0.44% max: 16.67% x̄: 2.93% x̃: 2.63%
95% mean confidence interval for quadwords value: -1.56 -0.45
95% mean confidence interval for quadwords %-change: -0.78% 0.25%
Inconclusive result (%-change mean confidence interval includes 0).

total registers in shared programs: 7081 -> 6771 (-4.38%)
registers in affected programs: 2217 -> 1907 (-13.98%)
helped: 193
HURT: 75
helped stats (abs) min: 1 max: 6 x̄: 2.04 x̃: 1
helped stats (rel) min: 6.25% max: 62.50% x̄: 24.32% x̃: 20.00%
HURT stats (abs)   min: 1 max: 3 x̄: 1.11 x̃: 1
HURT stats (rel)   min: 7.14% max: 50.00% x̄: 17.17% x̃: 14.29%
95% mean confidence interval for registers value: -1.37 -0.94
95% mean confidence interval for registers %-change: -15.53% -9.89%
Registers are helped.

total threads in shared programs: 5036 -> 5152 (2.30%)
threads in affected programs: 185 -> 301 (62.70%)
helped: 93
HURT: 19
helped stats (abs) min: 1 max: 2 x̄: 1.49 x̃: 1
helped stats (rel) min: 100.00% max: 100.00% x̄: 100.00% x̃: 100.00%
HURT stats (abs)   min: 1 max: 2 x̄: 1.21 x̃: 1
HURT stats (rel)   min: 50.00% max: 50.00% x̄: 50.00% x̃: 50.00%
95% mean confidence interval for threads value: 0.82 1.25
95% mean confidence interval for threads %-change: 63.96% 85.14%
Threads are helped.

total loops in shared programs: 19 -> 19 (0.00%)
loops in affected programs: 0 -> 0
helped: 0
HURT: 0

total spills in shared programs: 2 -> 0
spills in affected programs: 2 -> 0
helped: 1
HURT: 0

total fills in shared programs: 15 -> 0
fills in affected programs: 15 -> 0
helped: 1
HURT: 0

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10652>
 2021-05-06 16:40:50 +00:00
Alyssa Rosenzweig
ad6e53da5c pan/mdg: Remove midgard_opt_copy_prop_reg 
This is known broken code, and the fix is highly nontrivial. It isn't
doing terribly much for perf, so just rip off the band-aid. Prevents a
dEQP regression, and likely fixes bugs in real apps too.

total instructions in shared programs: 96640 -> 97089 (0.46%)
instructions in affected programs: 27831 -> 28280 (1.61%)
helped: 33
HURT: 301
helped stats (abs) min: 1 max: 6 x̄: 3.64 x̃: 5
helped stats (rel) min: 1.96% max: 10.00% x̄: 6.48% x̃: 7.94%
HURT stats (abs)   min: 1 max: 18 x̄: 1.89 x̃: 1
HURT stats (rel)   min: 0.46% max: 15.00% x̄: 3.17% x̃: 2.38%
95% mean confidence interval for instructions value: 1.09 1.59
95% mean confidence interval for instructions %-change: 1.80% 2.63%
Instructions are HURT.

total bundles in shared programs: 45615 -> 45612 (<.01%)
bundles in affected programs: 11257 -> 11254 (-0.03%)
helped: 121
HURT: 146
helped stats (abs) min: 1 max: 7 x̄: 2.34 x̃: 1
helped stats (rel) min: 1.22% max: 23.33% x̄: 7.85% x̃: 5.26%
HURT stats (abs)   min: 1 max: 17 x̄: 1.92 x̃: 2
HURT stats (rel)   min: 0.42% max: 25.00% x̄: 5.17% x̃: 3.85%
95% mean confidence interval for bundles value: -0.34 0.31
95% mean confidence interval for bundles %-change: -1.69% 0.23%
Inconclusive result (value mean confidence interval includes 0).

total quadwords in shared programs: 76662 -> 77275 (0.80%)
quadwords in affected programs: 20148 -> 20761 (3.04%)
helped: 28
HURT: 275
helped stats (abs) min: 1 max: 4 x̄: 1.54 x̃: 1
helped stats (rel) min: 0.43% max: 25.00% x̄: 4.89% x̃: 2.50%
HURT stats (abs)   min: 1 max: 12 x̄: 2.39 x̃: 2
HURT stats (rel)   min: 0.51% max: 28.57% x̄: 5.18% x̃: 4.26%
95% mean confidence interval for quadwords value: 1.80 2.25
95% mean confidence interval for quadwords %-change: 3.64% 4.86%
Quadwords are HURT.

total registers in shared programs: 7078 -> 7081 (0.04%)
registers in affected programs: 1028 -> 1031 (0.29%)
helped: 62
HURT: 70
helped stats (abs) min: 1 max: 2 x̄: 1.11 x̃: 1
helped stats (rel) min: 8.33% max: 50.00% x̄: 15.03% x̃: 12.50%
HURT stats (abs)   min: 1 max: 2 x̄: 1.03 x̃: 1
HURT stats (rel)   min: 8.33% max: 66.67% x̄: 20.13% x̃: 11.25%
95% mean confidence interval for registers value: -0.17 0.21
95% mean confidence interval for registers %-change: -0.14% 7.38%
Inconclusive result (value mean confidence interval includes 0).

total threads in shared programs: 5032 -> 5036 (0.08%)
threads in affected programs: 31 -> 35 (12.90%)
helped: 12
HURT: 6
helped stats (abs) min: 1 max: 2 x̄: 1.08 x̃: 1
helped stats (rel) min: 100.00% max: 100.00% x̄: 100.00% x̃: 100.00%
HURT stats (abs)   min: 1 max: 2 x̄: 1.50 x̃: 1
HURT stats (rel)   min: 50.00% max: 50.00% x̄: 50.00% x̃: 50.00%
95% mean confidence interval for threads value: -0.43 0.87
95% mean confidence interval for threads %-change: 13.82% 86.18%
Inconclusive result (value mean confidence interval includes 0).

total loops in shared programs: 19 -> 19 (0.00%)
loops in affected programs: 0 -> 0
helped: 0
HURT: 0

total spills in shared programs: 0 -> 2
spills in affected programs: 0 -> 2
helped: 0
HURT: 1

total fills in shared programs: 0 -> 15
fills in affected programs: 0 -> 15
helped: 0
HURT: 1

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10652>
 2021-05-06 16:40:50 +00:00
Alyssa Rosenzweig
4d9c0a32e7 pan/mdg: Use _output_ type for outmod printing 
Fixes incorrect outmods printed for conversions.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10652>
 2021-05-06 16:40:50 +00:00
Alyssa Rosenzweig
a9621c4493 pan/mdg: Don't print zero 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
cb10a8e2f1 pan/mdg: Reduced printed parens 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
0b45d4b6b5 pan/mdg: Don't print mem addr brackets 
Already comma separated.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
d0fc23b7bd pan/mdg: Don't print explicit .rte 
Default round mode.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
9bd3ebf829 pan/mdg: Suppress most attribute tables 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
07740e78b4 pan/mdg: Don't print zero shifts 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
fa68c8bca9 pan/mdg: More concise RMU name 
Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
7752b09ade pan/mdg: Hide units behind MIDGARD_MESA_DEBUG=verbose 
Not usually interesting unless debugging bundling.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
03e3e65cd4 panfrost: Assume lower_fragcolor has been called 
Allows us to clean up quite a bit.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10393>
 2021-05-04 20:04:03 +00:00
Alyssa Rosenzweig
0ec27d02e1 panfrost: Don't unroll loops in GLSL 
GLSL loop analysis is trouble. Just use NIR.

Signed-off-by: Alyssa Rosenzweig <alyssa@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10392>
 2021-05-03 15:10:20 +00:00
Icecream95
f85b7aa5d4 pan/mdg: Fix calculation of available work registers 
Make the rmu variable signed; otherwise the MAX2 has no effect and
work_count can end up being larger than 16.

Fixes INSTR_OPERAND_FAULTs in SuperTuxKart.

Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/4707
Fixes: c6ed8bf77c ("panfrost: Fix uniform_count on Midgard")
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10507>
 2021-04-30 10:38:35 +00:00