diff --git a/src/asahi/compiler/agx_compile.c b/src/asahi/compiler/agx_compile.c
index 7b9c900ac7d..3c74afaddcd 100644
--- a/src/asahi/compiler/agx_compile.c
+++ b/src/asahi/compiler/agx_compile.c
@@ -882,6 +882,9 @@ 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_sample_mask_in:
+      return agx_get_sr_to(b, dst, AGX_SR_INPUT_SAMPLE_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?
diff --git a/src/asahi/compiler/agx_opcodes.py b/src/asahi/compiler/agx_opcodes.py
index bbd264e14ee..22b9c34b73f 100644
--- a/src/asahi/compiler/agx_opcodes.py
+++ b/src/asahi/compiler/agx_opcodes.py
@@ -134,6 +134,7 @@ SR = enum("sr", {
    80: 'thread_position_in_grid.x',
    81: 'thread_position_in_grid.y',
    82: 'thread_position_in_grid.z',
+   124: 'input_sample_mask',
    144: 'opfifo_cmd',
    146: 'opfifo_data_l',
    147: 'opfifo_data_h',