OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

agx: Don't emit silly barriers

Browse Source 
Trust in the scoped_barrier.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24258>
This commit is contained in:
Alyssa Rosenzweig
2023-06-07 16:53:11 -04:00
committed by Marge Bot
parent b618ba9330
commit c1afe26be6
1 changed files with 6 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

21
src/asahi/compiler/agx_compile.c
View File

@@ -1073,37 +1073,28 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
case nir_intrinsic_scoped_barrier: { case nir_intrinsic_scoped_barrier: {
assert(!b->shader->is_preamble && "invalid"); assert(!b->shader->is_preamble && "invalid");
bool needs_threadgroup_barrier = false;
bool needs_image_barriers = 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) { if (nir_intrinsic_memory_scope(instr) != SCOPE_NONE) {
nir_variable_mode modes = nir_intrinsic_memory_modes(instr); nir_variable_mode modes = nir_intrinsic_memory_modes(instr);
if (modes & (nir_var_mem_global | nir_var_image)) if (modes & (nir_var_mem_global | nir_var_image))
agx_memory_barrier(b); agx_memory_barrier(b);
if (modes & nir_var_mem_shared)
needs_threadgroup_barrier = true;
if (modes & nir_var_image) { if (modes & nir_var_image) {
agx_image_barrier_1(b); agx_image_barrier_1(b);
agx_image_barrier_2(b); agx_image_barrier_2(b);
needs_image_barriers = true; 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); agx_threadgroup_barrier(b);
}
if (needs_image_barriers) { if (needs_image_barriers) {
agx_image_barrier_3(b); agx_image_barrier_3(b);