Module: Mesa Branch: main Commit: e785ae61258439f930de7581d4aa686160a3daa2 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=e785ae61258439f930de7581d4aa686160a3daa2
Author: Alyssa Rosenzweig <[email protected]> Date: Thu Feb 9 22:58:59 2023 -0500 agx: Implement load_helper_invocation Passes dEQP-GLES31.functional.shaders.helper_invocation.* Signed-off-by: Alyssa Rosenzweig <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/21265> --- src/asahi/compiler/agx_compile.c | 8 ++++++++ src/asahi/compiler/agx_opcodes.py | 1 + 2 files changed, 9 insertions(+) diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c index 58448ddb917..9aa8132cee8 100644 --- a/src/asahi/compiler/agx_compile.c +++ b/src/asahi/compiler/agx_compile.c @@ -763,6 +763,14 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr) case nir_intrinsic_load_back_face_agx: return agx_get_sr_to(b, dst, AGX_SR_BACKFACING); + case nir_intrinsic_load_helper_invocation: + /* Compare special register to zero. We could lower this in NIR (letting + * us fold in an inot) but meh? + */ + return agx_icmpsel_to(b, dst, agx_get_sr(b, 32, AGX_SR_IS_ACTIVE_THREAD), + agx_zero(), agx_immediate(1), agx_zero(), + AGX_ICOND_UEQ); + case nir_intrinsic_load_vertex_id: return agx_mov_to(b, dst, agx_abs(agx_vertex_id(b))); diff --git a/src/asahi/compiler/agx_opcodes.py b/src/asahi/compiler/agx_opcodes.py index 1e138978020..28601af125b 100644 --- a/src/asahi/compiler/agx_opcodes.py +++ b/src/asahi/compiler/agx_opcodes.py @@ -138,6 +138,7 @@ SR = enum("sr", { 56: 'active_thread_index_in_quad', 58: 'active_thread_index_in_subgroup', 62: 'backfacing', + 63: 'is_active_thread', 80: 'thread_position_in_grid.x', 81: 'thread_position_in_grid.y', 82: 'thread_position_in_grid.z',
