From 76641762cee4d6997a5bed1b5bf6d1f6f0e10b22 Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Mon, 29 May 2023 19:30:16 -0400 Subject: [PATCH] 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 Part-of: --- src/asahi/compiler/agx_compile.c | 12 ++++++++++++ src/asahi/compiler/agx_opcodes.py | 14 ++++++++++++-- 2 files changed, 24 insertions(+), 2 deletions(-) diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c index 02302683a2a..943637ce859 100644 --- 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 5c6915231df..12e8a105a1d 100644 --- 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)