diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h
index b2365c0a758..eb58a6da9c2 100644
--- a/src/amd/compiler/aco_ir.h
+++ b/src/amd/compiler/aco_ir.h
@@ -152,11 +152,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 before the barrier */
+    * atomics/control_barriers/sendmsg_gs_done/position-primitive-export 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 after the barrier */
+    * atomics/control_barriers/sendmsg_gs_done/position-primitive-export after the barrier */
    semantic_release = 0x2,
 
    /* the rest are for load/stores/atomics only */
diff --git a/src/amd/compiler/aco_scheduler.cpp b/src/amd/compiler/aco_scheduler.cpp
index 3531c0562da..5316878840d 100644
--- a/src/amd/compiler/aco_scheduler.cpp
+++ b/src/amd/compiler/aco_scheduler.cpp
@@ -428,6 +428,16 @@ is_done_sendmsg(amd_gfx_level gfx_level, const Instruction* instr)
    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)
 {
@@ -483,6 +493,7 @@ add_memory_event(amd_gfx_level gfx_level, memory_event_set* set, Instruction* in
                  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)