Repository navigation
fix(metal): port the Metal twins to the exact designs of their CUDA, HIP and SYCL twins (ADR-1498) - #1921
Merged
Conversation
lusoris
marked this pull request as ready for review
October 3, 2026 13:10
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 14:05
94fd613 to
8b073b6
Compare
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 14:18
e1e1dad to
c82f3b5
Compare
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
| VmafMtlFadmGainLimit out; | ||
| memset(&out, 0, sizeof(out)); | ||
| out.value = (float)limit; | ||
| out.is_float = ((double)out.value == limit) ? 1 : 0; |
| * every product and every sum rounded to fp32: convolution_edge_s(), the | ||
| * interior loops of convolution_y_c_s() / convolution_x_c_s() and the AVX2 | ||
| * scanlines alike. */ | ||
| VMAF_MTL_FUNC float vmaf_mtl_fm_blur(VmafMtlFmTaps t, VmafMtlFmWindow win) |
| int exponent = 0; | ||
| const double fraction = frexp(limit, &exponent); | ||
| out.value = vmaf_mtl_soft_make((vmaf_mtl_u64)ldexp(fraction, 53), exponent - 53); | ||
| const bool is_integer = limit == floor(limit) && limit < 2048.0; |
| for (unsigned wrong = 0; wrong <= WRONG_PLAN_KINDS; wrong++) { | ||
| const uint64_t sum = model_sum(&pic, wrong, m); | ||
| exact = wrong == 0u ? model_exact(m) : exact; | ||
| if (moment_of(sum, &pic) != cpu) { |
| moment_of(sum, &pic), cpu); | ||
| } | ||
| } | ||
| if ((moment_of(exact, &pic) != cpu) != c->rounds) { |
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 15:04
c82f3b5 to
18b19c2
Compare
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 15:14
18b19c2 to
f735c0d
Compare
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 15:44
f735c0d to
7cd06b1
Compare
5 of 6 tasks
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…t results by bits (ADR-1502) (#1927) * fix(codeql): fix the open CodeQL alerts in code and compare exact test results by bits (ADR-1502) Fixes 67 of the 71 open CodeQL alerts without changing a library or tool object file (GCC and clang, x86-64 and aarch64, -Db_lto=false, cmp of every object). No score moves. The exact-twin, replay and recorded-value tests asserted bit identity with `==`, which calls +0 equal to -0. They now use core/test/float_bits.h: vmaf_test_identical_f64/_f32 hold when the bit patterns match and the value is not a NaN, so the helper is never weaker than the `==` it replaces. test_float_bits pins it and fails for `a == b` and for a bit compare without the NaN rule. Other fixes: the three SpEED base-entropy products convert the float product's result explicitly (ADR-1477's arithmetic unchanged); adm.h and motion.h get include guards and adm_tools.h / adm_csf_tools.h open theirs before the includes; adm_options.h names the dead ADM_OPT_DEBUG_DUMP switch in prose; run_netflix_value_tests() keeps one body in both meson variants of test_speed_upstream_form; the CLI test fake keeps a comparison result instead of a caller's pointer. Left open: 1338 and 1339 (white-box tests that include ciede.c and integer_psnr.c, dismissed before as 88 and 96) and 1392 and 1393 (test_float_moment_sum.c, which open PR #1921 moves). * test(float_bits): split the signed-zero and neighbour cases to stay under the branch budget clang-tidy readability-function-size counted 16 branches (threshold 15) in test_signed_zeros_and_neighbours_differ() in all five lanes. The cases are unchanged; they now run as two tests. * test(moment): compare the SIMD moments with the scalar by bits The check_frame() comparison #1923 added asserts "a SIMD moment is not the scalar's bits" with `!=`, which CodeQL reports (cpp/equality-on-floats) and which accepts +0 for -0. It now uses vmaf_test_identical_f64() (ADR-1502). AVX2, AVX-512, NEON and SVE2 (qemu) runs pass.
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 18:19
751068c to
3dc3d70
Compare
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 20:20
3dc3d70 to
eb2af7e
Compare
lusoris
added a commit
that referenced
this pull request
Oct 3, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 3, 2026 23:02
eb2af7e to
f6a6c57
Compare
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 00:34
f6a6c57 to
b37ae99
Compare
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 02:57
b37ae99 to
8aed4c9
Compare
This was referenced Oct 4, 2026
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 14:55
8aed4c9 to
716d359
Compare
10 of 18 tasks
19 of 21 tasks
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 17:27
716d359 to
cc3fda6
Compare
lusoris
added a commit
that referenced
this pull request
Oct 4, 2026
…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.
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 19:50
cc3fda6 to
b19b5eb
Compare
… ship the FUSE mount tools in the node image (ADR-1593) (#2041) * feat(helm)!: grant node pods FUSE and the eBPF tracker per value, and ship the FUSE mount tools in the node image (ADR-1593) The published node image had no FUSE helper, so a node in mount mode refused to start, and the chart could neither give a node pod /dev/fuse nor turn on the eBPF descriptor tracker, while still rendering storage.mode: mount. - docker/Dockerfile.node: a fuse-tools stage copies Debian's setuid fusermount3 and the util-linux mount and umount it runs (libfuse 3.17 calls /bin/mount whenever /etc/mtab exists, which Docker creates), with their libraries at their Debian paths, into the runtime base. The record lands in /usr/local/share/vmafx/fuse-tools/ and licensing.json lists it as the fuse-tools dpkg-copied component; record-copied-debian-libs.sh accepts a per-line destination directory. - Helm: node.fuse (device-plugin resource, bounding set SYS_ADMIN and DAC_READ_SEARCH, allowPrivilegeEscalation for the setuid helper, UID 65532 kept, optional AppArmor profile) and node.ebpf (VMAFX_EBPF_BYPASS, mount prefix, UID 0 with BPF, PERFMON and SYS_ADMIN, host tracefs read-only). templates/node-validate.yaml refuses mount mode without node.fuse and every tracker combination the node would stop on. - Tests: test_helm_node_fuse_ebpf.py, test-record-copied-debian-libs.sh (rule-enforcement), the runtime-contract test requires the FUSE tools and a mount in the publish smoke test, which now mounts an rclone remote as UID 65532. Migration: storage.mode: mount now needs node.fuse.enabled: true and node.fuse.resourceName (the extended resource of a FUSE device plugin, for example devic.es/fuse); without them the chart refuses to render. Use storage.mode: http-serve or auto on clusters without FUSE.
…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…
lusoris
force-pushed
the
fix/metal-twins-exact
branch
from
October 4, 2026 20:05
b19b5eb to
891b821
Compare
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
2 participants
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
Every open Metal row now has its port: each Metal twin runs the design that makes its CUDA, HIP or SYCL twin return the CPU's scores bit for bit (ADR-1498). No row is closed. They close when the second outside-tester report shows each row's cases passing on an Apple GPU. #1918 (merged) built that measurement.
.metalfile compiles with-fno-fast-math -ffp-contract=off, defined once incore/src/metal/meson.build. Under these flags Metal's fp32+ - * /,sqrtandfmaare correctly rounded (Metal Shading Language Specification 4.1, Table 8.1).core/src/feature/metal/metal_portable.h. The header compiles both as MSL and as host C/C++, and a host test holds it against the CPU extractor value by value.metal_soft_double.h/metal_soft_signed.h. These are the SYCL twins' fp64 operations in 64-bit integers, statement for statement.c82f3b577the hostedmacOS MetalandmacOS Clang+Metaljobs compiled all 18.metalfiles (-std=metal3.1 -mmacosx-version-min=14.0 -fno-fast-math -ffp-contract=off) and every.mmwithout error. The tester's run is the first execution.float_momentlane order. This head passestest_metal_selftest_float_moment,test_metal_integer_vif_gainandtest_metal_selftest_integer_psnrunder qemu-aarch64 with a clang cross build.Type
fix— bug fixsycl/cuda/simd— backend-specific (Metal)Checklist
make format && make lintis green locally. The pre-commit hooks ran on every commit (clang-format, black, ruff, markdownlint, semgrep, source ADR citations, praetor audit), and the clang-tidy cpu lane was measured on every new C/C++ test in the dev container (scripts/dev/tidy-lane.sh cpu --only ...).python3 scripts/ci/run_meson_test.py -- -C build-cpu --suite fastgives 323 OK and 0 fail, including everymetal-selftestcase.tools/rc1-testerpytest passes, andscripts/dev/preflight.sh --stage msvcismpasses./cross-backend-diffand the worst ULP is ≤ 2. — Not possible here: the twins are Metal and this host has no Apple device. The parity tests compare at==(ciede at1e-9) on the tester's device, and the bundle runs the gate with--hold-exact metal..c/.cpp/.cu/.h/.hpp, it has the appropriate license header (seeCONTRIBUTING.md).!orBREAKING CHANGE:and the migration path is documented below. — Not breaking. One option went away:integer_motion_metaldropsmotion_add_uv, which the CPUmotionnever had and which the twin accepted only asfalse.docs/adr/_index_fragments/<NNNN-slug>.mdand the slug is appended todocs/adr/_index_fragments/_order.txt.Bug-status hygiene (ADR-0165)
docs/state.mdis updated. Every ported Metal row records its port, its host proof and the parity test whose passing report closes it, and every row stays open.T-METAL-INTEGER-VIF-FP32-GAIN-2026-10-03, found and ported here, andT-GPU-ADM-DECOUPLE-FP32-RECIPROCAL-2026-10-03. The second covers CUDA and HIP; it was found while porting and is not measured on a device.T-METAL-SYCL-FP64-FREE-ARITHMETIC-COPIES-2026-10-03tracks the Metal copies of the SYCL fp64-free headers.Netflix golden-data gate (ADR-0024)
assertAlmostEqual(...)score in the Netflix golden Python tests. No CPU file changes its arithmetic. The shared headersff_pair.h,ff_math.h,ciede_ff_math.handadm_gain_limit.hgained Metal branches, and their SYCL, HIP and CUDA preprocessed output is byte-identical (-E -P).Cross-backend numerical results
There is no device result yet. The host proofs below hold each header against the CPU's own code. The device result is the tester's report: per-row verdicts in
metal_rows, gate cells inmetal_gate.Deep-dive deliverables (ADR-0108)
docs/research/1498-metal-shading-language-fp-semantics.md.## Alternatives considered.AGENTS.mdinvariant note —core/src/feature/metal/AGENTS.md, section "Exact designs of the Metal twins (ADR-1498)". This PR also corrects the stale motion and float_adm notes there.changelog.d/changed/metal-twins-exact.md(strict FP) andchangelog.d/fixed/metal-twins-exact.md(what a Metal user sees change).docs/rebase-notes.md, "The Metal twins run the exact designs of the other GPU backends".Reproducer
On any host:
meson setup build core && ninja -C build -j4 python3 scripts/ci/run_meson_test.py -- -C build --suite fast python3 scripts/ci/run_meson_test.py -- -C build --suite metal-selftestOn an Apple Silicon Mac, either build with
-Denable_metal=enabledand runmeson test -C build --suite metal, or run the macOS tester bundle (docs/usage/tester-image.md).Per row: port, host proof, report case
==unless noted)T-METAL-FLOAT-PSNR-FP32-BLOCK-SUMS-2026-10-02test_metal_float_psnr_mathtest_metal_float_psnr_parity(10/12/16-bit large diffs, past 2^53)T-GPU-FLOAT-MOMENT-16BIT-SQUARES-2026-10-02test_metal_float_moment_math,test_metal_float_moment_sumtest_metal_float_moment_parityincl.test_float_moment_16bit_past_2_53_exactT-METAL-ADM-GAIN-LIMIT-FLOAT32-2026-10-01adm_gain_limit_product(), integer reciprocaltest_metal_integer_adm_mathtest_metal_integer_adm_parityT-METAL-MOTION-BLUR-THEN-DIFF-2026-09-29,T-GPU-MOTION-SAD-SCORE-NOT-EMITTED-2026-10-02test_metal_integer_motion_math(96/96 equal), table contracttest_metal_integer_motion_parity(motion3 at==), gatemotioncellT-GPU-TWIN-PARITY-GAPS-OUTSIDE-CUDA-2026-09-30test_metal_float_ssim_mathtest_metal_integer_psnr_parity,test_metal_motion_v2_parity,test_metal_float_ssim_parityT-GPU-INTEGER-VIF-MIN-DIM-TWINS-2026-09-29test_metal_integer_vif_mathtest_metal_integer_vif_parity(model_boundary)T-METAL-INTEGER-VIF-FP32-GAIN-2026-10-03test_metal_integer_vif_gain(400k variances, 6 limits)test_metal_integer_vif_parityT-BUG048-GPU-OPTION-PARITY-REMAINDER-2026-09-26test_metal_twin_option_tables_contracttest_metal_twin_option_parityT-GPU-FLOAT-MOTION-CPU-FLOAT-SUM-2026-10-01,T-GPU-FLOAT-MOTION3-MISSING-2026-09-30test_metal_float_motion_mathtest_metal_float_motion_parityT-GPU-CIEDE-CPU-ARITHMETIC-2026-10-01ciede_ff_math.hon Metal primitives, raster sumtest_metal_ciede_mathtest_metal_integer_ciede_parity(1e-9)T-GPU-FLOAT-ADM-TINY-FRAME-FLOOR-2026-10-01,-FRAME-SUM-FLOOR-,-CPU-ARITHMETIC-test_metal_float_adm_math(10^6 samples x 5 limits)test_metal_float_adm_parityT-GPU-FLOAT-VIF-CPU-ARITHMETIC-2026-10-01test_metal_float_vif_math(2.4M operands)test_metal_float_vif_parityT-GPU-SSIM-FRAME-SUM-ORDER-2026-10-01test_metal_integer_ssim_mathtest_metal_integer_ssim_parityT-GPU-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01test_metal_float_ms_ssim_math(15 means, 8 frames)test_metal_float_ms_ssim_paritytools/rc1-tester/image/metal-rows.jsonmaps each row to these cases and to its fixture metrics and gate cells.Known follow-ups
Hosted
Ubuntu SYCL/Ubuntu SYCL+CUDAfailtest_icx_system_libm(vmaf --versionexits -11 after the loader relinks Intel'slibimf.sofor IFUNCcosf). This is not from this PR: the same failure is on the runs of fix(gpu): add float_psnr's row sums in the CPU's order so the CUDA, SYCL and HIP twins are bit-identical past 2^53 units #1922, fix(sycl): keep float_adm's term kernel out of scratch memory on Xe2 and order its probe's queue #1924 and fix(test): make the Metal self-tests build on MSVC, pass on macOS and fit the sanitizer job #1925, and master's last impacted SYCL leg did not run.The sanitizer job: this PR's integer ssim header and two host tests had UB under UBSan (fixed in
fix(metal): keep the integer ssim header ...), andtest_metal_float_moment_sumnow has the 600 s timeout oftest_float_moment_sum.core/src/metal/dispatch_strategy.clisted the pre-portmotion_yname and none of the motion SAD / motion / motion3 features, so model dispatch sent them to the CPU. The table now matches every Metal extractor's features, andtest_metal_twin_option_tables_contract.pykeeps it that way.docs/backends/metal/index.md: finished the truncated dispatch sentence, addedcoreto the configure commands, and replaced the "remaining kernels to come" roadmap with the 17 registered extractors and ADR-1498.test_metal_twin_option_tables_contractcaches each table parse (4.9 s to 1.6 s) and has a 120 s timeout. The hosted Sanitizers (thread) and (undefined) jobs had killed it at 30 s.T-INTEGER-VIF-SV-SQ-CONVERSION-UB-2026-10-03is closed by fix(vif): convert integer VIF's residual variance through one defined helper that keeps x86's bits #2040 (ADR-1561,vif_sv_sq()).test_metal_integer_vif_gainnow calls the helper, covers the whole range again, and passes on x86 and on aarch64 (qemu, clang).The
.mmoption tables (and some registrations) put designators out of declaration order (-Wreorder-init-list,-Wc99-designator). This is pre-existing on master; the ports add instances infloat_adm_metal.mm,integer_motion_metal.mmandinteger_psnr_metal.mm. A reorder is mechanical but touches the planted texts of five contracts, so it is left for a follow-up.T-GPU-ADM-DECOUPLE-FP32-RECIPROCAL-2026-10-03: the CUDA and HIP scale-0 decouple takes an fp32 reciprocal. This was read from source and not measured; it needs its own fix PR with a device run.integer_cambi_metaldoes not declareheatmaps_path, because the CPU's writers are static incambi.c. A model with that option keeps the CPU extractor.MSL risks that only the hosted compile and the device run settle:
constantbuffers;as_typeinside constant builders;size_t/int64_tin kernels;RC5 (
T-METAL-SYCL-FP64-FREE-ARITHMETIC-COPIES-2026-10-03): the Metal copies of the SYCL fp64-free headers become one implementation.