fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-24 15:20:10 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

aco/gfx12: implement workgroup barrier

Browse source 
Same sequence LLVM uses for llvm.amdgcn.s.barrier.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29330>
This commit is contained in:
Rhys Perry 2024-05-16 16:54:49 +01:00 • committed by Marge Bot
parent fae2a85d57
commit ae18c88409
2 changed files with 11 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
src/amd/compiler/aco_lower_to_hw_instr.cpp
View file

@ -2891,8 +2891,12 @@ lower_to_hw_instr(Program* program)
program->workgroup_size > program->wave_size;
bld.insert(std::move(instr));
if (emit_s_barrier)
if (emit_s_barrier && ctx.program->gfx_level >= GFX12) {
bld.sop1(aco_opcode::s_barrier_signal, Operand::c32(-1));
bld.sopp(aco_opcode::s_barrier_wait, UINT16_MAX);
} else if (emit_s_barrier) {
bld.sopp(aco_opcode::s_barrier);
}
} else if (instr->opcode == aco_opcode::p_cvt_f16_f32_rtne) {
float_mode new_mode = block->fp_mode;
new_mode.round16_64 = fp_round_ne;

6
src/amd/compiler/aco_scheduler_ilp.cpp
View file

@ -97,6 +97,12 @@ can_reorder(const Instruction* const instr)
case aco_opcode::s_set_gpr_idx_idx:
case aco_opcode::s_sendmsg_rtn_b32:
case aco_opcode::s_sendmsg_rtn_b64:
case aco_opcode::s_barrier_signal:
case aco_opcode::s_barrier_signal_isfirst:
case aco_opcode::s_get_barrier_state:
case aco_opcode::s_barrier_init:
case aco_opcode::s_barrier_join:
case aco_opcode::s_wakeup_barrier:
/* SOPK */
case aco_opcode::s_cbranch_i_fork:
case aco_opcode::s_getreg_b32: