diff --git a/src/amd/compiler/aco_ir.cpp b/src/amd/compiler/aco_ir.cpp
index 5d65d9e8d42..509c748f15f 100644
--- a/src/amd/compiler/aco_ir.cpp
+++ b/src/amd/compiler/aco_ir.cpp
@@ -244,6 +244,46 @@ is_wait_export_ready(amd_gfx_level gfx_level, const Instruction* instr)
                               : !(instr->salu().imm & wait_event_imm_dont_wait_export_ready_gfx11));
 }
 
+static bool
+is_done_sendmsg(amd_gfx_level gfx_level, const Instruction* instr)
+{
+   if (gfx_level <= GFX10_3 && instr->opcode == aco_opcode::s_sendmsg)
+      return (instr->salu().imm & sendmsg_id_mask) == sendmsg_gs_done;
+   return false;
+}
+
+static bool
+is_pos_prim_export(amd_gfx_level gfx_level, const Instruction* instr)
+{
+   /* Because of NO_PC_EXPORT=1, a done=1 position or primitive export can launch PS waves before
+    * the NGG/VS wave finishes if there are no parameter exports.
+    */
+   return gfx_level >= GFX10 && instr->opcode == aco_opcode::exp &&
+          instr->exp().dest >= V_008DFC_SQ_EXP_POS && instr->exp().dest <= V_008DFC_SQ_EXP_PRIM;
+}
+
+uint16_t
+is_atomic_or_control_instr(amd_gfx_level gfx_level, const Instruction* instr, memory_sync_info sync,
+                           unsigned semantic)
+{
+   bool is_acquire = semantic & semantic_acquire;
+   bool is_release = semantic & semantic_release;
+
+   bool is_atomic = sync.semantics & semantic_atomic;
+   // TODO: NIR doesn't have any atomic load/store, so we assume any load/store is atomic
+   is_atomic |= !(sync.semantics & semantic_private) && sync.storage;
+   if (is_atomic) {
+      bool is_load = !instr->definitions.empty() || (sync.semantics & semantic_rmw);
+      bool is_store = instr->definitions.empty() || (sync.semantics & semantic_rmw);
+      return ((is_release && is_store) || (is_acquire && is_load)) ? sync.storage : 0;
+   }
+
+   uint16_t cls = BITFIELD_MASK(storage_count);
+   if (is_release && (is_done_sendmsg(gfx_level, instr) || is_pos_prim_export(gfx_level, instr)))
+      return cls & ~storage_shared;
+   return (instr->isBarrier() && instr->barrier().exec_scope > scope_invocation) ? cls : 0;
+}
+
 memory_sync_info
 get_sync_info(const Instruction* instr)
 {
diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h
index 12f965271c6..5800bf9000e 100644
--- a/src/amd/compiler/aco_ir.h
+++ b/src/amd/compiler/aco_ir.h
@@ -64,11 +64,11 @@ enum memory_semantics : uint8_t {
    semantic_none = 0x0,
    /* for loads: don't move any access after this load to before this load (even other loads)
     * for barriers: don't move any access after the barrier to before any
-    * atomics/control_barriers/sendmsg_gs_done/position-primitive-export before the barrier */
+    * atomic_loads/control_barriers before the barrier */
    semantic_acquire = 0x1,
    /* for stores: don't move any access before this store to after this store
     * for barriers: don't move any access before the barrier to after any
-    * atomics/control_barriers/sendmsg_gs_done/position-primitive-export after the barrier */
+    * atomic_stores/control_barriers/sendmsg_gs_done/position-primitive-export after the barrier */
    semantic_release = 0x2,
 
    /* the rest are for load/stores/atomics only */
@@ -1876,6 +1876,10 @@ is_phi(aco_ptr<Instruction>& instr)
 }
 
 bool is_wait_export_ready(amd_gfx_level gfx_level, const Instruction* instr);
+
+uint16_t is_atomic_or_control_instr(amd_gfx_level gfx_level, const Instruction* instr,
+                                    memory_sync_info sync, unsigned semantic);
+
 memory_sync_info get_sync_info(const Instruction* instr);
 
 inline bool
diff --git a/src/amd/compiler/aco_scheduler.cpp b/src/amd/compiler/aco_scheduler.cpp
index c25fa74b8bc..fa03cb526eb 100644
--- a/src/amd/compiler/aco_scheduler.cpp
+++ b/src/amd/compiler/aco_scheduler.cpp
@@ -461,24 +461,6 @@ MoveState::upwards_skip(UpwardsCursor& cursor)
    cursor.verify_invariants(block);
 }
 
-bool
-is_done_sendmsg(amd_gfx_level gfx_level, const Instruction* instr)
-{
-   if (gfx_level <= GFX10_3 && instr->opcode == aco_opcode::s_sendmsg)
-      return (instr->salu().imm & sendmsg_id_mask) == sendmsg_gs_done;
-   return false;
-}
-
-bool
-is_pos_prim_export(amd_gfx_level gfx_level, const Instruction* instr)
-{
-   /* Because of NO_PC_EXPORT=1, a done=1 position or primitive export can launch PS waves before
-    * the NGG/VS wave finishes if there are no parameter exports.
-    */
-   return instr->opcode == aco_opcode::exp && instr->exp().dest >= V_008DFC_SQ_EXP_POS &&
-          instr->exp().dest <= V_008DFC_SQ_EXP_PRIM && gfx_level >= GFX10;
-}
-
 memory_sync_info
 get_sync_info_with_hack(const Instruction* instr)
 {
@@ -533,8 +515,6 @@ void
 add_memory_event(amd_gfx_level gfx_level, memory_event_set* set, Instruction* instr,
                  memory_sync_info* sync)
 {
-   set->has_control_barrier |= is_done_sendmsg(gfx_level, instr);
-   set->has_control_barrier |= is_pos_prim_export(gfx_level, instr);
    if (instr->opcode == aco_opcode::p_barrier) {
       Pseudo_barrier_instruction& bar = instr->barrier();
       if (bar.sync.semantics & semantic_acquire)
@@ -542,12 +522,14 @@ add_memory_event(amd_gfx_level gfx_level, memory_event_set* set, Instruction* in
       if (bar.sync.semantics & semantic_release)
          set->bar_release |= bar.sync.storage;
       set->bar_classes |= bar.sync.storage;
-
-      set->has_control_barrier |= bar.exec_scope > scope_invocation;
    }
 
-   if (!sync->storage)
+   if (!sync->storage) {
+      set->has_control_barrier |=
+         is_atomic_or_control_instr(gfx_level, instr, *sync, semantic_acquire | semantic_release) !=
+         0;
       return;
+   }
 
    if (sync->semantics & semantic_acquire)
       set->access_acquire |= sync->storage;