fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-27 06:10:13 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

agx: Handle group_memory_barrier

Browse source 
A combination of control_barrier + memory_barrier but it's always seen with
those. This would be safer with scoped barriers...

Fixes dEQP-GLES31.functional.synchronization.inter_invocation.ssbo

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/21326>
This commit is contained in:
Alyssa Rosenzweig 2023-02-13 09:00:50 -05:00 • committed by Marge Bot
parent e9cec96633
commit e93a221024
1 changed files with 1 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/asahi/compiler/agx_compile.c
View file

@ -956,6 +956,7 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
case nir_intrinsic_control_barrier:
return agx_threadgroup_barrier(b);
case nir_intrinsic_group_memory_barrier:
case nir_intrinsic_memory_barrier_shared:
/* Always seen with a control_barrier */
return NULL;