From c1afe26be626d2a798a787584c56eabbb3b2d8ce Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Wed, 7 Jun 2023 16:53:11 -0400 Subject: [PATCH] agx: Don't emit silly barriers Trust in the scoped_barrier. Signed-off-by: Alyssa Rosenzweig Part-of: --- src/asahi/compiler/agx_compile.c | 21 ++++++--------------- 1 file changed, 6 insertions(+), 15 deletions(-) diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c index 168fc12808c..c95589066a4 100644 --- a/src/asahi/compiler/agx_compile.c +++ b/src/asahi/compiler/agx_compile.c @@ -1073,37 +1073,28 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr) case nir_intrinsic_scoped_barrier: { assert(!b->shader->is_preamble && "invalid"); - bool needs_threadgroup_barrier = false; bool needs_image_barriers = false; - if (nir_intrinsic_execution_scope(instr) != SCOPE_NONE) { - assert(nir_intrinsic_execution_scope(instr) > SCOPE_SUBGROUP && - "todo: subgroup barriers"); - - needs_threadgroup_barrier = true; - } - if (nir_intrinsic_memory_scope(instr) != SCOPE_NONE) { nir_variable_mode modes = nir_intrinsic_memory_modes(instr); if (modes & (nir_var_mem_global | nir_var_image)) agx_memory_barrier(b); - if (modes & nir_var_mem_shared) - needs_threadgroup_barrier = true; - if (modes & nir_var_image) { agx_image_barrier_1(b); agx_image_barrier_2(b); needs_image_barriers = true; } - - if (nir_intrinsic_memory_scope(instr) >= SCOPE_WORKGROUP) - needs_threadgroup_barrier = true; } - if (needs_threadgroup_barrier) + if (nir_intrinsic_execution_scope(instr) != SCOPE_NONE) { + assert(nir_intrinsic_execution_scope(instr) > SCOPE_SUBGROUP && + "todo: subgroup barriers"); + assert(gl_shader_stage_is_compute(b->shader->nir->info.stage)); + agx_threadgroup_barrier(b); + } if (needs_image_barriers) { agx_image_barrier_3(b);