From 417990b19eca6dfc46657a7a12af32031891828d Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Tue, 28 Feb 2023 17:44:01 +0000 Subject: [PATCH] aco: consider position/primitive exports around memory barriers MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This is needed to create barriers which ensure stores finish before position/primitive exports. Signed-off-by: Rhys Perry Reviewed-by: Timur Kristóf Part-of: --- src/amd/compiler/aco_ir.h | 4 ++-- src/amd/compiler/aco_scheduler.cpp | 11 +++++++++++ 2 files changed, 13 insertions(+), 2 deletions(-) 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)