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)