VMAFx · lusoris · Sep 6, 2026 · Sep 4, 2026
@@ -23015,6 +23015,13 @@ of returning unsupported for every query.
   change is confined to the Metal `.mm`).
 
 
+- Corrected the `integer_motion_v2.metal` header comment to document the reflect-101
+  mirror fold implemented in #1223, closed out the deferred Metal mirror status
+  from ADR-1166 via ADR-1176, updated `test_metal_motion_v2_parity` to return exit code
+  77 (meson skip) on `-ENODEV` and log active device execution to stdout for CI observability,
+  and updated `docs/state.md`.
+
+
 **Metal dispatch table + ARC retain-balance fixes (PR #117 audit MT-1 + MT-2)**
 
 - MT-1 (`dispatch_strategy.c`): `g_metal_features[]` was missing

@@ -0,0 +1,5 @@
+- Corrected the `integer_motion_v2.metal` header comment to document the reflect-101
+  mirror fold implemented in #1223, closed out the deferred Metal mirror status
+  from ADR-1166 via ADR-1176, updated `test_metal_motion_v2_parity` to return exit code
+  77 (meson skip) on `-ENODEV` and log active device execution to stdout for CI observability,
+  and updated `docs/state.md`.
@@ -165,6 +165,7 @@ a real kernel lands; they are removed from `metal_sources` in
 
 ## Governing ADRs
 
+- [ADR-1176](../../../../docs/adr/1176-metal-motion-v2-mirror-closeout.md) — Metal motion_v2 mirror closeout and reflect-101 parity
 - [ADR-0490](../../../../docs/adr/0490-float-ms-ssim-metal-port.md) — T8-2b: float_ms_ssim_metal port
 - [ADR-0421](../../../../docs/adr/0421-metal-first-kernel-motion-v2.md) — T8-1c through T8-1k batch specification
 - [ADR-0420](../../../../docs/adr/0420-metal-backend-runtime-t8-1b.md) — runtime (T8-1b), prerequisite
@@ -182,3 +183,11 @@ a real kernel lands; they are removed from `metal_sources` in
   flush blend/clip/seed/average logic must be mirrored into all four GPU
   twins (cuda/sycl/hip/metal) in the same PR to keep the `places=4`
   `test_metal_motion_v2_parity` gate green.
+
+## mv2_mirror cross-twin invariant (ADR-1176)
+
+- **mv2_mirror is reflect-101, identical across backends**: `integer_motion_v2.metal::mv2_mirror`
+  uses iterated reflect-101 `idx = (idx < 0) ? -idx : 2 * (sup - 1) - idx`, bit-identical
+  to CPU `integer_motion_v2.c::mirror`, CUDA `motion_v2_score.cu::mv2_mirror`,
+  SYCL `integer_motion_v2_sycl.cpp::dev_mirror_mv2`, and HIP `motion_v2_score.hip::mv2_mirror`.
+  Do not revert to the single-bounce or `- 1` edge-replicating form.
@@ -17,10 +17,9 @@
  *                    + 32768) >> 16
  *    4. SAD: atomic-add |h[i,j]| into a single ulong accumulator.
  *
- *  Mirror padding: edge-replicating reflective mirror
- *  (`2 * size - idx - 1` for idx >= size), matching the CUDA twin.
- *  DIFFERS from motion_v1's `2 * size - idx - 2`; see CUDA file
- *  header for the bring-up note.
+ *  Mirror padding: reflect-101 (`2 * (sup - 1) - idx`, iterated),
+ *  identical to CPU integer_motion_v2.c::mirror, CUDA mv2_mirror,
+ *  SYCL dev_mirror_mv2 and HIP mv2_mirror; see mv2_mirror below.
  *
  *  Threadgroup layout: 16 × 16 threads per group, +2 pixel halo on
  *  each side → 20 × 20 shared tile. Inner pitch padded to 21 to
