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

agx: Implement barriers

Browse Source 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/21062>
This commit is contained in:
Alyssa Rosenzweig
2022-12-03 14:34:44 -05:00
committed by Marge Bot
parent 251f6fb224
commit f857795e83
2 changed files with 16 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -776,6 +776,16 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
return agx_load_compute_dimension( return agx_load_compute_dimension(
b, dst, instr, AGX_SR_THREAD_POSITION_IN_THREADGROUP_X); b, dst, instr, AGX_SR_THREAD_POSITION_IN_THREADGROUP_X);
case nir_intrinsic_memory_barrier_buffer:
return agx_memory_barrier(b);
case nir_intrinsic_control_barrier:
return agx_threadgroup_barrier(b);
case nir_intrinsic_memory_barrier_shared:
/* Always seen with a control_barrier */
return NULL;
default: default:
fprintf(stderr, "Unhandled intrinsic %s\n", fprintf(stderr, "Unhandled intrinsic %s\n",
nir_intrinsic_infos[instr->intrinsic].name); nir_intrinsic_infos[instr->intrinsic].name);

6
src/asahi/compiler/agx_opcodes.py
View File

@@ -308,6 +308,12 @@ op("writeout", (0x48, 0xFF, 4, _), dests = 0, imms = [WRITEOUT], can_eliminate =
op("block_image_store", (0xB1, 0xFF, 10, _), dests = 0, srcs = 2, op("block_image_store", (0xB1, 0xFF, 10, _), dests = 0, srcs = 2,
imms = [FORMAT, DIM], can_eliminate = False) imms = [FORMAT, DIM], can_eliminate = False)
# Barriers
op("threadgroup_barrier", (0x0068, 0xFFFF, 2, _), dests = 0, srcs = 0,
can_eliminate = False)
op("memory_barrier", (0x96F5, 0xFFFF, 2, _), dests = 0, srcs = 0,
can_eliminate = False)
# Convenient aliases. # Convenient aliases.
op("mov", _, srcs = 1) op("mov", _, srcs = 1)
op("not", _, srcs = 1) op("not", _, srcs = 1)