fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 05:18:08 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 2

aco: consider position/primitive exports around memory barriers

Browse source 
This is needed to create barriers which ensure stores finish before
position/primitive exports.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/21624>
This commit is contained in:
Rhys Perry 2023-02-28 17:44:01 +00:00 • committed by Marge Bot
parent dfb6d3e443
commit 417990b19e
2 changed files with 13 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
src/amd/compiler/aco_ir.h
View file

@ -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 */

11
src/amd/compiler/aco_scheduler.cpp
View file

@ -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)