@@ -55,7 +54,7 @@ constant int MV2_FILTER[5] = {3571, 16004, 26386, 16004, 3571};
  *     Reflect-101 skips it: `2 * (sup - 1) - idx`. Every other backend
  *     already carries the corrected form -- CPU
  *     (integer_motion_v2.c::mirror), CUDA (fixed in PR #120 / T7-15),
- *     SYCL (integer_motion_sycl.cpp::dev_mirror_motion) and HIP
+ *     SYCL (integer_motion_v2_sycl.cpp::dev_mirror_mv2) and HIP
  *     (integer_motion_v2/motion_v2_score.hip) -- and the SYCL fix records
  *     the measured impact of the `- 1` form as a systematic ~2.6e-3 motion
  *     drift vs CPU on every frame after the first. Metal was the last

@@ -3202,7 +3202,12 @@ test_metal_motion_v2_parity = executable('test_metal_motion_v2_parity',
     dependencies: [pthread_dependency, math_lib],
     c_args: ['-DHAVE_METAL=1']
 )
-test('test_metal_motion_v2_parity', test_metal_motion_v2_parity, suite : ['fast', 'gpu'])
+test('test_metal_motion_v2_parity', test_metal_motion_v2_parity,
+    suite : ['fast', 'gpu'],
+    should_fail : false,
+    protocol : 'exitcode',
+    verbose : true,
+)
 
 # integer_ssim_metal — fixed-point SSIM (feature "ssim") on Metal; CPU vs.
 # Metal parity at places=4 (1e-4), skips cleanly when no Metal device.

@@ -40,8 +40,9 @@
  * also asserts motion3_v2 is finite on both paths.
  *
  * Skip behaviour: when `vmaf_metal_state_init` returns -ENODEV (Linux,
- * Windows, Intel Mac), the test emits "[skip: no Metal device]" and passes
- * cleanly. Mirrors the test_sycl_motion3_parity.c skip pattern.
+ * Windows, Intel Mac), the test sets `mu_skipped = 1`, emits
+ * "[skip: no Metal device]" to stderr, and exits 77 (meson skip).
+ * On Apple Silicon with a Metal device, it executes and exits 0 on pass.
  *
  * Cross-references:
  *   - core/test/test_cuda_motion_v2_parity.c (motion3_v2 parity, CUDA twin)
@@ -160,9 +161,13 @@ static char *run_metal_motion_v2(double scores_out[NUM_MOTION_V2_FEATURES])
     if (err != 0 || mstate == NULL) {
         /* No Apple-Family-7+ Metal device — skip cleanly. */
         (void)fprintf(stderr, "[skip: no Metal device] ");
+        mu_skipped = 1;
         return NULL;
     }
 
+    (void)fprintf(stdout, "[metal device active: motion_v2 parity run on device]\n");
+    (void)fflush(stdout);
+
     VmafConfiguration cfg = {.log_level = VMAF_LOG_LEVEL_NONE};
     VmafContext *vmaf = NULL;
     err = vmaf_init(&vmaf, cfg);

@@ -1,7 +1,8 @@
 <!-- markdownlint-disable MD013 MD060 -->
 # ADR-1166: Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites
 
