fix(sycl): make float_ms_ssim_sycl the CPU's arithmetic so its per-scale means are bit-identical

[lusoris]

Summary

float_ms_ssim_sycl now computes the CPU extractor's arithmetic. On an Arc A380 every per-scale mean of every measured frame is the CPU's bit for bit, with enable_lcs and enable_chroma. Before, the twin matched the CPU on none of 104 frames (up to 2.98e-6). It costs time: 42.6 ms per 3840x2160 frame against 31.4.

This is the SYCL part of T-GPU-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01, which #1695 opened when it fixed the CUDA twin. The SYCL twin had the same four differences: the decimate added sample * tap in two roundings where the reference fuses each tap; the Gaussian window sums were fp32 running sums where iqa_convolve() adds fp32 products in double; l / c / s were fp32 quotients where the CPU divides double numerators by fp32 denominators; the host combined unrounded means without fabs() on l and c.

A SYCL kernel may not use double (ADR-0220), so the CUDA fix does not carry over as written. float_ssim_sycl already had the same arithmetic without fp64, and this PR shares it instead of copying it.

What changed

• core/src/feature/sycl/sycl_ssim_terms.h (new): the per-pixel SSIM arithmetic, moved out of integer_ssim_sycl.cpp unchanged: window sums and the l / c quotients as exact fp32 pairs, each term to int64 units of 2^-52, the exact host sum. Both SSIM twins include it.
• core/src/feature/sycl/integer_ms_ssim_sycl.cpp: sycl::fma() per decimate tap, pair window sums, ssim_terms() for l / c / s, int64 group partials, fp32 means and fabs() of all three on the host.
• scripts/ci/cross_backend_calibration.py: EXACT_TWINS lists float_ms_ssim and float_ms_ssim_lcs: sycl.
• ADR-1414 records the design and the alternatives.

Type

• feat — new feature
• fix — bug fix
• perf — performance improvement
• refactor — no behavior change
• docs — documentation only
• test — test-only
• build / ci — tooling / infra
• port — cherry-pick from upstream Netflix/vmaf
• sycl / cuda / simd — backend-specific

Checklist

