agx: Implement image barriers

Or cache flushes or whatever these actually are. Probably could be optimized
once we understand what the 4 individual instructions are actually doing. Fixes
dEQP-GLES31.functional.image_load_store.2d.qualifiers.*.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24258>

diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c
index 0230268..943637c 100644 (file)
--- a/src/asahi/compiler/agx_compile.c
+++ b/src/asahi/compiler/agx_compile.c
@@ -970,6 +970,7 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
       assert(!b->shader->is_preamble && "invalid");
 
       bool needs_threadgroup_barrier = false;
+      bool needs_image_barriers = false;
 
       if (nir_intrinsic_execution_scope(instr) != SCOPE_NONE) {
          assert(nir_intrinsic_execution_scope(instr) > SCOPE_SUBGROUP &&
@@ -987,6 +988,12 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
          if (modes & nir_var_mem_shared)
             needs_threadgroup_barrier = true;
 
+         if (modes & nir_var_image) {
+            agx_image_barrier_1(b);
+            agx_image_barrier_2(b);
+            needs_image_barriers = true;
+         }
+
          if (nir_intrinsic_memory_scope(instr) >= SCOPE_WORKGROUP)
             needs_threadgroup_barrier = true;
       }
@@ -994,6 +1001,11 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
       if (needs_threadgroup_barrier)
          agx_threadgroup_barrier(b);
 
+      if (needs_image_barriers) {
+         agx_image_barrier_3(b);
+         agx_image_barrier_4(b);
+      }
+
       return NULL;
    }
 

diff --git a/src/asahi/compiler/agx_opcodes.py b/src/asahi/compiler/agx_opcodes.py
index 5c69152..12e8a10 100644 (file)
--- a/src/asahi/compiler/agx_opcodes.py
+++ b/src/asahi/compiler/agx_opcodes.py
@@ -365,8 +365,18 @@ op("block_image_store", (0xB1, 0xFF, 10, _), dests = 0, srcs = 2,
 # Barriers
 op("threadgroup_barrier", (0x0068, 0xFFFF, 2, _), dests = 0, srcs = 0,
    can_eliminate = False)
-op("memory_barrier", (0x96F5, 0xFFFF, 2, _), dests = 0, srcs = 0,
-   can_eliminate = False)
+
+def memory_barrier(name, a, b, c):
+    op(name, (0xF5 | (a << 10) | (b << 8) | (c << 12), 0xFFFF, 2, _), dests = 0, srcs = 0,
+       can_eliminate = False)
+
+memory_barrier("memory_barrier", 1, 2, 9)
+
+# TODO: Not clear what these individually are. Some might be cache flushes?
+memory_barrier("image_barrier_1", 2, 2, 10)
+memory_barrier("image_barrier_2", 3, 2, 10)
+memory_barrier("image_barrier_3", 2, 1, 10)
+memory_barrier("image_barrier_4", 3, 1, 10)
 
 # Convenient aliases.
 op("mov", _, srcs = 1)