Repository navigation
fix(hip): bring motion_hip and the HIP option twins onto the CPU's arithmetic - #1636
Merged
Merged
Conversation
lusoris
added a commit
that referenced
this pull request
Sep 30, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
September 30, 2026 13:49
fdb0d99 to
c00be3b
Compare
lusoris
added a commit
that referenced
this pull request
Sep 30, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
September 30, 2026 14:20
2bd6ca5 to
816233c
Compare
This was referenced Sep 30, 2026
lusoris
added a commit
that referenced
this pull request
Sep 30, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
September 30, 2026 23:53
816233c to
2b1325a
Compare
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 00:00
2b1325a to
8f2ed42
Compare
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 00:28
8f2ed42 to
ed98569
Compare
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 00:36
ed98569 to
9470d07
Compare
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 00:44
9470d07 to
9bf25d0
Compare
8 tasks done
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 06:48
9bf25d0 to
3139555
Compare
…ithmetic motion_hip blurred each frame and differenced the blurred frames, while the CPU motion differences the frames first and rounds after each filter pass, so motion2 / motion3 were 1.26e-5 off on the Netflix pair on a gfx1036 (T-HIP-MOTION-BLUR-THEN-DIFF-2026-09-29). motion_hip now runs the diff-first kernel motion_v2_hip already used, through one shared launcher (integer_motion_sad_hip.c); motion_score.hip is gone. Both motion twins copy the reference luma into pinned memory and upload it without a host wait, so collect() is the only wait of a frame, and motion_hip's debug motion score carries motion_fps_weight / motion_max_val like the CPU's (ADR-1377). The motion tile loads and the integer ADM scale-0 vertical DWT now clamp their reflected rows into the plane. A host replay of every launched thread shows the bare reflection escapes only below 9 rows, so the clamp is the identity for every frame the twins accept. vif_hip needs 16x16 frames like vif_sycl: model dispatch sends smaller frames to the CPU vif and a direct request fails at init (ADR-1381). psnr_hip, integer_ssim_hip, float_ssim_hip and float_motion_hip take the CPU option tables: the PSNR options through psnr_score.h with apsnr_* from a new flush, enable_db / clip_db on both SSIM twins, enable_lcs as a device kernel variant, motion_max_val on float_motion_hip, and identical SSIM windows scoring exactly 1 (ADR-1382). The HIP twins share vmaf_hip_rc_to_errno(), which common.h declared but nothing defined. Built for gfx90a, gfx1030, gfx1036 and gfx1100 in vmaf-dev-mcp; every touched kernel produces all four code objects, the fast suite passes apart from one container-only timeout in untouched code (see the PR), and the HIP device tests skip without an AMD GPU. Not yet run on AMD hardware: the docs/state.md rows carry the verify-and-time commands for ryzen-4090-arc.
Two reviews of #1636 found these; each fix has a device-free check that fails on the previous code. - float_motion_score.hip reflected tile indices once, so padding threads read before the input plane on 3x3 to 9x9 and 17x17 frames. The loads now use the clamped index that motion_v2 uses (hip_tile_index.h), which is the identity for every consumed sample. - float_ssim_hip forced an identical window to exactly 1. The CPU reports 1 - 2^-24 (72.247 dB) on some identical frames because its fp32 luminance denominator rounds. Pass 2 now forms each pixel's term as ssim_accumulate_default_scalar() does (l * c * s in double from the CPU-typed factors) and the host rounds the frame mean to fp32 like iqa_ssim(). integer_ssim_hip keeps its exact-1 rule, which matches the CPU from 3x3 up; 1x1 and 2x2 are recorded as a low RC3 row. - motion_v2_hip stored the unweighted, unclipped SAD and emitted nothing for a one-frame run. It now stores MIN(score * mfw, mmxv) like the CPU and folds motion2_v2 / motion3_v2 from that. - psnr_hip is now TEMPORAL like the CPU psnr, so --subsample keeps every frame in apsnr_*. - motion_hip now defaults debug to false and emits VMAF_integer_feature_motion_sad_score, matching the CPU motion. - The motion twins bound the staged upload by the allocated size and no longer rewrite the geometry in submit(). An enqueue error after a copy started now drains the stream before returning. vif_hip now returns -ENOSYS first in scaffold builds (ADR-1264). - The parity gate takes the hip backend and a float_ssim_lcs cell, and test_hip_twin_option_parity uses the gate's 5e-5 for float_ssim. - .standards-baseline.json is re-recorded with the pinned engine (the deleted motion_score.hip is gone, the float_motion kernels moved). - The duplicated RC3 disposition row in docs/state.md is merged into one. The float_motion motion3 and option gap is opened as an RC3 row (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30). The motion3 moving average at two frames was reported but refuted: (x + x) / 2 == x exactly.
…ontract current submit() no longer reads the picture geometry, so the scaffold branch of motion_hip and motion_v2_hip left ref_pic unused (-Wunused-parameter in an enable_hipcc=false build). float_ssim_hip now emits the fp32-rounded CPU mean through vmaf_ssim_emit_score_named() after validating the ratio, so test_nonfinite_collector_wiring names that emitter for it.
…the README block The review fix pass re-recorded .standards-baseline.json with the pinned engine (185 -> 182: the deleted motion_score.hip and the refactored motion kernels), but the managed README governance block still said 185, so standardsctl audit failed the README governance check. The pinned adopt cannot render the block on this tree (it aborts on the editor scan bound), so the count is set to the value the renderer derives from the baseline; standardsctl audit now reports the block verified.
libvmaf chooses submit()/collect() from an extractor's callbacks before init() runs (read_pictures_dispatch_one()), and vmaf_feature_extractor_context_submit() checks fex->submit before it calls init(). Under motion_force_zero, motion_hip and float_motion_hip switched to extract() inside init() and set submit and collect to NULL, so frame 0 called a NULL submit(): on master 10f27ef both --feature motion_hip=motion_force_zero=true and --feature float_motion_hip=motion_force_zero=true exit 139 on a gfx1036. The twins now keep a no-op submit() and a collect() that writes the zeros extract() writes; motion_hip frees its device objects at the switch and keeps close() for the name dictionary. Both runs exit 0 on the gfx1036 and emit the CPU's force-zero features, all 0. test_integer_motion_force_zero drives motion_hip through vmaf_read_pictures() and passes on the device. State row T-HIP-MOTION-FORCE-ZERO-NULL-SUBMIT-2026-09-30 (closed). The CUDA motion twins crash the same way; that fix is separate.
Both cases had never run on an AMD device and failed on the gfx1036;
the twins were right, the expectations were stale.
- test_hip_float_motion_parity expected the flushed tail motion2 to be
1.5 times the last debug motion score. Since ADR-1382 the debug score
goes through motion_clip() like the CPU's, so both carry
motion_fps_weight once and the tail equals the last debug score.
- test_hip_twin_option_parity looked up the one-frame motion2 / motion3
scores under their JSON names (integer_motion2); the feature collector
holds them as VMAF_integer_feature_motion{2,3}_score.
The HIP twins of this branch ran on a gfx1036 (Ryzen 9950X3D iGPU, ROCm 7.2.4) for the first time: - every HIP device test passes (row tests and the whole gpu suite); - motion_hip equals the CPU motion on every frame (1.26e-5 apart on master), and on all 9600 frames of 20 looped runs; - the parity gate passes every HIP cell (float_ssim and float_ssim_lcs 1.0e-5, psnr 0, motion_v2 0, vif 1.0e-6), psnr_hip with all four options matches the CPU, apsnr_* included; - float_motion_hip exits 0 on 3x3 and 17x17 frames. T-HIP-MOTION-BLUR-THEN-DIFF-2026-09-29 and T-HIP-FLOAT-MOTION-TILE-OOB-2026-09-30 close; the HIP halves of the ADM-DWT, VIF-min-dim and BUG048 rows are recorded as verified, those rows stay open for CUDA and Metal. Two findings go on record. The staged upload of the motion twins costs throughput on this iGPU: at 4K a single motion_v2_hip goes from 10.17 to 13.24 ms/frame, while the waiting upload with the same kernel gives 10.70 (T-HIP-UPLOAD-WAIT-THROUGHPUT-2026-09-19 has the numbers, runs with several twins included). And the gfx1036 loses a run of a stream's commands about once per 10^4 frames, master included, which is what made vif_hip report the previous frame's sums now and then. The accumulator memset was never the cause, so the device-test commit no longer replaces it. scripts/dev/hip_dispatch_drop_probe.hip reproduces the loss without vmafx (T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01, explicitly deferred).
The branch now sits on #1637, which fixed the CUDA halves of the shared rows and changed the engine's motion_force_zero dispatch. - The metric pages said the HIP twins were not yet measured on an AMD device; they state the gfx1036 results the state rows already record (motion2/motion3 and the PSNR option outputs equal the CPU's). - motion_force_zero: the engine initialises an extractor before it picks the dispatch path since #1637, so a cleared submit()/collect() pair no longer crashes. The HIP invariant note, the rebase note and the CUDA changelog entry say that the HIP twins keep their asynchronous pair.
lusoris
force-pushed
the
fix/hip-rc3-parity
branch
from
October 1, 2026 06:56
3139555 to
b6d0877
Compare
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
cambi_hip, speed_chroma_hip and speed_temporal_hip no longer run any stage of the CPU extractors on the host. Each frame is one staged upload (vmaf_hip_picture_upload_staged()), the whole pipeline on the extractor's stream and one small readback; collect() is the only host wait. This ports the SYCL designs of ADR-1357 (CAMBI) and ADR-1358 (SpEED) to HIP as ADR-1378 and ADR-1384. CAMBI keeps cambi.c's sliding column histograms and sums the top-K c-values exactly, so it is bit-identical to the CPU wherever the CPU's own double sum is exact, and it now applies cambi.c's reciprocal-table window guard. The window, mask index, resize tables, contrast weights, top-K mean and the guard are new shared helpers in cambi.c, which the CPU init and the SYCL twin call too. SpEED reproduces speed.c in fp32 operation for operation. On HIP that is a per-TU build flag (-ffp-contract=off, correctly rounded divide and sqrt), not per-operation intrinsics: HIP's __fmul_rn/__fadd_rn/__fdiv_rn are the plain contracting operators and __fsqrt_rn the native approximation. The init-time configure is one shared routine, speed_internal_gpu_configure(), used by the SYCL twins as well. Both kernels keep their per-work-item arithmetic in one header that a host test replays against the CPU extractors in the fast suite (test_hip_cambi_device_math, test_hip_speed_device_math), with planted regressions; test_hip_device_resident_contract.py pins one upload, one readback and one wait per frame. Built for gfx90a, gfx1030, gfx1036 and gfx1100; not yet run on an AMD device (verify commands in docs/state.md). Depends on #1636 (vmaf_hip_picture_upload_staged(), vmaf_hip_rc_to_errno()).
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
cambi_hip, speed_chroma_hip and speed_temporal_hip no longer run any stage of the CPU extractors on the host. Each frame is one staged upload (vmaf_hip_picture_upload_staged()), the whole pipeline on the extractor's stream and one small readback; collect() is the only host wait. This ports the SYCL designs of ADR-1357 (CAMBI) and ADR-1358 (SpEED) to HIP as ADR-1378 and ADR-1384. CAMBI keeps cambi.c's sliding column histograms and sums the top-K c-values exactly, so it is bit-identical to the CPU wherever the CPU's own double sum is exact, and it now applies cambi.c's reciprocal-table window guard. The window, mask index, resize tables, contrast weights, top-K mean and the guard are new shared helpers in cambi.c, which the CPU init and the SYCL twin call too. SpEED reproduces speed.c in fp32 operation for operation. On HIP that is a per-TU build flag (-ffp-contract=off, correctly rounded divide and sqrt), not per-operation intrinsics: HIP's __fmul_rn/__fadd_rn/__fdiv_rn are the plain contracting operators and __fsqrt_rn the native approximation. The init-time configure is one shared routine, speed_internal_gpu_configure(), used by the SYCL twins as well. Both kernels keep their per-work-item arithmetic in one header that a host test replays against the CPU extractors in the fast suite (test_hip_cambi_device_math, test_hip_speed_device_math), with planted regressions; test_hip_device_resident_contract.py pins one upload, one readback and one wait per frame. Built for gfx90a, gfx1030, gfx1036 and gfx1100; not yet run on an AMD device (verify commands in docs/state.md). Depends on #1636 (vmaf_hip_picture_upload_staged(), vmaf_hip_rc_to_errno()).
lusoris
added a commit
that referenced
this pull request
Oct 1, 2026
* perf(hip): run cambi and SpEED entirely on the device, matching the CPU cambi_hip, speed_chroma_hip and speed_temporal_hip no longer run any stage of the CPU extractors on the host. Each frame is one staged upload (vmaf_hip_picture_upload_staged()), the whole pipeline on the extractor's stream and one small readback; collect() is the only host wait. This ports the SYCL designs of ADR-1357 (CAMBI) and ADR-1358 (SpEED) to HIP as ADR-1378 and ADR-1384. CAMBI keeps cambi.c's sliding column histograms and sums the top-K c-values exactly, so it is bit-identical to the CPU wherever the CPU's own double sum is exact, and it now applies cambi.c's reciprocal-table window guard. The window, mask index, resize tables, contrast weights, top-K mean and the guard are new shared helpers in cambi.c, which the CPU init and the SYCL twin call too. SpEED reproduces speed.c in fp32 operation for operation. On HIP that is a per-TU build flag (-ffp-contract=off, correctly rounded divide and sqrt), not per-operation intrinsics: HIP's __fmul_rn/__fadd_rn/__fdiv_rn are the plain contracting operators and __fsqrt_rn the native approximation. The init-time configure is one shared routine, speed_internal_gpu_configure(), used by the SYCL twins as well. Both kernels keep their per-work-item arithmetic in one header that a host test replays against the CPU extractors in the fast suite (test_hip_cambi_device_math, test_hip_speed_device_math), with planted regressions; test_hip_device_resident_contract.py pins one upload, one readback and one wait per frame. Built for gfx90a, gfx1030, gfx1036 and gfx1100; not yet run on an AMD device (verify commands in docs/state.md). Depends on #1636 (vmaf_hip_picture_upload_staged(), vmaf_hip_rc_to_errno()). * test(hip): split the device-resident contract checks into shared helpers The CAMBI and SpEED checks repeated the same host-stage, frame-path and readback logic; one helper each removes the repeats and keeps ruff's branch limit. * fix(hip): report -ENOSYS first from the scaffold cambi and SpEED twins ADR-1264 has a build without hipcc report -ENOSYS and nothing else; the speed_temporal_hip ran their host configure first, so a bad window or format returned -EINVAL there. They now return -ENOSYS before any check, and the window-guard test checks only cambi.c's half on a scaffold build. The contract test pins the order, with a planted regression per family. * docs(hip): record measured gfx1036 results for device-resident cambi and SpEED * fix(hip): add assertions to cambi_hip_plan to satisfy assertion density * docs(hip): align the cambi notes with the short-frame fix on master described the CUDA, HIP and Metal twins as running that walk on the host. With cambi_hip device-resident that no longer holds for HIP (nor for CUDA since ADR-1379): only the Metal twin calls vmaf_cambi_calculate_c_values(). The feature notes, the metric pages and the changelog entry say so, and the state row records that cambi_hip equals the CPU on the gfx1036 on narrow frames as well. * chore(ci): regenerate the source ADR citation map after rebasing The rebase onto master took master's map for the conflicted file, which still listed the deleted core/src/feature/hip/speed/speed_score.hip under ADR-0567 and lacked this branch's new sites. * docs: regenerate the indexes and the citation map after rebasing
This was referenced Oct 1, 2026
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity.
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…HIP and SYCL twins (ADR-1498) (#1921) * build(metal): compile every Metal kernel without fast math or contraction (ADR-1498) The Metal compiler's default is fast math: it may divide through a reciprocal, reassociate and contract across statements, and its safe mode still contracts within a statement. Every .metal now takes one list, -fno-fast-math -ffp-contract=off (metal_shader_strict_fp_args), under which fp32 + - * /, sqrt and fma are correctly rounded (Metal Shading Language Specification 4.1, sections 1.6.3 and 8.4; Research-1498): the policy of the CUDA (ADR-1403), HIP and SYCL (ADR-1367) kernels. Kernels may include feature/ and feature/metal/. core/src/feature/metal/metal_portable.h is the prelude of the ports' arithmetic headers: one spelling of the types, fma, bit casts and clz that compiles as Metal Shading Language and as host C or C++, so a host test can hold a port's arithmetic against the CPU (test_metal_portable checks the host primitives). The SYCL twins' fp64-free forms are followed in Metal headers until RC5 merges the two (T-METAL-SYCL-FP64-FREE-ARITHMETIC-COPIES-2026-10-03). test_metal_shader_build_contract also holds the strict list: defined once between markers, taken by every kernel, no flag that turns fast math or contraction back on. * fix(metal): add the fp64 operations in 64-bit integers that exact Metal twins need Metal has no fp64 type, so a Metal twin whose CPU reference computes in double cannot return the CPU's bits without the fp64 operations in 64-bit integers that the SYCL twins use (sycl_soft_double.h, sycl_soft_signed.h). Those headers are written against sycl:: and C++20 and cannot be included from Metal Shading Language, so the Metal ports of integer_ssim, float_vif, float_ssim / float_ms_ssim and float_adm had nothing to follow their SYCL twins with. This is infrastructure of the rows T-GPU-SSIM-FRAME-SUM-ORDER-2026-10-01, T-GPU-FLOAT-VIF-CPU-ARITHMETIC-2026-10-01, T-GPU-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01, T-GPU-FLOAT-ADM-CPU-ARITHMETIC-2026-10-01 and the float_ssim part of T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30; it closes none of them. core/src/feature/metal/metal_soft_double.h and metal_soft_signed.h are the two SYCL headers statement for statement (ADR-1498): the same algorithms, each operation rounding to nearest even as the fp64 operation it stands for, written on metal_portable.h in the subset Metal Shading Language, C and C++ share (values in and out, typedef'd structs built by _make() functions instead of designated initializers, constants as macros, the 64x64 product in 32-bit limbs, no double, long long or 128-bit type). Every SYCL function, struct and constant has a counterpart under the mapping documented at the top of each header: vmaf_sycl_soft::f -> vmaf_mtl_f, struct S -> VmafMtlS, kName -> VMAF_MTL_SOFT_NAME. What proves it here: test_metal_soft_double (every host, fast suite) compiles both headers on the host and compares every operation bit for bit with the host's binary64 arithmetic, on at least 10^6 operands per operation from a fixed SplitMix64 sequence (ties and near-ties, sums and products that carry to a power of two, cancellations, far-apart exponents, zeros, integers around 2^53 and 2^64) plus pairwise edge tables, on every fp32 subnormal, and the 128-bit helpers against bit-by-bit references. Its 21 cases pass in 1.1 s; 30 arithmetic bugs planted in the headers one at a time were each caught (two further mutants return the same results by construction). test_metal_soft_double_contract.py pins the Metal subset and the mapping (signatures, struct fields, constant values, and per function the SYCL function's numeric literals and branch, loop, return, select and shift counts) and holds 21 planted-regression cases. The headers also type-check as C11, C17, C23 and C++14/17/20, and as C++ through the Metal branch of metal_portable.h against a stand-in for <metal_stdlib>. What only the macOS job and an Apple device prove: no kernel includes the headers yet, so the hosted macOS job compiles them first with the port that does, and the tester's device run of those twins shows the GPU returns the same bits, including vmaf_mtl_div_step() under the device's fp32 division (its corrections cover one digit either way). * fix(metal): add float_psnr's squared differences as exact integers float_psnr_metal added each 16x16 threadgroup's squared differences in fp32 with simd_sum() and the host added the float partials in double. float_psnr.c adds its float squares in double, which is exact; an fp32 group sum is exact at 8 bits only and rounds at 10, 12 and 16 bits once a group's differences are large (T-METAL-FLOAT-PSNR-FP32-BLOCK-SUMS-2026-10-02). The kernel now forms the CPU's term as an integer in units of 1 / scaler^2: one fp32 product of the raw sample difference with itself, converted to a 32-bit integer (vmaf_mtl_fpsnr_term() in the new metal_float_psnr_math.h, valid as MSL and as host C). Each threadgroup adds its 256 terms as ulong in threadgroup memory (MSL has no 64-bit simd_sum or usable 64-bit atomic) and stores one ulong; the host adds the group sums in uint64 and divides the exact total by scaler^2 and the pixel count, as float_psnr_cuda.c::float_psnr_noise() does. Design: ADR-1455 (CUDA), ADR-1450 (SYCL), ADR-1440 (HIP), under ADR-1498. Proven here: test_metal_float_psnr_math compiles the header on the host and holds it against float_psnr.c's float square through the CPU's picture_copy() for every sample difference at 8, 10, 12 and 16 bits (at the bottom and the top of the range; at 16 bits it returns the rounded squares, not the integer squares), and the integer group sums against the CPU's noise on 576x324 full-range noise frames at every depth (equal; the former fp32 group sum differs at 16 bits). test_metal_float_psnr_exact_contract fails on a planted return of the fp32 simd_sum, a float group total, a float partials buffer or readback, a host double sum, a scaled or integer square, a kernel-side float conversion and another host division order. Only the device run proves the kernel's results: test_metal_float_psnr_parity (the cases of float_psnr_twin_parity.h at ==) on an Apple device, and that the hosted macOS job compiles the kernel. * fix(metal): add float_moment's 16-bit second moments as the CPU's float squares float_moment_metal's 10/12/16-bit kernel added exact integer squares (rv * rv) where moment.c::compute_2nd_moment() forms each square in float, which at 16 bits is the integer square rounded to 24 bits; the second moments of full-range 16-bit content were therefore not the CPU's (T-GPU-FLOAT-MOMENT-16BIT-SQUARES-2026-10-02). The host also added the exact workgroup sums in double, which rounds once a sum passes 2^53. The kernel now adds vmaf_mtl_moment_float_square() (the new metal_float_moment_math.h, valid as MSL and as host C): one fp32 product of the sample with itself, converted to an integer below 2^32, as moment_float_square() of the CUDA, SYCL and HIP twins does (ADR-1453, ADR-1449, ADR-1447, under ADR-1498). The 8-bit kernel keeps the integer square, which is the CPU's float there. The host adds the reconstructed workgroup sums in uint64 and converts each total once; its divisions by n_pix * scaler (exact products) are the CPU's division of the same exact value, so test_metal_float_moment_contract's pins stay as they are. Proven here: test_metal_float_moment_math compiles the header on the host and holds it against picture_copy() and compute_2nd_moment() on a one-pixel plane for every sample value at 8, 10, 12 and 16 bits (53248 16-bit squares are rounded, and an integer square fails exactly those), and the twin's moments against compute_1st_moment() / compute_2nd_moment() on full-range noise at every depth and a bright 16-bit 1920x1080 frame (equal; exact integer squares give another second moment at 16 bits). test_metal_float_moment_exact_contract fails on a planted integer square in the kernel or the header, a scaled sample, a double host sum or double totals, and a changed reference square. Only the device run proves the kernel's results: test_metal_float_moment_parity (full-range noise at 8 to 16 bits and the bright 16-bit pair at ==) on an Apple device, and that the hosted macOS job compiles the kernel. * fix(metal): form integer ADM's gain-limited and restored samples in integer arithmetic integer_adm_metal multiplied the restored sample by adm_enhn_gain_limit narrowed to binary32 ((int)((float)rst * egl)), where the CPU's decouple stores the double product truncated toward zero; the binary32 product is another number for a non-integer limit such as 1.2 and 1.5, and at scales 1-3 the restored sample itself was converted to float, which rounds above 2^24 (T-METAL-ADM-GAIN-LIMIT-FLOAT32-2026-10-01). Read from source in the same functions: the scale-0 Q15 ratio took the reciprocal 2^30 / o as the fp32 quotient truncated, where div_lookup holds the integer division; the two differ for 686 of the 65535 operands (|o| = 3 is the first) and move k for 741176 (o, t) pairs. The CUDA and HIP decouple_r_s0() use the same fp32 reciprocal. The decouple now lives in the new metal_integer_adm_math.h (valid as MSL and as host C): the integer reciprocal, the CPU's Q15 ratio at both scales, the restored sample as the int64 expression narrowed to int32, and the gain limit as adm_gain_limit_product() of the shared adm_gain_limit.h, the integer form the SYCL twin uses (ADR-1413, under ADR-1498). The host splits the limit with adm_gain_limit_split() and passes its significand halves and binary point in three fields of the CSF uniform, in IadmCsf and IadmCsfHost alike. adm_gain_limit.h keeps its includes and the double-only split out of MSL behind __METAL_VERSION__; the guard replaces blank and comment lines, so the preprocessed output of every other includer is byte for byte the same (gcc -E and clang -E of test_adm_gain_limit.c, icpx -fsycl -E and icpx -fsycl -fsycl-device-only -E of integer_adm_sycl.cpp: cmp equal before and after). The option table already equals the CPU adm's. Proven here: test_metal_integer_adm_math compiles the header on the host and holds it against adm_decouple_band() and adm_decouple_band_s123() with the CPU's div_lookup: every scale-0 operand against 27 distorted values at limits 1, 1.2, 1.5 and 100 with the angle flag set and clear, every distorted value at the 686 operands with a wrong fp32 reciprocal (the former kernel is off on 318294 of them), and two million scales-1-3 pairs at six limits (the former binary32 product is off on 28149, 17382 and 23569 at 1.2, 1.5 and 100); a planted fp32 reciprocal or binary32 product fails it. The header's Metal branch compiles with -nostdinc against the MSL types only, without a double. test_metal_integer_adm_exact_contract fails on a planted fp32 reciprocal, binary32 gain product, float restored sample, kernel-side decouple, float limit on the host, IadmCsf / IadmCsfHost layout drift, an unguarded split in the shared header and a changed CPU reference. Only the device run proves the kernel's results: test_metal_integer_adm_parity (test_adm_gain_limit_1_2_exact, test_adm_gain_limit_1_5_exact and the other cases at ==, which also need the rest of the Metal ADM pipeline to be the CPU's) on an Apple device, and that the hosted macOS job compiles the kernel and its include of adm_gain_limit.h. * fix(metal): difference motion frames before the blur and emit the CPU motion's SAD score and table integer_motion_metal blurred each frame into a uint16 ping-pong and summed the absolute differences of the blurred frames, where integer_motion.c (since Netflix a4a1492d) blurs prev - cur, rounding after the vertical and after the horizontal pass; the two orders round differently (T-METAL-MOTION-BLUR-THEN-DIFF-2026-09-29). It also declared only motion_add_uv, an option the CPU motion does not have, emitted VMAF_integer_feature_motion_y_score and a motion2 of its own, and never VMAF_integer_feature_motion_sad_score or motion3 (T-GPU-MOTION-SAD-SCORE-NOT-EMITTED-2026-10-02). The kernel now stages prev - cur of a 16x16 threadgroup and its halo, filters it vertically and horizontally with the CPU's rounding and adds |h| as integers (the new metal_integer_motion_math.h, valid as MSL and as host C; design: core/src/feature/sycl/integer_motion_pipeline_sycl.cpp, ADR-1371, and the CUDA motion_v2_score.cu SAD kernel, ADR-1372, under ADR-1498). The mirror is the CPU's reflect-101 within two samples of the plane and clamps the tile positions that feed no pixel. The host keeps a ring of raw luma planes, three with motion_five_frame_window (ADR-1491): frame n writes slot n % ring and the SAD is taken against slot (n + 1) % ring. collect() appends the CPU's SAD score on every frame (0 before the first SAD and under motion_force_zero, otherwise MIN(sad / 256 / (w * h) * motion_fps_weight, motion_max_val), and the same value as VMAF_integer_feature_motion_score with debug); flush() derives motion2 and motion3 with the CPU's vmaf_motion_window_flush() (ADR-1478) for both windows. The option table and provided_features are integer_motion.c's (the SAD score first); motion_add_uv, accepted only as false before, is gone. Proven here: test_metal_integer_motion_math runs the header in the kernel's threadgroup layout and holds the SAD score against the CPU motion extractor's VMAF_integer_feature_motion_sad_score (public API) at 3x3, 4x5, 17x17, 19x23, 33x33, 31x9, 64x64 and 257x145, at 8, 10, 12 and 16 bits, on noise, a full-scale step and full-scale stripes: 96 of 96 equal, while blurring first gives another SAD on 53; the mirror stays in the plane for every tile index at sizes 3 to 80. A planted unrounded vertical pass or a wrong mirror fails it. test_metal_integer_motion_exact_contract fails on a planted blur-then-diff, a float reduction, an unrounded pass, a changed tap, a lost ring, the twin's own motion2, a missing SAD score, an option outside the CPU table, a changed default, a missing provided feature and a changed CPU score. The Metal selftests stay green. Only the device run proves the kernel's results: test_metal_integer_motion_parity (every option case, tiny and large frames, the five-frame window, the three SAD-score sites, at ==) and the motion cases of test_metal_twin_option_parity on an Apple device, and that the hosted macOS job compiles the kernel and the host. test_motion_five_frame_window still lists integer_motion_metal among the twins without the window and will fail on a Metal build until it moves to the twins that keep the option. * fix(metal): psnr twin takes the CPU options and sums an exact integer SSE Rows: T-BUG048-GPU-OPTION-PARITY-REMAINDER-2026-09-26 (psnr part) and item (2) of T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30. What was wrong: integer_psnr_metal declared two of the six options of the CPU `psnr` (enable_chroma, uncapped), so a request or model with min_sse, enable_mse, reduced_hbd_peak or enable_apsnr kept the CPU extractor. It had no flush (no apsnr_*) and no VMAF_FEATURE_EXTRACTOR_TEMPORAL, so --subsample would have summed apsnr over a subset of the frames. The host added the block sums in double and halved every chroma plane as 4:2:0, so a 4:2:2 or 4:4:4 chroma SSE read the wrong number of blocks. The kernel reduced each block with two uint32 simd_sum() calls on the halves of the squares; at 14 and 16 bits a 32-lane sum of squares passes 2^32 and the carry is lost. What the port does (the design of psnr_sycl, psnr_cuda and psnr_hip, ADR-1365, ADR-1373, ADR-1382; ADR-1498): each kernel thread forms its squared difference in 64 bits, thread 0 adds the threadgroup's 256 values in uint64 from threadgroup memory and stores one exact SSE per group (no simd_sum, no atomics, no float). The host adds the group sums in uint64, forms the MSE with the CPU's expression and scores through core/src/feature/psnr_score.h: vmaf_psnr_peak(), vmaf_psnr_max() per plane, vmaf_psnr_from_mse() and, from a new flush(), vmaf_psnr_aggregate() for apsnr_*. enable_mse emits mse_* after each psnr_*. The option table is the CPU's (names, types, defaults, ranges, flags); the twin is TEMPORAL; chroma planes take the CPU's ceiling subsampling per pixel format. init/close share one teardown helper. What proves it here: test_metal_integer_psnr_exact_contract.py (device-free) compares the twin's option table with integer_psnr.c's, read from both sources by the new core/test/metal_option_tables.py, and asserts the uint64 group reduction, the psnr_score.h calls, flush, TEMPORAL and the pixel-format geometry; ten planted regressions (a missing option, a range change, no TEMPORAL, no flush, a double block sum, an own log10 score, 4:2:0 chroma for every format, a simd_sum, a float term, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_integer_psnr passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. test_metal_kernel_registration.c asserted that the twin rejects min_sse; it now asserts the CPU options are taken and an unknown key is still refused. What only the device run proves: the kernel compiles on the hosted macOS runner, and test_metal_integer_psnr_parity (every case at ==) plus the psnr cases of test_metal_twin_option_parity pass on an Apple GPU. * fix(metal): motion_v2 twin stores the CPU's capped score and derives motion2/3 with its window Row: T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30, item (3) (motion_v2). What was wrong: motion_v2_metal stored the raw SAD score where integer_motion_v2.c stores MIN(SAD / 256 / (w * h) * motion_fps_weight, motion_max_val). Its own flush weighted the stored value again, left the motion_max_val cap off motion2_v2, blended the unweighted SAD into the motion3_v2 stamp, and returned without output for a one-frame input (the CPU emits motion2_v2 = motion3_v2 = 0). It added the block SADs in double and lacked two CPU options, motion_force_zero and motion_five_frame_window, so a request or model setting either kept the CPU extractor. What the port does (the design of motion_v2_hip / motion_v2_sycl / motion_v2_cuda, #1636, #1645, ADR-1373, ADR-1382, ADR-1491; ADR-1498): collect() adds the threadgroup SADs in uint64 and stores the CPU's value with the CPU's expression; flush() is the CPU extractor's own vmaf_motion_window_flush() (motion_window.h, ADR-1478) for the three- and the five-frame window, so a one-frame run gets 0 / 0. motion_five_frame_window takes the SAD against frame n - 2: the twin keeps the luma of the last `depth` frames (1 or 2) in Shared buffers, frame n reading and then replacing slot n % depth, and stores 0 below frame `depth` (the CPU's min_idx). motion_force_zero stores 0 for every frame and runs no kernel, as the CPU's extract(). The option table is the CPU's (order, aliases, ranges, FEATURE_PARAM flags), so feature names carry the same suffixes. The 3x3 floor stays the first check of init(). init/close share one teardown helper. The kernel is unchanged (its group SAD is an exact uint32, at most 256 x 65536); only its stale comments (atomic accumulator, places=4) are corrected. What proves it here: test_metal_motion_v2_exact_contract.py (device-free) compares the option table with integer_motion_v2.c's, asserts the stored score expression, the uint64 SAD, the window call, the n - 2 read, the force-zero store and the integer kernel reduction in both kernels; nine planted regressions (a missing option, an alias change, the raw SAD store, an own flush with an n_frames < 2 return, a double block sum, a three-frame-only depth, force zero ignored, a float group sum, a CPU reference drift) each fail it, and so do the pre-port sources. test_metal_selftest_motion_v2 passes. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_motion_v2_parity on an Apple GPU (weight and cap, weight, cap, blend, force zero, one frame, tiny/odd/large frames, the five-frame window at 8 and 10 bits, every output at ==) and the motion_v2 case of test_metal_twin_option_parity. * fix(metal): integer vif twin sends frames below 16 pixels to the CPU and keeps tile loads in bounds Row: T-GPU-INTEGER-VIF-MIN-DIM-TWINS-2026-09-29 (Metal). What was wrong: integer_vif_metal declared no minimum frame size; init() refused only a zero-sized scale 3 (frames below 8 pixels). Every scale of integer VIF reflects its filter taps once, which stays inside the plane only while floor(dim / 2^s) exceeds the tap half-width, so between 8 and 15 pixels the twin read other samples than the CPU, and a model run on such frames used the twin. Its compute kernels also reflected every sample of their 16-wide tile plus halo once, including samples beyond a small scale that no output reads, and one reflection sends those outside the buffer (a 16x16 frame at scale 1 loads index 19 of 8 samples, reflected to -5); a device read outside a Metal buffer is undefined. What the port does (the HIP guard, ADR-1381 on ADR-1324; ADR-1498): vif_metal_min_dim() = 16, derived from the CPU's filter widths (integer_vif.h vif_filter1d_width: scale filters need {9, 10, 12, 16}, decimation filters {5, 6, 8}); check_context_metal() returns -ENOTSUP below it and the twin names the CPU `vif` as context_fallback_name, so model dispatch computes those frames on the CPU; init() refuses a direct request below it with -EINVAL before any device work. The kernels' border index is vmaf_mtl_vif_mirror() in the new core/src/feature/metal/metal_integer_vif_math.h (on metal_portable.h): a fold over the period 2 * (sup - 1), which is the CPU's single reflection for every index that reflection brings into the plane and lands in the plane for any other index, so every tile load stays in bounds and no read tap changes. The arithmetic of the statistic is untouched. init() and close() share one teardown helper (no goto) and submit() encodes scale 0 and each of scales 1-3 in a helper, with the same kernels, order and arguments, and now reports a failed command buffer as -EIO. What proves it here: test_metal_integer_vif_math (host, every host, fast) compiles the same header and checks the fold equals the CPU's reflection for every in-range index of every plane length up to 4096; that from 16 up to 300 pixels (and at 853 and 4096) every tap a compute or decimation output reads is a single reflection; that every compute-tile load lands in the plane (the single-reflection form fails this case, planted and seen to fail); and that 16 is the smallest size whose reads are all single reflections. test_metal_integer_vif_min_dim_contract.py (device-free) pins the minimum's derivation, the context check and CPU fallback, the init() refusal before vmaf_metal_context_new(), the kernel's use of the shared fold and the fold itself; seven planted regressions (no context check, a wrong fallback, no init refusal, a minimum without the decimation filters, a single-bounce kernel mirror, a fold drift, a CPU filter width drift) each fail it, and so do the pre-port sources. test_metal_selftest_integer_vif passes. The header compiles as host C++14; the .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_metal_integer_vif_parity's declares_min_dim, direct_init_rejects_below_min and model_boundary on an Apple GPU. model_boundary also compares the twin at == from 16x16 up, which needs the statistic itself to be the CPU's: integer_vif.metal still forms the gain in fp32 where integer_vif.c uses double (the defect ADR-1432 fixed in vif_sycl), which this row does not cover. * fix(metal): cambi twin takes src_width, src_height and full_ref with the CPU's semantics Row: T-BUG048-GPU-OPTION-PARITY-REMAINDER-2026-09-26, option tables of integer_psnr_hvs_metal, integer_cambi_metal, ssimulacra2_metal and integer_vif_metal. What was wrong: of the four twins, integer_cambi_metal's table differed from the CPU `cambi`'s: it lacked src_width, src_height, full_ref and heatmaps_path, so a request or model with any of them kept the CPU extractor. It also never applied the CPU's source-size checks (an encode and a source scaling in opposite directions, a source below the CAMBI minimum) or the CPU's reciprocal-table bound on the adjusted windows, so a window the CPU refuses would have indexed past that table; and it capped the score with its own clamp (a NaN would have stayed NaN where the CPU's MIN() reports cambi_max_val). integer_psnr_hvs_metal, ssimulacra2_metal and integer_vif_metal already declare the CPU's tables, flags and features and execute their options; they are pinned here unchanged. What the port does (the CPU's own code through cambi_internal.h; ADR-1498): the table is the CPU's but heatmaps_path. Dimensions follow cambi.c::validate_and_setup_dimensions (unset sizes are the picture's, both validated, opposite scaling refused); the encode and source windows come from vmaf_cambi_adjust_window() and are bounded by vmaf_cambi_check_window_fits_lut(); buffers are sized for the larger of the encode and, under full_ref, the source picture, as cambi.c::init does. full_ref runs the same per-picture pipeline (host preprocessing, GPU mask, decimation and mode filter, the CPU's c-values and pooling) on the reference at the source size and window and emits cambi_source and cambi_full_reference = MIN(MAX(0, dist - src), cambi_max_val), as cambi.c::extract does; every cap uses the CPU's MIN() spelling. The local copies of adjust_window_size(), get_mask_index() and the contrast weights are replaced by vmaf_cambi_adjust_window(), vmaf_cambi_mask_index() and vmaf_cambi_contrast_weights(). The per-scale planes rotate through local handles, so s->d_image / d_mask / d_tmp never move; init/close share one teardown helper. The kernels are unchanged. Not done, reported: heatmaps_path. cambi.c writes heatmaps with open_heatmaps() / dump_c_values(), which it keeps static; the twin would need a second copy of them or an export from the CPU extractor, which this port may not edit. A request or model with heatmaps_path keeps the CPU extractor (ADR-1183), and test_twin_cambi of test_metal_twin_option_parity still reports that one difference. What proves it here: test_metal_twin_option_tables_contract.py (device-free) compares the four twins' option tables with their CPU extractors' (read by core/test/metal_option_tables.py, whose bracket helper is now public), the TEMPORAL / --subsample rule and the provided features, accepts exactly the recorded heatmaps_path gap, and asserts that each declared option is executed (psnr_hvs enable_chroma, ssimulacra2 yuv_matrix, vif's three options, cambi's dimension checks, source window, full_ref pass and outputs); eight planted regressions (no full_ref, a faked heatmaps_path declaration, a default drift, a missing provided feature, a TEMPORAL flag on one side only, full_ref not run, a local window helper, a CPU reference drift) each fail it, and so does the pre-port cambi source. The cambi, psnr_hvs, ssimulacra2, vif and twin_option Metal self-tests pass. The .mm passes a clang -fsyntax-only run with stub Foundation/Metal headers. What only the device run proves: test_twin_psnr_hvs, test_twin_ssimulacra2 and test_twin_vif (and their _provides cases) of test_metal_twin_option_parity on an Apple device, test_twin_cambi but for the heatmaps_path difference, and test_metal_integer_cambi_parity unchanged at the default options. * fix(metal): add float_motion's SAD row by row in the CPU's order and blur with the CPU's taps float_motion_metal summed |cur - prev| per SIMD group (simd_sum), the groups per threadgroup and the threadgroups in double on the host. The CPU adds each row left to right into one fp32 accumulator, the rows into another, and divides in fp32; every step rounds, so the block sum is not the CPU's score (1.36e-4 off on the 1080p checkerboards on CUDA with the same reduction). The twin also had no motion_add_scale1, motion_add_uv or motion_filter_size. The port takes float_motion_hip's design (ADR-1404, ADR-1419; CUDA and SYCL ADR-1409 / ADR-1411): the blur kernel stores |cur - prev| of every sample transposed in groups of 64 rows, float_motion_row_sum adds each row in order with one thread per row, and the host finishes each plane through vmaf_float_motion_score_from_row_sads() and adds the planes in double, as motion_score_pair() does. Every per-sample value is a function of the new core/src/feature/metal/metal_float_motion_math.h, valid as Metal Shading Language and as host C++ (ADR-1498): picture_copy()'s conversion, the filter motion_filter_size selects (taps that round to motion_tools.h's floats), convolution_f32_c_s()'s two 5-tap passes with every product and sum rounded, the reflect-101 fold in closed form, the transposed index, and motion_scale_bilinear() / motion_bilinear_interp() for the scale-1 term. The tile loader folds every index into the plane, so no load leaves it for any size. motion_add_uv runs the chain on U and V with the CPU's plane checks. The option table gains motion_add_scale1, motion_filter_size and motion_add_uv at the CPU's positions; the command buffer's status is checked. Proved here: test_metal_float_motion_math compiles the header as C++ and runs the kernels' loops on the host. The fold equals convolution_reflect101() for sizes 1 to 70; the taps equal FILTER_5_s / FILTER_3_s / FILTER_5_NO_OP_s; samples and blurred planes equal picture_copy() and convolution_f32_c_s() (dispatched and scalar) bit for bit on 8 sizes from 3x3 to 161x91 plus the 3-tap 2x2, 2x5 and 7x2 planes, at 8, 10, 12 and 16 bits with filter sizes 5, 0, 3 and 1; the row sums and plane score equal compute_motion() with and without motion_add_scale1 at 640x360, 1920x1080, 333x127, 3x3 and 2x2; and the frame score equals the CPU extractor's VMAF_feature_motion_score through libvmaf in seven cases (8 to 16 bits, motion_add_scale1 + motion_add_uv, motion_add_uv, motion_filter_size 3 and 1). A per-block sum is detected on the 1080p fixture; transposed blur passes, a fused tap, a double plane mean and a single-bounce fold each fail the test. test_metal_float_motion_exact_contract pins the kernels and the host to the header and catches eleven planted regressions. The selftests and Metal contracts stay green. Only the device run proves that the Apple GPU's fp32 + - * / are the correctly rounded operations the specification gives under -fno-fast-math and that the kernels compile (the hosted macOS job): test_metal_float_motion_parity at == and the gate's float_motion cell. No fp32 subnormal is produced on these inputs beyond exact zeros. * fix(metal): emit float_motion's motion3 with the CPU's blend options and indices float_motion_metal provided VMAF_feature_motion_score and VMAF_feature_motion2_score only and declared neither motion_blend_factor nor motion_blend_offset, so a run on the twin lost the CPU's motion3 without a warning and a model reading motion3 kept the CPU extractor. Its motion_force_zero path wrote no motion3 either. The twin now follows float_motion_hip (ADR-1404) and float_motion_cuda: it provides VMAF_feature_motion3_score and declares motion_blend_factor (mbf) and motion_blend_offset (mbo) in the CPU table's order. motion3 is motion_blend_clip() of motion2, through the CPU's motion_blend() from motion_blend_tools.h: at index 0 from the first SAD alone, then the blended min(prev, cur) at index - 1 with motion2; flush() writes the tail motion2 / motion3 of the last SAD at the last index, and motion3 = 0 for a one-frame run; the probe that keeps a repeated flush idempotent now looks for motion3 at the last index, which no collect writes. motion_force_zero writes motion2 = motion3 = 0 (and the debug motion) as motion_append_forced_zero() does. motion_blend_tools.h is included ahead of Foundation, whose MIN is defined only when MIN is not, so the build keeps one MIN and no redefinition warning. Proved here: test_metal_float_motion_exact_contract now also requires the provided motion3, the blend through motion_blend(), two motion3 emissions through motion_blend_clip, the one-frame motion3, the force-zero motion3, motion3 at frame 0 from the first SAD, and the option names in the CPU's order (all but motion_max_val, which the next change adds); seven new planted regressions are each detected. The Objective-C++ passes a syntax check against stub frameworks; the float_motion and twin_option selftests and the Metal contracts stay green. Only the device run proves the values: test_metal_float_motion_parity (test_float_motion_motion3 with the default and the blend options, test_float_motion_one_frame, test_float_motion_force_zero) at ==, and the provides case of test_metal_twin_option_parity. * fix(metal): cap float_motion with motion_max_val and clip the debug score as the CPU does float_motion_metal lacked motion_max_val and wrote the debug VMAF_feature_motion_score unweighted and uncapped, where the CPU writes motion_clip(score) = MIN(score * motion_fps_weight, motion_max_val); its motion2 and motion3 carried no cap either. A request or model with motion_max_val kept the CPU extractor (ADR-1183), and --feature float_motion_metal=motion_max_val=... failed with an unknown option. The twin declares motion_max_val (alias mmxv, default 10000, range 0 to 10000) last, as float_motion.c does, so its option table is now the CPU's: the same nine names, aliases, types, defaults, ranges and flags in the same order. Every motion and motion2 value goes through the CPU's motion_clip() (the fps weight, then the cap), the debug score of frames 1 and on included (frame 0's stays 0, as on the CPU), and motion3 through motion_blend_clip() with the cap, using MIN from motion_blend_tools.h. The design is float_motion_hip's (ADR-1382, ADR-1404) and float_motion_cuda's (ADR-1373). Proved here: test_metal_float_motion_exact_contract requires the option names in the CPU's order with none missing, motion_clip() with the cap, the debug score through it, and the capped blend; four new planted regressions (an unclipped debug score, an uncapped motion2, an uncapped motion3, a missing motion_max_val) are each detected. A field-by-field comparison of the two option tables (names, aliases, types, defaults, ranges, flags) finds no difference. The Objective-C++ passes a syntax check against stub frameworks; the float_motion and twin_option selftests and the Metal contracts stay green. Only the device run proves the values: test_metal_float_motion_parity (test_float_motion_max_val with a midpoint and a zero cap, test_float_motion_fps_weight_debug_score, test_float_motion_max_val_out_of_range) at ==, and test_twin_float_motion of test_metal_twin_option_parity. * fix(metal): compute ciede2000 with ciede.c's fp32-pair arithmetic, summed in raster order ciede_metal computed CIEDE2000 in fp32 with another form of the formula (pow(x, 2.4f), pow(t, 1/3) as the cube root, 7.787 t + 16/116 for the linear Lab branch, hue in degrees) and added the values per threadgroup: the construct that put the CUDA, SYCL and HIP twins up to 1.1e-5 from the CPU (T-GPU-CIEDE-CPU-ARITHMETIC-2026-10-01). The kernel now runs the arithmetic of the SYCL and HIP twins (ADR-1436, ADR-1448): pixel() of core/src/feature/ciede_ff_math.h, ciede.c statement for statement with every fp64 value an fp32 pair, on the Metal primitives of the new core/src/feature/metal/metal_ciede_math.h (fma, correctly rounded / and sqrt under the kernels' -fno-fast-math -ffp-contract=off, exp(log(x) / 3) and exp(0.2 log(x)) as root estimates: Metal has no cbrt). It stores one float per pixel at its raster position; the host adds the plane with ciede_frame_sum() and applies extract()'s score expression. ciede.c's constants are fp64 expressions of the bit depth: the host evaluates make_constants() and passes the result (buffer 8); the two tables of ff_math.h are program constants. The host chroma upscale stays: it reproduces scale_chroma_planes(), and its cost is T-METAL-CIEDE-HOST-UPSCALE-2026-09-29 (RC8). The option table stays the CPU's (none). Metal is C++14-based, has no fp64 type and wants an address space on every reference and program-scope variable, so ff_pair.h, ff_math.h and ciede_ff_math.h gain a VMAF_FF_MSL_SUBSET branch (positional initializers, constants in the backend's program-scope space, address spaces on the table pointers and on the Constants and Tables references, no C++ library and no fp64 host helpers under __METAL_VERSION__); the generator writes ff_math.h's constants through VMAF_FF_CONSTANT and VMAF_FF_PAIR_INIT. For the SYCL and HIP twins the preprocessed output is byte-identical before and after (-E -P): the SYCL extractor and probe (icpx -fsycl, host and device), the HIP kernel (hipcc, device and host) and the HIP probe (clang++ and g++). Proven here: test_metal_ciede_math compiles the same header text on the host and holds the pair functions to the host's extended-precision math and the pixel to the reference's fp64 statements (0 of 600 000 pixels differ at 8, 10, 12 and 16 bits); its _coarse build, with root estimates 2^-18 off, stays inside the same bounds (1 pixel differs); test_metal_ciede_exact_contract (16 planted regressions); test_sycl_ciede_math (host and Arc A380) and the HIP host replay pass on the edited headers. The kernel and the .mm were type-checked only against stub Metal and Foundation headers. Only the device run proves the scores (test_metal_integer_ciede_parity at 1e-9, the report's ciede2000 on four fixtures), that Metal's /, sqrt and fma behave as the specification says, and the hosted macOS build that hex float literals and struct copies out of the constant address space compile. * fix(metal): refuse frames below 17x17 in float_adm_metal init, as the CPU float_adm does float_adm_metal's init() accepted frames below 17x17. The CPU float_adm and the CUDA, SYCL and HIP twins refuse them with -EINVAL through adm_frame_size_check(): below 17 pixels the scale-3 bands of the four-level DWT have a single sample. init_fex_metal() now calls adm_frame_size_check("float_adm_metal", w, h) first, before it reads its state, creates a Metal context or allocates a buffer, as float_adm_cuda.c, float_adm_hip.c and float_adm_sycl.cpp do (ADR-1374, ADR-1420, ADR-1498). Proof on this host: core/test/test_metal_float_adm_exact_contract.py (device-free, suite fast) requires the call and its place ahead of fex->priv, vmaf_metal_context_new() and every buffer, and fails on a planted removal and on the check moved after the context. The Metal self-tests (test_metal_selftest_float_adm, test_metal_selftest_twin_option) pass. Only the device run proves the rest: the .mm is compiled by the hosted macOS job and test_metal_float_adm_parity's test_float_adm_metal_rejects_frames_below_17 runs on an Apple device. Row: T-GPU-FLOAT-ADM-TINY-FRAME-FLOOR-2026-10-01 (stays open until the device report). * test(metal): hold the Metal motion twins to the five-frame window and follow the vif host's table line Two device-free tests still described the Metal twins before the ports: test_motion_five_frame_window listed integer_motion_metal and motion_v2_metal as twins that leave motion_five_frame_window to the CPU, which both now compute (T-METAL-MOTION-BLUR-THEN-DIFF-2026-09-29, item 3 of T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30); and test_hip_vif_log2_table_contract planted its Metal regression on the line the integer vif host had before it moved into a helper (T-GPU-INTEGER-VIF-MIN-DIM-TWINS-2026-09-29). The planted regression is still detected. * docs(metal): record what each ported Metal twin computes and keep its rows open for the tester The Metal guide gains a table of the ported twins: what changed for a user (option tables, exact sums, the vif and float_adm frame floors, the motion SAD score and motion3) and the host test that holds each header against the CPU. core/src/feature/metal/AGENTS.md records the exact designs (ADR-1498) and replaces the notes the ports made stale: motion_fps_weight is applied per frame in collect(), and motion2 / motion3 come from the CPU's vmaf_motion_window_flush(). docs/state.md: each ported row records the port on fix/metal-twins-exact and stays open until a tester's report shows its cases passing on an Apple GPU. Two rows are opened: T-METAL-INTEGER-VIF-FP32-GAIN-2026-10-03 (the twin forms the VIF gain in fp32) and T-GPU-ADM-DECOUPLE-FP32-RECIPROCAL-2026-10-03 (the CUDA and HIP scale-0 decouple takes an fp32 reciprocal; found while porting the Metal decouple, not measured on a device). metal-rows.json maps the new vif row to its report cases. * fix(metal): name the soft-double rounding bit halfway, since half is an MSL type metal_soft_double.h and metal_soft_signed.h copied the SYCL headers' local variables, parameter and struct field named `half`. In Metal Shading Language `half` is a scalar type (Specification 4.1, Table 2.1), so every kernel that includes these headers would fail to compile on macOS; the host builds accept the name and no Linux lane runs the Metal compiler. The field, parameter and locals are now `halfway`; the soft-double contract maps the SYCL field name to it. test_metal_shader_build_contract now scans every .metal file, every Metal header and every header they include for a declaration or member access named after an MSL type or address space (half, bfloat, device, constant, thread, threadgroup, kernel, ...). It finds the seven uses in the previous headers and the planted ones. * fix(metal): compute integer vif's gain terms as the CPU's truncated double products integer_vif.metal formed the VIF gain, sv_sq and the vif_enhn_gain_limit clamp in fp32. integer_vif.c::vif_accumulate_pixel() computes the gain in double and truncates sigma2_sq - g * sigma12 and g * g * sigma1_sq to integers, so an fp32 gain moved those integers by one on a share of the pixels and the scores off the CPU's. The twin now takes the two integers from core/src/feature/metal/metal_integer_vif_gain.h, the Metal spelling of the SYCL twin's sycl_integer_vif_math.h (ADR-1432, ported by ADR-1498): one integer division decides both and, for a sample within the fp64 chain's rounding error of an integer, the reference's fp64 operations are replayed in 64-bit integers (metal_soft_double.h, 32-bit limbs, no 64-bit multiply-high). No kernel names double. The host forms the limit's fp64 parts once (vmaf_mtl_ivif_make_gain_limit) and binds a VmafMtlGainLimit in place of the float4; the host tail already rounds each scale's sums to float as vif_store_residuals() does. Proof here: test_metal_integer_vif_gain compiles the kernel header on the host and compares the replay, the integer evaluation and the selected value with the CPU's lines value by value on 340 000 random, boundary and pixel-window variances (8, 10, 12 and 16 bit) for six gain limits; dropping the replay, the sigma2 boundary zone or the gain zone makes it fail. test_metal_integer_vif_gain_contract pins the construct with planted regressions (fp32 gain, fp32 clamp, dropped replay, double, 64-bit literal, multiply-high, reserved identifier, fp32 limit binding, changed host tail). Only the device run proves the kernel's results (the == cases of test_metal_integer_vif_parity and test_vif_metal_model_boundary) and that the GPU's fp32 division and conversions are what the spec says. * fix(metal): floor float_adm_metal frame sums at the CPU's 1e-10 area limit float_adm_metal floored the frame numerator and denominator at 1e-2 * area / 1080p, the value of an ADM_OPT_SINGLE_PRECISION branch no build defines. compute_adm() (adm.c) and float_adm_cuda's fadm_final_scores() floor at 1e-10 * (w * h) / (1920.0 * 1080.0), so the twin floored sums the CPU keeps (adm_noise_weight = 0 on a flat frame). collect_fex_metal() now uses the CPU's expression. The device-free contract test reads the CPU expression from adm.c, requires the twin to match it, and fails on planted returns of the 1e-2 floor and of a changed reference floor. The score itself is only measured by test_metal_float_adm_parity on a device. Row: T-GPU-FLOAT-ADM-FRAME-SUM-FLOOR-2026-10-01 * fix(metal): run float_adm_metal on the CPU's arithmetic, in the CPU's order float_adm_metal copied the reference's constants (the DWT quant step, the 1/30 and 1/15 weights as fp32, the angle constant), multiplied the gain in fp32, summed the terms per threadgroup with simd_sum and pooled with its own powf code: the constructs that put the CUDA, SYCL and HIP twins away from the CPU (T-GPU-FLOAT-ADM-CPU-ARITHMETIC-2026-10-01). The host now takes the CSF weights, the reduced region, the pooling and the angle threshold from adm_float_reference.h (adm_csf_rfactor_s(), adm_border_s(), adm_pool_bands_s(), adm_decouple_cos_1deg_sq_s()). The kernels are the SYCL twin's design (ADR-1434, ADR-1420): decouple + CSF of both signals, the per-sample terms of the reduced region, and one work item per (slot, row) adding its row left to right in fp32; the host adds the rows in fp32. The per-sample arithmetic is the new metal_float_adm_math.h, sycl_float_adm_math.h on metal_portable.h and metal_soft_double.h: the IEEE fp32 quotient (never a reciprocal, ADR-1442), (cos^2 |o|^2) |t|^2, and the three fp64 expressions as exact fp32 pairs with a 64-bit integer replay. The pair operations are written in the header statement for statement as sycl_exact_fp.h, because ff_pair.h is namespace and designated-initializer C++ that MSL does not accept. Proven here: test_metal_float_adm_math compiles the header as C and holds the three fp64 expressions (10^6 samples x 5 limits, and the replay alone) and a composed scale (decouple, terms, rows, fold, pooling; 8 sizes x 4 option sets, 64 decouple trials) to adm_tools.c bit for bit; six planted regressions (reciprocal quotient, regrouped angle test, fp32 1/30, fp32 1/15, centre tap last, fp32 gain) each fail it. The contract test pins the host calls, the kernels' shape and the header, with planted returns of every removed construct. Only the device run proves that the Metal kernels compile and that the device's / and fma are the correctly rounded ones (test_metal_float_adm_parity). * fix(metal): run float_vif's CPU arithmetic on the device and take every CPU option float_vif_metal filtered with a decimal tap table fixed at kernelscale 1.0, took log2 from the device library, kept vif_sigma_nsq in fp32, summed the terms per threadgroup and refused every other kernelscale; its option table lacked vif_prescale, vif_prescale_method and the per-scale floors. Its scores differ from the CPU extractor's by up to 3.8e-5 on video, which the CUDA, SYCL and HIP twins removed in ADR-1412, ADR-1422 and ADR-1444. The twin now follows the SYCL design (ADR-1422, ADR-1498). The host takes the four filters from vif_get_filter() and passes them as kernel arguments, so no kernel holds a tap and every kernelscale runs (up to 69 taps). The arithmetic is core/src/feature/metal/metal_float_vif_math.h: vif_pixel_statistic_s() and log2f_approx() operation for operation, with the two fp64 expressions as exact fp32 pairs and the reference's fp64 operations replayed in 64-bit integers (metal_soft_double.h) next to a rounding boundary. Four kernels (vertical pass, horizontal pass with the statistic, one thread per row, decimate) replace the tiled kernel; the host adds the row sums in fp32 as vif_statistic_s() does. A prescale that changes the frame size runs on the host through picture_copy() and vif_scale_frame_s(). The option table equals float_vif.c's, flags included, and the per-scale floors are applied. Proof here: test_metal_float_vif_math compiles the header on the host and compares it with vif_statistic_s(), picture_copy() and compute_vif() value by value (statistic, the two fp64 expressions with 2e6 random, 4e5 boundary and 12 witness operands, row sums, the whole pipeline at 8 to 16 bits and kernelscales 0.1 to 4.0); test_metal_float_vif_exact_contract plants the removed constructs. Replaying float_vif.metal itself on the host through an MSL shim equals compute_vif() on every score of 12 frame and kernelscale cases. Only an Apple device proves the kernels' results, the device's fp32 division and fma, denormal handling and the host prescale (test_metal_float_vif_parity, test_metal_twin_option_parity). * fix(metal): give float_adm_metal the CPU's adm_fNsM overrides and adm_skip_aim_scale range float_adm_metal's option table lacked the eight per-scale CSF overrides adm_f1s0..3 / adm_f2s0..3 of float_adm.c and passed -1.0 for each to adm_csf_rfactor_s(), so a feature string or model that sets one ran a different CSF on Metal; its adm_skip_aim_scale range started at -1 where the CPU's starts at 0. test_metal_twin_option_parity would have reported both on the tester's device. The table now declares the eight options with the CPU's alias, default, range and flags and hands them to adm_csf_rfactor_s(), and adm_skip_aim_scale has the CPU's range. adm_csf_mode stays default-only, as on the CUDA, SYCL and HIP twins (ADR-1316): the device test now accepts that one flag, as it already did for float_ssim's scale. float_vif_metal runs every vif_kernelscale since its port, so test_gpu_option_value_capability_contract lists it as a full-range option. test_metal_twin_option_tables_contract compared four Metal twins' tables with their CPU extractors' from the sources; it now compares all seventeen, with the two recorded gaps (cambi's heatmaps_path, float_adm's adm_csf_mode), and fails on a planted removal of adm_f2s3. metal_option_tables.py reads a sized table (float_psnr.c's options[2]) and a file without one. * fix(metal): form integer_ssim_metal's per-pixel term as the CPU's fp64 term in integers Wrong: the twin reduced fp32 terms on the device, so its score differed from the CPU's `ssim` by rounding and by summation order. Now: the kernel forms integer_ssim.c::ssim_reduce_row_range()'s fp64 term operation for operation on soft-signed values held in 64-bit integers (metal_integer_ssim_math.h, the design of the SYCL twin, ADR-1443; CUDA ADR-1424, HIP ADR-1438), stores every term's bit pattern unreduced, and the host adds the plane in calc_ssim()'s raster order as a double. enable_db and clip_db go through the CPU helpers; the option table equals the CPU's. Proof here: test_metal_integer_ssim_math holds the term and the whole twin against the CPU at ==; test_metal_integer_ssim_exact_contract pins the constructs with planted regressions. Only a Metal device proves test_metal_integer_ssim_parity. * fix(metal): form float_ssim_metal's window terms as the CPU's and add them in raster order Wrong: the kernel summed the taps in plain fp32, divided l, c and s in fp32 and forced 1 on a zero denominator, then reduced the scores per threadgroup. An identical flat 64x64 frame scored +inf dB with enable_db where the CPU scores 72.247198959355487 dB, and every textured frame differed in the last bits. The scale error also named a pin (`float_ssim_metal:scale=1`) the CLI does not accept. Now: the convolution taps are exact fp32 pairs rounded once, as the CPU's fp64 sum of fp32 products rounds; the window values are the CPU's fp32 operations in its order with no forced 1; lv and cv are the CPU's fp64 quotients computed on values held in 64-bit integers (metal_ssim_terms.h, the design of the SYCL twin, ADR-1463). Every window's term (or lv, cv and sv under enable_lcs) is stored at its raster position and the host adds the plane in iqa_ssim()'s order (ADR-1464 for CUDA), divides by the window count and rounds the mean to fp32. The planes come from picture_copy(). The scale error says `float_ssim_metal=scale=1`. Scale 1 only stays. Proof here: test_metal_float_ssim_math holds the window terms against ssim_tools.c bit for bit and the whole twin against compute_ssim() at ==, including the flat 72.247198959355487 dB case; the source contract pins the constructs with planted regressions. Only a Metal device proves test_metal_float_ssim_parity. * fix(metal): run float_ms_ssim_metal's decimation and windows as the CPU's and add terms in raster order Wrong: the decimation was one 9x9 kernel with unfused partial sums, the window sums were plain fp32, l, c and s were fp32 quotients, and the scores were reduced per threadgroup, so the twin differed from the CPU in the last bits of the scale means. Now: the decimation is ms_ssim_decimate.c's two separable passes with one explicit fma per tap in the CPU's tap order (metal_ms_ssim_math.h, ADR-1414 for SYCL, ADR-1403 for CUDA); the window sums, the fp32 window values and the fp64 lv and cv quotients on values held in 64-bit integers are metal_ssim_terms.h's, shared with float_ssim_metal as the SYCL twins share sycl_ssim_terms.h. Every window's lv, cv and sv is stored at its raster position of its (plane, scale) region (ADR-1465 for CUDA, ADR-1466 for SYCL); the host adds each region in index order, rounds each mean to fp32 and combines the scales as ms_ssim.c does, with fabs() on all three means. The planes come from picture_copy(). The option table is the CPU's. Proof here: test_metal_float_ms_ssim_math holds the decimation against ms_ssim_decimate_scalar() and ms_ssim_decimate() and the whole twin against compute_ms_ssim() at ==, score and all fifteen means; the source contract pins the constructs with planted regressions. test_metal_ms_ssim_options_contract now looks for the combine call where it looked for pow(). Only a Metal device proves test_metal_float_ms_ssim_parity. * fix(metal): rename float_vif's half_unit local, since half is an MSL type metal_float_vif_math.h declared `const float half` in the rounding-boundary test of the fp32 pair sum; `half` is a Metal Shading Language scalar type, so float_vif.metal would not compile on macOS. test_metal_shader_build_contract found it once the float_vif port landed on this branch. * test(metal): follow the exact float_adm and float_ssim twins in two device-free contracts test_float_adm_csf_upstream_contract pinned float_adm_metal's own copy of dwt_quant_step() to upstream's float arithmetic (ADR-1489). The port deleted the copy: the twin takes its CSF weights from adm_csf_rfactor_s(). The contract now requires that call and fails on a planted local copy or a replaced call. test_nonfinite_collector_wiring required raw L/C/S isfinite() guards in float_ssim.metal, which stood in for the forced 1 on a zero denominator. The port forms each window's terms as the CPU's through metal_ssim_terms.h and the host emits through vmaf_ssim_emit_scores_named(), which fails closed on a non-finite mean; the test now requires the header, and the forced-1 forms stay forbidden. core/src/feature/metal/AGENTS.md: the float_adm quantisation-step note now says the weights come from the CPU's routine. * docs(metal): record the float_adm, float_vif, vif gain and ssim-family ports and keep their rows open docs/state.md: the float_adm floor and arithmetic, integer ssim, float_ms_ssim, float_vif and integer vif gain rows record their Metal port, its host proof and the parity test whose passing report closes them; the BUG048 row records float_adm_metal's new options and the all-twin table contract; the twin parity gaps row records float_ssim (items 1 and 4). The Metal guide's table gains the five twins and states the three options a Metal twin runs differently from the CPU (cambi's heatmaps_path, float_adm's adm_csf_mode, float_ssim's scale). core/src/feature/metal/AGENTS.md records the exact designs of these twins and the MSL name rule. A changelog fragment lists what a Metal user sees change; docs/rebase-notes.md records the shared headers and host routines a rebase must keep. * test(metal): compare the recorded option-table gaps with set order normalised test_metal_twin_option_tables_contract matched float_adm's recorded adm_csf_mode gap against the repr of a frozenset of flags, whose order follows the interpreter's per-run string hash, so the fast suite failed on some runs. Each difference now lists a frozenset's items sorted before it is compared; the test passes under PYTHONHASHSEED 1 to 5. * fix(metal): form float_moment's rounded second-moment sum past 2^53 units as the CPU does float_moment_metal added the 16-bit float squares as exact uint64 units and the host converted the total once. That equals the CPU's second moment only while the plane sum stays at most 2^53 units of 1 / scaler^2 (16-bit frames up to 2^21 pixels). moment.c adds one float square per pixel into a double in raster order, so past 2^53 every add rounds and the exact total, rounded once, is another number. On a 16-bit frame of 3840x2160 the twin returned a different ref2nd / dis2nd than the CPU, and test_float_moment_16bit_past_2_53_exact (ADR-1497) fails on an Apple device until the twin forms the CPU's sum. The port follows the CUDA / HIP / SYCL twins (ADR-1497, ADR-1498). A frame that can pass 2^53 units (vmaf_mtl_msum_may_round()) enqueues five more kernels on the frame's encoder, after the frame kernel and before the one wait: float_moment_plane_sums (the four exact sums from the partials), float_moment_row_totals, float_moment_row_plans, float_moment_row_units (increments composed over lanes with the ordered tree) and float_moment_ordered_totals (one checked walk per plane with the run and term fallback). Each of the last four returns while the plane's exact sum is at most 2^53, where it is the CPU's sum. collect() takes the two second-moment sums from the walk and fails closed without them; the host refuses a pipeline that cannot run 256 lanes. Integers only, strict FP unchanged. metal_float_moment_sum.h is a copy of float_moment_sum.h and the four ordered_sum.h functions it uses, not an include: every pointer parameter needs a Metal address space, and a hook macro on each would change the preprocessed text of the CUDA, SYCL and HIP twins that compile the shared header, which stays untouched. The copy is statement for statement under a rename rule written at its top (and `half` -> `half_way`, a Metal type name); test_metal_float_moment_exact_contract.py maps it back and requires every function body and constant to equal the shared one, and test_float_moment_sum_contract.py runs that check on the shared header it edits, so a change to one without the other fails. Proof, on the host: the model of test_float_moment_sum.c moved to float_moment_sum_model.h and runs against both headers; test_metal_float_moment_sum lays the copy out as the kernels do (256 lanes, ordered tree, batches of 256 rows) and matches picture_copy() + compute_2nd_moment() bit for bit on the ADR-1497 parity cases, 3840x2160 noise and bright frames, a frame whose sum is 2^54 plus seven ones and 7680x4320 (past 2^56), also with four kinds of wrong plans. Planted mutations fail it: dropped term fallback, reordered tree, ignored parity of the sum, tie to the odd value, unchecked plan. The device-free contract has 16 new planted regressions (exact sum used past 2^53, missing walk, missing store, fp64, ULL literal, a device reduction in place of the walk, a copy that no longer follows the shared header). * docs(metal): record float_moment's past-2^53 port and measure it in the tester's report docs/state.md: T-GPU-FLOAT-MOMENT-16BIT-SQUARES-2026-10-02 records the past-2^53 port (ADR-1497 on Metal) and its host proof, and stays open. metal-rows.json adds test_float_moment_16bit_past_2_53_exact to the row's cases, so the tester's report measures it. The Metal guide, the Metal agent page and the changelog fragment describe the five extra kernels and the 256-thread pipeline requirement. * refactor(metal): bring the Metal port tests and headers to the clang-tidy cpu standard The required Tidy Ratchet check (cpu lane, clang-tidy 22) failed on #1921: 151 findings in 18 files with an allowance of zero. Re-measured with scripts/dev/tidy-lane.sh on the 19 changed host test translation units: 151 before, 0 in every file outside the baseline after, and no baseline file above its count. Refactored, no change of assertions or coverage: long test functions and run_tests() split into helpers (test.h's mu_assert_msg carries the messages), braces, isolated declarations, designated initializers in the C++ tests (C++23), using aliases, loop conversions, const locals, bit-pattern compares instead of memcmp on floats, leak-free early exits where the analyzer saw a malloc without a free, <c...> headers, and index arithmetic widened before the multiplication. Plane and Filters in test_metal_float_vif_math.cpp lose their member functions for free functions. Metal headers: the typedef structs and the brace initialisers stay as they are because the headers compile as MSL and as host C; each carries a NOLINTNEXTLINE citing ADR-1498 (modernize-use-using, modernize-use-designated-initializers). Other suppressions, each with its reason inline: - bugprone-incorrect-roundings in metal_float_motion_math.h: the reference's own (int)(width * 0.5 + 0.5) of motion.c::vmaf_image_sad_c(). - bugprone-misplaced-widening-cast in metal_integer_adm_math.h: the reference's (int64_t)(1u << (14 + k_shift)) of integer_adm_kernels.h. - clang-analyzer-core.BitwiseShift in metal_integer_vif_gain.h, metal_soft_double.h and metal_soft_signed.h: the shift count is bounded by a count-leading-zeros result the analyzer cannot see (17..49, 30..52 and 0..52). - modernize-use-std-numbers on the log2 polynomial coefficient of vif_tools.c, which is not log2(e). - performance-enum-size on vif_tools.h's enum: a plain C header, an enumeration has no underlying type before C23. metal_float_vif_math.h's NaN test x != x became a sign-free bit-pattern test (same result, no redundant-expression finding). convolution_internal.h takes three const locals; no arithmetic changed. No real defect found: every flagged expression in the headers mirrors the CPU or SYCL statement it is a port of. * docs(adr): regenerate the tag pages and the citation registry after the rebase onto ADR-1499 The rebase took master's generated files at each stop; this regenerates the by-tag pages and scripts/ci/source-adr-citations.json at the tip, with ADR-1498 counted next to ADR-1499. * fix(metal): add float_psnr_metal's rows in the CPU's order so frames past 2^53 units match (ADR-1499) float_psnr.c adds each row's sum into one double, and past 2^53 units of 1 / scaler^2 those adds round. float_psnr_metal summed each 16x16 threadgroup exactly and rounded the exact frame total once, which is another number on a 16-bit frame past that point; since #1922 the shared fixture's test_float_psnr_16bit_past_2_53_exact compares it with ==. As the CUDA, SYCL and HIP twins now do (ADR-1499), a threadgroup covers 256 pixels of one row and the host adds each row's segments exactly in 64 bits and the rows into a double in order, through vmaf_float_psnr_row_noise() of core/src/feature/float_psnr_rows.h. The kernel is unchanged: its group index is already row-major. test_metal_float_psnr_exact_contract.py requires the row dispatch and the helper and fails on a frame total rounded once and on a 16x16 threadgroup. metal-rows.json adds the past-2^53 case to the row. Not compiled for Metal on this host. * test(metal): hold the vif gain header to the CPU only where the CPU's conversion is defined test_metal_integer_vif_gain failed on the hosted Ubuntu ARM clang job and under qemu-aarch64 with a clang cross build: on some samples the CPU's `int32_t sv_sq = sigma2_sq - g * sigma12` converts a double below INT32_MIN, which is undefined in C. x86 returns INT32_MIN (0 after the MAX()), clang for aarch64 keeps the low 32 bits of a 64-bit conversion, and the header, like the SYCL twin, returns x86's 0. Random variances reach that range, and so do windows drawn through the CPU's own fixed-point moments (1869 of 100000 at 8 bits). The test now draws both kinds of sample inside the defined range and passes on x86 and under qemu-aarch64 (clang 23 cross build). The CPU defect is recorded as T-INTEGER-VIF-SV-SQ-CONVERSION-UB-2026-10-03 (open, RC3): an aarch64 clang build scores those pixels differently from x86. * test(metal): key the shader contract's sources by POSIX paths so it passes on Windows test_metal_shader_build_contract keyed each scanned source by str(path.relative_to(ROOT)), which is a backslash path on Windows, so the Windows ARM64 MSVC job could not find core/src/feature/metal/metal_soft_double.h among them. The key is the path's as_posix() form. * fix(metal): keep the integer ssim header and two host tests free of undefined behaviour under UBSan The hosted ASan + UBSan job failed on three tests of this branch. metal_integer_ssim_math.h computed the integer product sums before it knew they were in range, as the SYCL header does, and discarded them otherwise; above 2^53 vmaf_mtl_signed_from_exact() then shifts by a negative count, which C leaves undefined (UBSan: shift exponent 4294967292). The term now takes the integer path only when vmaf_mtl_issim_products_are_exact() holds and the rounded path otherwise; the values are the same, and the selection line the contract pins is unchanged. test_metal_integer_adm_math's model of the former fp32 decouple and test_metal_integer_vif_gain's model of the former fp32 gain converted floats beyond the int range. The adm model saturates there, as the device's conversion does; the vif count treats such a sample as differing without converting. The sanitizer…
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
This brings the HIP twins onto the CPU's arithmetic for the four RC3 rows the SYCL twins already closed, and it has now run on an AMD GPU.
motion_hipblurred each frame instead of the frame difference and was 1.26e-5 off the CPU; it now runs the diff-first kernelmotion_v2_hipalready used, with no host wait insubmit()(ADR-1377). The motion tile loads and the integer ADM scale-0 vertical DWT clamp their reflected rows, andvif_hiphands frames below 16 pixels to the CPU (ADR-1381).psnr_hip,integer_ssim_hip,float_ssim_hipandfloat_motion_hiptake the CPU option tables (ADR-1382).On
ryzen-4090-arc(Ryzen 9950X3D iGPU, gfx1036, ROCm 7.2.4) every HIP device test passes,motion_hipequals the CPUmotionon every frame, and the parity gate passes every HIP cell. The device run also found a crash (both HIP motion twins died on frame 0 withmotion_force_zero, fixed here), two stale test expectations (fixed), a throughput cost of the staged upload on this iGPU (recorded), and a platform fault: the gfx1036 now and then never runs a run of a stream's commands, on master as well (recorded with a probe, explicitly deferred).T-HIP-MOTION-BLUR-THEN-DIFF-2026-09-29,T-HIP-FLOAT-MOTION-TILE-OOB-2026-09-30andT-HIP-MOTION-FORCE-ZERO-NULL-SUBMIT-2026-09-30close. The HIP halves of the ADM-DWT, VIF-min-dim and BUG048 rows are verified; #1637 landed their CUDA halves, soT-CUDA-HIP-ADM-DWT-VERT-TINY-HEIGHT-OOB-2026-09-29closes here as well and the other two rows stay open for Metal only.Device run on gfx1036 (2026-10-01)
Build:
meson setup build-hip core -Denable_hip=true -Denable_hipcc=true -Dhip_gfx_targets=gfx1036 -Db_lto=false && ninja -C build-hip, host build onryzen-4090-arc, every device run under the gfx1036 lock.run_meson_test.py -- -C build-hip --num-processes 1with the 15 tests the rows nameMemory access fault by GPU noderun_meson_test.py -- -C build-hip --suite gpu --num-processes 1test_hip_float_ssim_parity_large(float_ssim_hipdoes not decimate yet, pre-existing)motion_hipvs CPUmotion(motion row)--precision=maxmotion2/motion31.26e-5 apart on master 10f27ef, identical on this branch; identical on all 9600 frames of 20 runs of the pair looped ten timescross_backend_parity_gate.py ... --backends cpu hip --features float_ssim float_ssim_lcs psnr motion_v2 vifpsnr_hipoptions (BUG048 row)psnr_hip=enable_mse=true:enable_apsnr=true:reduced_hbd_peak=true:min_sse=0.5vs the CPUpsnr_*/mse_*andapsnr_y/cb/cridenticalfloat_motion_hip3x3 / 17x17 (tile row)ffmpegcrops,--feature float_motion_hipvif_hipminimum (VIF row)test_hip_vif_min_dim test_hip_vif_parityvif, from 16 up within 5e-5test_hip_adm_dwt2_rows test_hip_adm_tiny_frames test_hip_adm_parity test_hip_adm_small_border--model version=vmaf_v0.6.1, Netflix pairRe-run on the head rebased onto master
bcb45e6cb(which includes #1637): the device suite (52 OK, same single skip marker), the parity gate (1.0e-5, 1.0e-5, 0, 0, 1.0e-6),motion_hipagainst the CPU (0 on 48 frames), thepsnr_hipoption run (identical, also under--subsample 2), themotion_force_zeroruns (exit 0) and the default model (76.667849 against 76.667831) give the numbers above. The fast suite passes on the HIP build (247 OK) and on a CPU-only build (193 OK, 1 skipped).Found on the device
motion_force_zero(T-HIP-MOTION-FORCE-ZERO-NULL-SUBMIT-2026-09-30, closed). libvmaf pickssubmit()/collect()beforeinit()runs; the twins cleared them ininit(), so frame 0 called a NULLsubmit()(exit 139 on master). They now keep a no-opsubmit()and acollect()that writes the CPU's zeros;test_integer_motion_force_zerocovers it. The CUDAmotiontwins crashed the same way on the RTX 4090; fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637 fixed that in the engine, which now initialises an extractor before it picks the dispatch path. The HIP twins keep their asynchronous pair and do not depend on that order.test_hip_float_motion_parityexpected the flushed tail to be 1.5 times the debug score, but since ADR-1382 both carrymotion_fps_weightonce;test_hip_twin_option_parity's one-frame case looked motion2/3 up under their JSON names.(median t(22) - median t(2)) / 20, 3 interleaved runs:motion_hip14.25 ms/frame on master, 12.95 here;motion_v2_hip10.17 on master, 13.24 here. The kernel is not the cause (hipEvent timing of the two code objects on the same planes: 9.76 vs 10.03 ms): with the waitingvmaf_hip_picture_upload()ininteger_motion_sad_hip.cboth twins run at 10.7 ms/frame. On an integrated GPU the runtime copies a pageable picture without a host copy, so the staged path's host copy into pinned memory, which runs while the GPU idles, costs more than the wait it removes. With several twins in one process and with thevmaf_v0.6.1model the two uploads are within noise. The design stays as ADR-1377 decided (the device-resident cambi/SpEED PR builds on it, and a discrete GPU is unmeasured); the numbers are inT-HIP-UPLOAD-WAIT-THROUGHPUT-2026-09-19and the HIP guide.T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01, explicitly deferred). Over repeated 480-frame runs,vif_hipreports a frame whose four scales carry the sums of that frame and the one before, a scale with numerator and denominator 0 (invalid ratio), or wrong scales 1 to 3, about once per 10^4 frames, on master as well. A previous handoff commit on this branch replaced the accumulatorhipMemsetAsyncwith a zeroing kernel; that did not help (placed before or after the upload), so this branch keeps master's memset.scripts/dev/hip_dispatch_drop_probe.hipreproduces the fault without vmafx: five runs of 100000 frames lost 382 of 6.0 million kernel dispatches, each time a run of commands from the start of a frame.HIP_FORCE_DEV_KERNARG=0,HSA_ENABLE_INTERRUPT=0,AMD_DIRECT_DISPATCH=0and a spinning wait do not stop it;AMD_SERIALIZE_KERNEL=3andHSA_ENABLE_SDMA=0make it more frequent. The verification above therefore repeats runs and counts mismatches.What changes, per row:
T-HIP-MOTION-BLUR-THEN-DIFF-2026-09-29(closed):integer_motion/motion_score.hipandinteger_motion_hip.hare deleted; both motion twins callvmaf_hip_motion_sad_submit()(core/src/feature/hip/integer_motion_sad_hip.{h,c}), the only loader and launcher ofinteger_motion_v2/motion_v2_score.hip.motion_hipkeeps a raw-luma ping-pong. Its debugmotionscore now carriesmotion_fps_weight/motion_max_vallike the CPU's, and a one-frame run reportsmotion3 = 0. The luma is copied into a pinned buffer on the host and uploaded without waiting (vmaf_hip_picture_upload_staged(),core/src/hip/picture_hip.{h,c}), socollect()is the only host wait of a frame.T-CUDA-HIP-ADM-DWT-VERT-TINY-HEIGHT-OOB-2026-09-29(HIP half verified; closed, the CUDA half landed in fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637):adm_dwt2_load_column()serves scale 0 only; replaying every thread row of its launch for heights 1 to 8192 shows the single reflection leaves the plane only below 9 rows, so the 17x17 minimum kept every accepted frame inside. The load now readsadm_dwt2_source_row()(integer_adm/adm_dwt2_rows.h), clamped into the plane, which is the identity for every accepted frame.T-GPU-INTEGER-VIF-MIN-DIM-TWINS-2026-09-29(HIP half verified; open for Metal):vif_hipdid not fault (itsmirror2_i()clamps), but below 16 pixels it read other samples than the CPU. It now declares a 16-pixel ADR-1324 context check with the CPUvifas fallback and refuses direct requests below that at init, asvif_sycldoes.T-BUG048-GPU-OPTION-PARITY-REMAINDER-2026-09-26(HIP part verified; open for Metal):psnr_hiptakesmin_sse,enable_mse,reduced_hbd_peakandenable_apsnrthrough the CPU'spsnr_score.h, withapsnr_*from a newflush().integer_ssim_hiptakesenable_db/clip_db, andfloat_ssim_hiptakesenable_lcs(a pass-2 kernel variant that reduces L, C and S on the device),enable_dbandclip_db.float_motion_hiptakesmotion_max_valand appliesmotion_clip()to every score it emits. Withenable_db, identical frames report what the CPU reports, including 72.247 dB (1 - 2^-24) on a flat identical frame forfloat_ssim_hip.The HIP twins now share
vmaf_hip_rc_to_errno().core/src/hip/common.hdeclared it, but nothing defined it;kernel_template.cdoes now, and four private copies are gone.Review fixes (second pass, before the device run)
Two reviews of the first head raised eight items; each fix came with a device-free check that fails on the unfixed code (
test_hip_kernel_source_contract.py: 26 planted regressions).float_motion_hipread outside its input plane at extents 3 to 9 and 17. It now uses the same clamped index asmotion_v2_score.hip(hip_tile_index.h).T-HIP-FLOAT-MOTION-TILE-OOB-2026-09-30, closed.float_ssim_hipforced identical windows to 1, which the CPU does not do. Pass 2 now forms(l * c) * sin double from the CPU-typed L, C and S and the host rounds each frame mean to fp32 likeiqa_ssim().integer_ssim_hipon 1x1 / 2x2 identical frames reports+infwhere the CPU reports 156.54 / 159.55 dB:T-HIP-INTEGER-SSIM-TINY-IDENTICAL-DB-2026-09-30(RC3, low).motion_v2_hipstored the raw SAD; it now storesMIN(score * motion_fps_weight, motion_max_val)likeinteger_motion_v2.c, and a one-frame run emitsmotion2_v2/motion3_v2.psnr_hipwas not TEMPORAL, so--subsample N > 1summedapsnr_*over 1/N of the frames.cross_backend_parity_gate.pyandcross_backend_vif_diff.pytakehip, with afloat_ssim_lcscell.motion3moving average at two frames (the twin skips(x + x) / 2, which equalsxexactly).motion_hipdefaulteddebugto true and never emittedVMAF_integer_feature_motion_sad_score. Opened, not fixed:float_motion_hipstill has nomotion3and lacks five CPU options (T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30, RC3).vif_hipreturns the scaffold-ENOSYSbefore its minimum-size check (ADR-1264);.standards-baseline.jsonre-recorded with the pinned engine (185 to 182 findings).Type
fix— bug fixtest— test-onlyfeat— new featureperf— performance improvementrefactor— no behavior changedocs— documentation onlybuild/ci— tooling / infraport— cherry-pick from upstream Netflix/vmafsycl/cuda/simd— backend-specificChecklist
make format && make lintis green locally. What ran for this head (rebased onto masterbcb45e6cb):meson setupon the mergedcore/test/meson.build,make lint-md(0 issues in the 26 changed pages),make docs-fragments-check,make verify-all(audit, context, HISS replay: pass),scripts/ci/check-state-md-rows.sh, and the deliverables, state.md-touch and silent-revert gates againstbcb45e6cb. The pre-commit hooks (clang-format 23.1.1 among them) and HIP-lane clang-tidy 22.1.8 (tidy-ratchet.py --lane hip --only, 0 findings in the four C files the device fixes touch) ran on the previous head; the rebase changes no HIP source or test, onlycore/src/feature/hip/AGENTS.md. The full CPUmake lintdid not run; this head changes no CPU source.python3 scripts/ci/run_meson_test.py -- -C build-hip --suite gpu --num-processes 1on the gfx1036, 52 OK (see the device table)./cross-backend-diffand the worst ULP is ≤ 2. Measured on the gfx1036 with the parity gate and per-feature comparisons (table above):motion_hip,motion_v2_hipandpsnr_hipidentical to the CPU,vif_hip1.0e-6,float_ssim1.0e-5..c/.cpp/.cu/.h/.hpp, it has the appropriate license header (seeCONTRIBUTING.md). The newscripts/dev/hip_dispatch_drop_probe.hipcarries it too.!orBREAKING CHANGE:and the migration path is documented below. Not a breaking change.docs/adr/_index_fragments/<NNNN-slug>.mdand the slug is appended todocs/adr/_index_fragments/_order.txt— do not editdocs/adr/README.mddirectly (regenerated byscripts/docs/concat-adr-index.sh; see ADR-0221).Bug-status hygiene (ADR-0165)
docs/state.mdupdated in this PR. Closed:T-HIP-MOTION-BLUR-THEN-DIFF-2026-09-29,T-HIP-FLOAT-MOTION-TILE-OOB-2026-09-30,T-HIP-MOTION-FORCE-ZERO-NULL-SUBMIT-2026-09-30, andT-CUDA-HIP-ADM-DWT-VERT-TINY-HEIGHT-OOB-2026-09-29(HIP half here, CUDA half in fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637). HIP halves verified, rows open for Metal:T-GPU-INTEGER-VIF-MIN-DIM-TWINS-2026-09-29,T-BUG048-GPU-OPTION-PARITY-REMAINDER-2026-09-26.T-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30(opened by fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637) records which of its HIP items this PR already fixes; thefloat_ssim_hipscale hint stays open there. Opened:T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01(explicitly deferred),T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30,T-HIP-INTEGER-SSIM-TINY-IDENTICAL-DB-2026-09-30(RC3).T-HIP-UPLOAD-WAIT-THROUGHPUT-2026-09-19carries the staged-upload measurements.Netflix golden-data gate (ADR-0024)
assertAlmostEqual(...)score in the Netflix golden Python tests.Cross-backend numerical results
gfx1036, ROCm 7.2.4, Netflix 576x324 pair unless noted,
--precision=max, largest absolute difference against--backend cpu:Rare single-frame outliers from
T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01appear on master and on this branch alike: in one interleaved session on a quiet host,vif_hiphad one wrong frame in 2 of 20 master runs and 3 of 20 branch runs (480 frames each);motion_hiphad none in 20 runs,psnr_hipnone in 60,motion_v2_hip2 frames in 60.Performance
BBB 3840x2160,
(median t(22) - median t(2)) / 20over 3 interleaved runs per variant, single feature, gfx1036, load average 8 to 9:motion_hipgets faster because it no longer runs its own blur pipeline;motion_v2_hipgets slower because of the staged upload, as explained under "Found on the device". Multi-twin runs (motion_v2_hip+psnr_hip+float_motion_hip: 37.8 / 37.5 / 38.1 ms/frame;motion_hip+motion_v2_hip: 23.6 / 24.6 / 22.0) and thevmaf_v0.6.1model (216 / 217 / 217, ADM on the CPU) are within noise across the three variants.Deep-dive deliverables (ADR-0108)
docs/research/1377-hip-rc3-cpu-parity.md: motion order and SAD equivalence, host replays of the tile and ADM row loads, the VIF bound, identical-window SSIM exactness, and a "Device results (gfx1036, 2026-10-01)" section with the measurements, the staged-upload analysis and the dropped-dispatch probe.## Alternatives consideredin ADR-1377, ADR-1381 and ADR-1382.AGENTS.mdinvariant note —core/src/feature/hip/AGENTS.md: the shared diff-first motion kernel and launcher, staged uploads, tile and ADM-row clamps, thevif_hip16x16 minimum, the CPU option tables, never clearingsubmit/collectundermotion_force_zero, and how the gfx1036 command loss shows up;scripts/dev/AGENTS.md: keep the probe standalone.docs/state.md.changelog.d/fixed/hip-motion-diff-first.md,changelog.d/fixed/hip-integer-tiny-frame-guards.md,changelog.d/changed/hip-twin-cpu-option-parity.md,changelog.d/fixed/hip-motion-force-zero-async.md,changelog.d/added/hip-dispatch-drop-probe.md.docs/rebase-notes.md, "ADR-1377 / ADR-1381 / ADR-1382 — HIP RC3 CPU parity", including themotion_force_zerocallbacks and the probe.Reproducer
On an AMD host (
ryzen-4090-arc, gfx1036):meson setup build-hip core -Denable_hip=true -Denable_hipcc=true -Dhip_gfx_targets=gfx1036 ninja -C build-hip python3 scripts/ci/run_meson_test.py -- -C build-hip --suite gpu --num-processes 1 python3 scripts/ci/cross_backend_parity_gate.py --vmaf-binary build-hip/tools/vmaf \ --reference testdata/ref_576x324_48f.yuv --distorted testdata/dis_576x324_48f.yuv \ --width 576 --height 324 --backends cpu hip \ --features float_ssim float_ssim_lcs psnr motion_v2 vif Y=python/test/resource/yuv build-hip/tools/vmaf -r $Y/src01_hrc00_576x324.yuv -d $Y/src01_hrc01_576x324.yuv -w 576 -h 324 \ -p 420 -b 8 --no_prediction --backend hip --feature motion_hip=motion_force_zero=true \ --frame_cnt 3 --json -q -o /tmp/fz.json # exit 0; exit 139 on master hipcc -O2 --offload-arch=gfx1036 scripts/dev/hip_dispatch_drop_probe.hip -o /tmp/probe /tmp/probe 100000 12 0 # bad_frames=0 lost=0 on a healthy stack; not on this gfx1036Device-free, as before:
python3 scripts/ci/run_meson_test.py -- -C build-hip --suite=fastandpython3 core/test/test_hip_kernel_source_contract.py.CI note
The previous head failed
Windows MinGW64on a 30 s timeout oftest_gpu_public_header_docs, which runs the doxygen that the Windows runner happens to have onPATH(C:\Strawberry\c\bin\doxygen.exe) and took 3.4 s on master's run.fix/cambi-short-frame-oob(#1642) hit the same timeout; this PR does not touch public headers or that test. #1642 has since raised that timeout to 120 s on master, so this branch's own commit for it dropped out of the rebase.Known follow-ups
submit/collectundermotion_force_zero; since fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637 the engine handles that. A comment incore/src/libvmaf.c(init_before_dispatch()) still lists the HIP twins among them.T-HIP-UPLOAD-WAIT-THROUGHPUT-2026-09-19); the other HIP extractors still wait on the host insubmit().T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01: re-run the probe after a ROCm or amdgpu update.T-HIP-FLOAT-MOTION-MOTION3-OPTIONS-2026-09-30andT-HIP-INTEGER-SSIM-TINY-IDENTICAL-DB-2026-09-30.bcb45e6cb(after fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637, perf(cuda): run cambi and SpEED entirely on the device #1639, fix(speed): size speed_temporal and speed_chroma frame buffers correctly #1643, perf(cuda): read the device pictures directly in psnr_hvs_cuda #1646, fix(sycl): accept bare native images in AOT image check #1651, fix(ssimulacra2): refuse 4:0:0 input at init instead of reading a NULL plane #1654 and fix(cambi): keep the c-values walks inside short frames #1642). Where fix(cuda): match the CPU motion order, option tables and tiny-frame guards #1637 and this PR describe the same paragraph or row (docs/metrics/{features,motion,psnr,ssim,vif}.md, threedocs/state.mdrows), both backends' statements are kept, and one new commit aligns the HIP notes with the engine'smotion_force_zerofix and replaces "not yet measured on an AMD device" in the metric pages with the gfx1036 results.perf/hip-rc3-device-resident(perf(hip): run cambi and SpEED entirely on the device #1661, cambi + SpEED on the device) is rebased onto this head.float_ssimscale, the fp-contract port, psnr_hvs, cambi and SpEED.