• Commits follow Conventional Commits (the commit-msg hook enforces this).
• make format && make lint is green locally. (clang-format on the touched files; the SYCL clang-tidy wrapper on integer_ms_ssim_sycl.cpp, integer_ssim_sycl.cpp, the new header and the test: 0 warnings in them, nine older ones in the test file removed; the commit hooks.)
• Unit tests pass: python3 scripts/ci/run_meson_test.py -- -C build. (--suite=fast on a GCC CPU build: 208 of 208. On the icx SYCL build: --suite sycl on the Arc A380 58 of 60, and 207 of 208 without a device. The three failures fail on master too: test_sycl_float_adm_parity and its large variant because float_adm_sycl still uses scratch memory on this GPU (ADR-1395), and test_integer_adm_simd on icx builds since fix(adm): keep the scale-0 masking centre tap in int32 (ADR-1402) #1700, fixed by fix(build): build every x86 SIMD library without FP contraction #1706.)
• If I touched any SIMD/GPU code path, I ran /cross-backend-diff and the worst ULP is ≤ 2. (0 on the device; tables below.)
• If I touched a feature extractor with SIMD/GPU twins, I either updated every twin or listed the gap under "Known follow-ups" below. (CUDA is done by fix(cuda): build every kernel without FMA contraction; make float_ms_ssim_cuda the CPU's arithmetic #1695; HIP and Metal are listed there.)
• If I added a new .c / .cpp / .cu / .h / .hpp, it has the appropriate license header (see CONTRIBUTING.md).
• If this is a breaking change, the commit message uses ! or BREAKING CHANGE: and the migration path is documented below. (Not a breaking change.)
• If this PR adds an ADR, the ADR row lives in docs/adr/_index_fragments/<NNNN-slug>.md.

Bug-status hygiene (ADR-0165)

• docs/state.md updated in this PR: T-SYCL-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01 added under Recently closed with the evidence; T-GPU-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01 now names Metal only and stays open (HIP landed in fix(hip): make float_ms_ssim_hip the CPU's arithmetic so the twin is bit-identical #1710).

Netflix golden-data gate (ADR-0024)

• I did not modify any assertAlmostEqual(...) score in the Netflix golden Python tests.
• If I believe a golden value must change, I have explained why below AND pinged @lusoris for a CODEOWNERS exception. (None changes: no CPU code is touched.)

Cross-backend numerical results

Arc A380 (dg2-g11, xe driver, Level Zero, icx 2026.0), --precision max, float_ms_ssim on --backend cpu of a GCC build against float_ms_ssim_sycl. Frames identical and largest absolute difference of the score. "Before" is master dff31c445.

fixture                               frames  before            after
Netflix 576x324                       48      0/48, 6.9e-8      48/48, 0
checkerboard 1 px, 1920x1080          3       0/3,  1.06e-6     3/3,   0
checkerboard 10 px, 1920x1080         3       0/3,  2.98e-6     3/3,   0
BBB 3840x2160                         200     0/50, 1.23e-6     199/200, 1.1e-16
enable_lcs, 16 outputs, four fixtures 104     -                 all 1664 outputs identical
enable_chroma, checkerboards + BBB    56      up to 6.8e-7      all identical (Y, Cb, Cr)
Netflix 10-, 12-, 16-bit, 4:2:2 10-bit 3 each  -                all identical
enable_db + clip_db, Netflix          48      -                 45/48, 3.6e-15

The device arithmetic matches on every frame: all per-scale means are identical. The three places that are not zero are host math calls. The twin's host tail calls the pow() and log10() of its own icx build, which links Intel's libimf, and the reference here is a GCC build (T-ICX-LIBIMF-HOST-MATH-2026-10-01 in #1706).

Against the CPU extractor of the same icx binary, which is what the gate compares, scripts/ci/cross_backend_parity_gate.py --features float_ms_ssim float_ms_ssim_lcs --backends cpu sycl:

fixture                         binary as built here         with #1706's flags
Netflix 576x324, 48 frames      FAIL 7.7e-9 (lcs 5.96e-8)    OK 0, OK 0
checkerboard 1 px, 3 frames     OK 0                         OK 0, OK 0
checkerboard 10 px, 3 frames    FAIL 7.9e-9 (lcs 5.96e-8)    OK 0, OK 0
BBB 3840x2160, 200 frames       FAIL 1.4e-8 (lcs 5.96e-8)    OK 0, OK 0

The failures are the CPU side, not the twin: on this AVX-512 host the icx build contracted the scalar tails of ssim_avx512.c, so the CPU extractor itself was one fp32 unit off in a per-scale mean on 4 of 104 frames. #1706 fixes that; with its two flag changes applied to this tree the cell is 0 on all four fixtures. The twin equals the scalar CPU path of its own binary on every frame either way.

Bit-identical in practice, not by construction: the CPU forms l and c in double and adds them into a running double in raster order; the twin carries them as pairs good to about 2^-46 and adds them exactly. The fp32 rounding of each per-scale mean absorbs the difference.

test_sycl_kernel_scratch on the A380 audits 110 kernels; none of this PR's uses scratch memory and the ratchet list is unchanged.

Performance (if perf or feat)

Not a performance change, and it costs time. vmaf tool on the Arc A380, alternating before/after runs, medians, host load 10 to 20:

twin, fixture                              before   after    pairs
float_ms_ssim_sycl, BBB 3840x2160          31.4     42.6     15 x 100 frames
float_ms_ssim_sycl, Netflix 576x324        0.84     1.20     15 x 46 frames
float_ssim_sycl scale=1, BBB 3840x2160     22.9     23.0     9 x 100 frames   helpers only moved
float_ms_ssim (CPU), BBB 3840x2160         117      -        3 x 20 frames

ms per frame. A cheaper tap accumulation (one two-sum per tap instead of a full pair addition) gave the same values and 42.2 ms, so the taps are not where the time goes, and the shared arithmetic was left as float_ssim_sycl had it. Both passes read every tap from device memory (22 and 55 samples per pixel), as before.

Deep-dive deliverables (ADR-0108)

• Research digest — no digest needed: the four differences were analysed in Research-1403 for the CUDA twin, and the measurements of this fix are in ADR-1414 and the docs/state.md row.
• Decision matrix — ADR-1414 ## Alternatives considered.
• AGENTS.md invariant note — core/src/feature/sycl/AGENTS.md (the four rules, the shared header), scripts/ci/AGENTS.md (EXACT_TWINS), and docs/development/rebase-sensitive-invariants.md.
• Reproducer / smoke-test command — pasted below under "Reproducer".
• CHANGELOG fragment — changelog.d/fixed/sycl-float-ms-ssim-cpu-arithmetic.md.
• Rebase note — docs/rebase-notes.md, ADR-1414.

Reproducer

source /opt/intel/oneapi/setvars.sh
CC=icx CXX=icpx meson setup build-sycl core -Denable_sycl=true -Denable_cuda=false --buildtype=release
ninja -C build-sycl
# 53 and 54 of 54 outputs differ on master; need a SYCL GPU.
build-sycl/test/test_sycl_ms_ssim_parity
build-sycl/test/test_sycl_ms_ssim_parity_large
# float_ssim_sycl shares the arithmetic and must not move.
build-sycl/test/test_sycl_float_ssim_parity
# No kernel uses scratch memory (Intel GPU).
build-sycl/test/test_sycl_kernel_scratch
# No device needed.
python3 -m pytest core/test/test_sycl_kernel_source_contract.py scripts/ci/test_cross_backend_parity_gate.py
# The gate cell, tolerance 0 (needs #1706 on an AVX-512 host).
ONEAPI_DEVICE_SELECTOR=level_zero:0 python3 scripts/ci/cross_backend_parity_gate.py \
  --vmaf-binary build-sycl/tools/vmaf \
  --reference python/test/resource/yuv/checkerboard_1920_1080_10_3_0_0.yuv \
  --distorted python/test/resource/yuv/checkerboard_1920_1080_10_3_1_0.yuv \
  --width 1920 --height 1080 --features float_ms_ssim float_ms_ssim_lcs --backends cpu sycl

Known follow-ups

• float_ms_ssim_metal keeps the old arithmetic: T-GPU-FLOAT-MS-SSIM-CPU-ARITHMETIC-2026-10-01 (integer_ms_ssim_hip landed in fix(hip): make float_ms_ssim_hip the CPU's arithmetic so the twin is bit-identical #1710).
• Time: the 11 ms at 3840x2160 is the pair arithmetic on a kernel that reads every tap from device memory. Staging the window through local memory, as float_ssim_sycl does in its horizontal pass, is tuning left for after the exactness work.
• fix(build): build every x86 SIMD library without FP contraction #1706 makes the same-binary gate cell pass on AVX-512 hosts. This PR does not depend on it to build or to pass its own tests (the exact test compares against the scalar CPU path).
• float_ms_ssim_cuda is bit-identical too (ADR-1403) and not listed in EXACT_TWINS; this PR does not touch that entry.