From f857795e833f6a99056db849f87d2a25c9b4b5f3 Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Sat, 3 Dec 2022 14:34:44 -0500 Subject: [PATCH] agx: Implement barriers Signed-off-by: Alyssa Rosenzweig Part-of: --- src/asahi/compiler/agx_compile.c | 10 ++++++++++ src/asahi/compiler/agx_opcodes.py | 6 ++++++ 2 files changed, 16 insertions(+) diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c index 69b39cc3b38..560b1ffa469 100644 --- a/src/asahi/compiler/agx_compile.c +++ b/src/asahi/compiler/agx_compile.c @@ -776,6 +776,16 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr) return agx_load_compute_dimension( b, dst, instr, AGX_SR_THREAD_POSITION_IN_THREADGROUP_X); + case nir_intrinsic_memory_barrier_buffer: + return agx_memory_barrier(b); + + case nir_intrinsic_control_barrier: + return agx_threadgroup_barrier(b); + + case nir_intrinsic_memory_barrier_shared: + /* Always seen with a control_barrier */ + return NULL; + default: fprintf(stderr, "Unhandled intrinsic %s\n", nir_intrinsic_infos[instr->intrinsic].name); diff --git a/src/asahi/compiler/agx_opcodes.py b/src/asahi/compiler/agx_opcodes.py index f86602ed451..1e138978020 100644 --- a/src/asahi/compiler/agx_opcodes.py +++ b/src/asahi/compiler/agx_opcodes.py @@ -308,6 +308,12 @@ op("writeout", (0x48, 0xFF, 4, _), dests = 0, imms = [WRITEOUT], can_eliminate = op("block_image_store", (0xB1, 0xFF, 10, _), dests = 0, srcs = 2, imms = [FORMAT, DIM], can_eliminate = False) +# 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) + # Convenient aliases. op("mov", _, srcs = 1) op("not", _, srcs = 1)