agx: Plumb in coverage mask

This is internally used by the hardware when writing to the tilebuffer. We need
to use it externally to spill multisample render targets.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24258>

src/asahi/compiler/agx_compile.c

@@ -1027,6 +1027,9 @@ agx_emit_intrinsic(agx_builder *b, nir_intrinsic_instr *instr)
    case nir_intrinsic_load_sample_mask_in:
       return agx_get_sr_to(b, dst, AGX_SR_INPUT_SAMPLE_MASK);
 
+   case nir_intrinsic_load_sample_mask:
+      return agx_get_sr_to(b, dst, AGX_SR_COVERAGE_MASK);
+
    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?

src/asahi/compiler/agx_opcodes.py

@@ -128,6 +128,7 @@ SR = enum("sr", {
    53: 'subgroup_index_in_threadgroup',
    56: 'active_thread_index_in_quad',
    58: 'active_thread_index_in_subgroup',
+   60: 'coverage_mask',
    62: 'backfacing',
    63: 'is_active_thread',
    80: 'thread_position_in_grid.x',