From e93a221024d40a32ab41ae354a3775e0a10953a3 Mon Sep 17 00:00:00 2001
From: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Date: Mon, 13 Feb 2023 09:00:50 -0500
Subject: [PATCH] agx: Handle group_memory_barrier

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>
---
 src/asahi/compiler/agx_compile.c | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c
index fddfe2056e5..7b135a297be 100644
--- a/src/asahi/compiler/agx_compile.c
+++ b/src/asahi/compiler/agx_compile.c
@@ -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;