-- **Status**: Accepted
+- **Status**: Accepted (Superseded-in-part 2026-09-04 by [ADR-1176](1176-metal-motion-v2-mirror-closeout.md) for Metal motion_v2 mirror deferral)
+- **Superseded-in-part (2026-09-04)**: Metal motion_v2 mirror deferral superseded by [ADR-1176](1176-metal-motion-v2-mirror-closeout.md) (landed in PR #1223, 71da046db).
 - **Date**: 2026-09-03
 - **Deciders**: Lusoris
 - **Tags**: process, upstream, bug, build, windows, api, docs

@@ -0,0 +1,91 @@
+<!-- markdownlint-disable MD013 MD060 -->
+# ADR-1176: Metal motion_v2 mirror closeout and reflect-101 parity contract
+
+- **Status**: Accepted
+- **Date**: 2026-09-04
+- **Deciders**: Lusoris
+- **Tags**: `metal`, `gpu`, `motion-v2`, `parity`, `boundary`, `closeout`, `fork-local`
+
+## Context
+
+[ADR-1166](1166-upstream-issue-harvest.md) recorded the upstream triage findings from
+Netflix/vmaf#1580 and noted that the Metal `integer_motion_v2` mirror kernel used
+`2 * sup - idx - 1` at the high boundary rather than reflect-101 (`2 * (sup - 1) - idx`).
+At triage time, ADR-1166 recorded this item as deferred under the premise that fixing
+it would move Metal scores and required dedicated Apple Silicon GPU verification and
+snapshot regeneration.
+
+PR #1223 (commit `71da046db`) subsequently implemented the reflect-101 fold in
+`core/src/feature/metal/integer_motion_v2.metal:74-80`:
+
+```metal
+inline int mv2_mirror(int idx, int sup)
+{
+    if (sup <= 1) return 0;
+    while (idx < 0 || idx >= sup)
+        idx = (idx < 0) ? -idx : 2 * (sup - 1) - idx;
+    return idx;
+}
+```
+
+This implementation iterated the bounce to handle out-of-bounds indices and corrected
+the boundary formula to `2 * (sup - 1) - idx`, matching:
+
+- CPU `core/src/feature/integer_motion_v2.c:157` (`mirror`)
+- CUDA `core/src/feature/cuda/integer_motion_v2/motion_v2_score.cu:51` (`mv2_mirror`)
+- SYCL `core/src/feature/sycl/integer_motion_v2_sycl.cpp:109` (`dev_mirror_mv2`)
+- HIP `core/src/feature/hip/integer_motion_v2/motion_v2_score.hip:67` (`mv2_mirror`, [ADR-1106](1106-hip-motion-v2-mirror-reflect101-correction.md))
+
+However, several closeout conditions remained unmet on `master`:
+
+1. The kernel header comment in `core/src/feature/metal/integer_motion_v2.metal:20-23`
+   contradicted the code by continuing to claim `2 * size - idx - 1` padding.
+2. `docs/state.md` and ADR-1166 still described the Metal fix as open/deferred.
+3. No Metal score snapshots exist in `testdata/` (`git ls-tree origin/master testdata | grep -i metal`
+   returns 0), so no snapshot regeneration (`/regen-snapshots`) was actually required.
+4. In `core/test/test_metal_motion_v2_parity.c`, non-Apple hosts returning `-ENODEV` from
+   `vmaf_metal_state_init` emitted `[skip: no Metal device]` to stderr but exited with 0.
+   Because Meson hides passing stderr in standard test runs, CI logs could not prove whether
+   the test executed on live hardware or skipped.
+
+## Decision
+
+1. **Close out the Metal motion_v2 mirror fix**: Formally record that the reflect-101
+   kernel fix landed in PR #1223 (`71da046db`). This ADR supersedes in part the deferral
+   recorded in ADR-1166 §Neutral / follow-ups.
+2. **Align documentation with implementation**: Correct the header comment in
+   `core/src/feature/metal/integer_motion_v2.metal` to document the iterated reflect-101
+   fold (`2 * (sup - 1) - idx`). Record the invariant in `core/src/feature/metal/AGENTS.md`.
+3. **No snapshot regeneration needed**: Confirm that no fork-added Metal reference
+   snapshots exist under `testdata/`. CPU golden assertions remain unaffected.
+4. **Make test skip observable**: In `core/test/test_metal_motion_v2_parity.c`, set
+   `mu_skipped = 1` on the `-ENODEV` branch so `test.c` exits 77 (Meson's standard skip code),
+   log explicit device confirmation to stdout when hardware is active, and configure the test in
+   `core/test/meson.build` with `should_fail : false`, `protocol : 'exitcode'`, and `verbose : true`.
+
+## Alternatives considered
+
+| Option | Pros | Cons | Why not chosen |
+|---|---|---|---|
+| Leave ADR-1166 deferral open | Zero documentation work | Leaves `state.md` and ADR-1166 out of sync with actual code in tree | Rejected: code was already merged in #1223; state tracking must reflect reality. |
+| Regenerate testdata snapshots | Follows standard GPU change playbook | No Metal snapshots exist in `testdata/` | Rejected: `testdata/` only contains CPU and select CUDA/SYCL snapshots. |
+| Keep exit 0 on `-ENODEV` skip | Matches legacy Metal parity tests | Indistinguishable from real device pass in CI logs | Rejected: exits 77 with verbose stdout gives clear observability of real hardware runs. |
+
+## Consequences
+
+- **Positive**: Code, comments, ADRs, and `docs/state.md` are unified; cross-backend parity
+  contract is documented; CI logs unambiguously distinguish device execution from skip.
+- **Negative**: None. Kernel code is unchanged from PR #1223; CPU golden data is untouched.
+- **Neutral / follow-ups**: Other `test_metal_*_parity.c` tests may adopt the observable
+  `mu_skipped = 1` exit 77 pattern in future maintenance sweeps.
+
+## References
+
+- PR #1223 (commit `71da046db`) — upstream issue harvest batch implementation
+- [ADR-1166](1166-upstream-issue-harvest.md) — upstream issue harvest triage (superseded in part)
+- [ADR-1106](1106-hip-motion-v2-mirror-reflect101-correction.md) — HIP motion_v2 mirror correction
+- [ADR-0421](0421-metal-first-kernel-motion-v2.md) — Metal motion_v2 initial kernel specification
+- [ADR-0214](0214-gpu-parity-ci-gate.md) — cross-backend parity gate (places=4)
+- `core/src/feature/metal/integer_motion_v2.metal`
+- `core/test/test_metal_motion_v2_parity.c`
+- Source: req
@@ -995,9 +995,10 @@ ADRs may exist there for local session continuity, but the tracked
 | [ADR-1154](1154-hip-backend-gaps.md) | AMD ROCm HIP backend gap closure: promotes 11 feature extractors to active GPU execution bringing active HIP extractors to 17/19 on AMD hardware, fixes pointer-to-pointer kernel parameter packaging in float_psnr_hip and float_moment_hip, fixes parameter ordering in float_moment_hip, fixes parameter types and partial buffer sizes to double in integer_ms_ssim_hip, fixes option dictionary serialization in integer_cambi_hip, fixes chroma plane copying in integer_psnr_hip, implements vmaf_hip_dispatch_supports with VMAF_HIP_DISPATCH env support, drains gpu_pending in libvmaf serial flush, prunes dead uncompiled files (adm_decouple.hip, integer_moment_hip.h/moment_score.hip), and formally defers integer_ssim_hip and integer_adm_hip. | Accepted | 2026-09-03 | hip, rocm, gpu, dispatch, parity, docs |
 | [ADR-1168](1168-default-model-single-source.md) | The default VMAF model is defined in exactly one place | Accepted | model, ci-gate, c-api, go, python, single-source |
 | [ADR-1169](1169-default-model-v1-0-16.md) | The fork's default VMAF model is `vmaf_v1.0.16_3d0h` | Accepted | model, default, v1.0.16, breaking-scores, upstream-divergence |
-| [ADR-1166](1166-upstream-issue-harvest.md) | Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites | Accepted | process, upstream, bug, build, windows, api, docs |
+| [ADR-1166](1166-upstream-issue-harvest.md) | Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites | Accepted (Superseded-in-part by [ADR-1176](1176-metal-motion-v2-mirror-closeout.md)) | process, upstream, bug, build, windows, api, docs |
 | [ADR-1172](1172-bound-lto-link-parallelism.md) | Bound per-link LTO parallelism to four partitions by default | Accepted | build, meson, developer-experience, lto |
 | [ADR-1171](1171-release-please-credential-gate-warning.md) | release-please credential gate warns on push, errors on dispatch | Accepted | ci, release, release-please |
 | [ADR-1179](1179-sycl-v1-model-crash-fix.md) | Fix Intel Arc SYCL Crashes and Default Model Resolution Divergence | Accepted | sycl, gpu, cambi, speed, model, default, arc, fp64, adr-0220 |
 | [ADR-1183](1183-model-options-gate-gpu-twin-selection.md) | Model options gate GPU twin selection | Accepted | core, feature, options, cuda, gpu-twins, dispatch |
+| [ADR-1176](1176-metal-motion-v2-mirror-closeout.md) | Metal motion_v2 mirror closeout and reflect-101 parity contract | Accepted | metal, gpu, motion-v2, parity, boundary, closeout, fork-local |
 | [ADR-1197](1197-gpu-threaded-flush-ownership.md) | The threaded flush leaves GPU extractors to their own backend flush | Proposed | cuda, sycl, threading, cli, testing |
@@ -1 +1 @@
-| [ADR-1166](1166-upstream-issue-harvest.md) | Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites | Accepted | process, upstream, bug, build, windows, api, docs |
+| [ADR-1166](1166-upstream-issue-harvest.md) | Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites | Accepted (Superseded-in-part by [ADR-1176](1176-metal-motion-v2-mirror-closeout.md)) | process, upstream, bug, build, windows, api, docs |
@@ -0,0 +1 @@
+| [ADR-1176](1176-metal-motion-v2-mirror-closeout.md) | Metal motion_v2 mirror closeout and reflect-101 parity contract | Accepted | metal, gpu, motion-v2, parity, boundary, closeout, fork-local |
@@ -909,4 +909,5 @@
 1171-release-please-credential-gate-warning
 1179-sycl-v1-model-crash-fix
 1183-model-options-gate-gpu-twin-selection
+1176-metal-motion-v2-mirror-closeout
 1197-gpu-threaded-flush-ownership
@@ -189,6 +189,9 @@ no rebase impact: fork-only test wiring in `core/test/meson.build` and documenta
 - `core/src/feature/sycl/speed_chroma_sycl.cpp`, `core/src/feature/sycl/speed_temporal_sycl.cpp`: Replaced `double` accumulators and workgroup local accessors with `float` to satisfy ADR-0220 on fp64-less Intel Arc devices.
 - `core/src/meson.build`: Passed `_x86_simd_strict_fp_extra` (`-fp-model=precise`) to `x86_avx2_static_lib` and `x86_avx512_static_lib` when compiling with `icx`.
 - `python/test/sycl_default_model_test.py`: Wholly fork-added regression test gating `--backend sycl` default model execution. No upstream rebase conflict.
+## fix/metal-motion-v2-mirror-closeout — Metal motion_v2 mirror closeout and test observability (2026-09-04)
+
+no rebase impact: fork-only Metal backend (`core/src/feature/metal/integer_motion_v2.metal`, `core/test/test_metal_motion_v2_parity.c`, `core/src/feature/metal/AGENTS.md`, ADR-1176). All touched files are fork-added surfaces with no upstream Netflix/vmaf counterpart.
 
 ## fix/vmaftune-state-bugs — libx264 two-pass CRF conflict fix (2026-09-03)
Original file line number	Diff line number	Diff line change
		@@ -1 +1 @@
		\| [ADR-1166](1166-upstream-issue-harvest.md) \| Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites \| Accepted \| process, upstream, bug, build, windows, api, docs \|
		\| [ADR-1166](1166-upstream-issue-harvest.md) \| Harvest stale upstream Netflix/vmaf reports, verify each against the fork, fix what still bites \| Accepted (Superseded-in-part by [ADR-1176](1176-metal-motion-v2-mirror-closeout.md)) \| process, upstream, bug, build, windows, api, docs \|
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1 @@
		\| [ADR-1176](1176-metal-motion-v2-mirror-closeout.md) \| Metal motion_v2 mirror closeout and reflect-101 parity contract \| Accepted \| metal, gpu, motion-v2, parity, boundary, closeout, fork-local \|