fix(sycl): give every SYCL feature kernel the CPU's fp32 arithmetic

[lusoris]

Summary

Every SYCL feature kernel now does fp32 arithmetic the way the CPU reference does. Each SYCL feature TU compiles with -fp-model=precise -ffp-contract=off -foffload-fp32-prec-div -foffload-fp32-prec-sqrt (ADR-1367). So a * b + c is never fused into an FMA, and / and sqrt are correctly rounded, on the AOT images and on the SPIR-V image other devices compile at first launch.

Before this, -fp-model=precise alone left 12%, 29% and 8% of random multiply-adds, divisions and square roots different from the host, while the SYCL guide called the kernels IEEE-754 strict. This closes T-SYCL-FP-MODEL-PRECISE-CONTRACTS-2026-09-29.

The branch is one commit on master 10f27efe2; #1627, which it was stacked on, is merged. On the Arc A380 this change fails the same tests as master, and its own contract test passes on the JIT image and on a dg2-g11 AOT image (A380 section below).

What changes:

• One flag set, one variable. sycl_strict_fp_args is defined once in core/src/meson.build, between new policy markers, and reaches every feature TU. perf(sycl): run ssimulacra2 on the device and wait once per frame in ms_ssim #1627's sycl_exact_fp_args / sycl_exact_fp_sources fold into it; sycl_exact_fp.h is unchanged apart from its comments.
• Every link that generates device images carries the precision pair.
  • ADR-1358 found that the pair only worked at the link. That was before ADR-1360's -fno-sycl-rdc. The AOT images are now finished per TU at compile time and need the pair there, while the SPIR-V JIT image is still device-linked at the final icpx link and needs it there.
  • A JIT-only build (the SYCL parity CI lane) without the link pair keeps approximate division.
  • The MSVC build's explicit device link (ADR-1364) generates every Windows image and takes sycl_strict_fp_args whole.
  • The contract test's probe gets its own device link under MSVC, because link.exe never wraps device code.
• integer_vif_sycl writes its fused operation explicitly. It approximated the CPU's fp64 sigma2_sq - g * sigma12 through an implicit contraction; with contraction off it drifted 6x further from the CPU. It is now sycl::fma(-g, sigma12, sigma2_sq) and ends up closer than before.
• Tests.
  • test_sycl_fp_arith_contract (new) compiles a probe kernel with the extractors' own arguments and checks 1,048,576 random and 2,197 boundary operands on the device.
  • test_strict_fp_compiler_args.py executes the policy for icpx and AdaptiveCpp.
  • test_sycl_kernel_source_contract.py now checks that no TU gets a private FP list.
• clang-tidy. gen-sycl-compile-commands.py drops the icpx-only precision flags, which stock clang rejects as unknown arguments. Without this, Tidy SYCL would fail on every SYCL TU.
• Docs. The SYCL guide states what the line guarantees and what still differs: transcendentals, summation order, fp32 for CPU fp64 expressions, and float_ssim's formula. These are corrected to match:
  • docs/development/sycl-toolchains.md, docs/metrics/features.md and the AGENTS files;
  • core/src/feature/cuda/AGENTS.md, which overstated which CUDA kernels use --fmad=false.

Type

• fix — bug fix
• sycl / cuda / simd — backend-specific

Checklist

• Commits follow Conventional Commits (the commit-msg hook enforces this).
• make format && make lint is green locally. Partial, run on the rebased head on Linux:
  • pre-commit run --from-ref origin/master --to-ref HEAD: all 39 applicable hooks pass.
  • make lint-py lint-md docs-fragments-check lint-reuse python-locks-check passes.
  • make lint-sh passes every step except scripts/githooks/tests/test_install.py. That test fails only while a .venv without pre_commit is first on PATH, and it fails the same way on master. It is a hook-installer bug outside this change; the test passes on this branch without that .venv on PATH.
  • make verify-all passes with the engine standards-gate.yml pins (f41e74d8).
  • A scoped SYCL tidy ratchet over the six SYCL TUs and two test files this change touches or reaches through its headers finds no uncited NOLINT and no file above master's local measurement. Two files carry findings: speed_sycl_pipeline.cpp (11) and ssimulacra2_sycl.cpp (13), all misc-static-assert. Master measures the same counts in the same configuration, which is the drift T-TIDY-RATCHET-GPU-LANES-UNREPRODUCIBLE-2026-09-22 tracks.
  • The full make lint-c was not run.
• Unit tests pass.
  • --suite sycl 52/52 on the Arc B580 and 52/52 on the UHD 770 (office box, before the rebase).
  • test_strict_fp_compiler_args, test_sycl_kernel_source_contract and test_gen_sycl_compile_commands pass.
  • On the Arc A380, test_sycl_fp_arith_contract passes on both image paths. The A380's 16 failing tests fail identically on master; see below.
• If I touched any SIMD/GPU code path, I ran /cross-backend-diff and the worst ULP is ≤ 2. I ran the per-twin CPU comparison and the ADR-0214 gate instead (tables below). Several twins were never within 2 ULP; none moved outside its tolerance.
• If I touched a feature extractor with SIMD/GPU twins, I either updated every twin or listed the gap under "Known follow-ups" below.
• If I added a new .c / .cpp / .cu / .h / .hpp, it has the appropriate license header (see CONTRIBUTING.md).
• If this PR adds an ADR, the ADR row lives in docs/adr/_index_fragments/<NNNN-slug>.md and the slug is appended to docs/adr/_index_fragments/_order.txt — do not edit docs/adr/README.md directly.

Bug-status hygiene (ADR-0165)

• docs/state.md updated:
  • Moved to Recently closed: T-SYCL-FP-MODEL-PRECISE-CONTRACTS-2026-09-29.
  • Opened, RC3, each with a verify-and-time command for ryzen-4090-arc: T-CUDA-FP-CONTRACT-DEFAULT-2026-09-29 and T-HIP-FP-CONTRACT-DEFAULT-2026-09-29.
  • Opened, RC3: T-CLI-FLOAT-MOMENT-NO-TWIN-2026-09-29.
  • Opened, RC2: T-CI-PARITY-GATE-STALE-METRIC-KEYS-2026-09-29.

Netflix golden-data gate (ADR-0024)

• I did not modify any assertAlmostEqual(...) score in the Netflix golden Python tests.
• No golden value needs to change: CPU code is untouched.

Cross-backend numerical results

Arc B580 and UHD 770 (office box, measured before the rebase)

Max abs diff against --backend cpu --precision max, before (#1627 head) -> after. The Arc B580 and the UHD 770 give the same result unless noted. 576x324 is the Netflix pair (48 frames); 4K is BBB 3840x2160 (22 frames).

twin            576x324 max            3840x2160 max
vif             3.87e-7 -> 3.49e-7     1.22e-7 -> 1.49e-7
adm             2.27e-7 (unchanged)    1.34e-6 (unchanged)
motion          1.26e-5 (unchanged)    5.56e-6 (unchanged)
motion_v2       0 (unchanged)          0 (unchanged)
float_psnr      0 (unchanged)          0 (unchanged)
float_motion    3.05e-6 -> 3.09e-6     2.67e-5 -> 2.29e-5
float_vif       2.71e-5 -> 3.81e-5     3.18e-6 -> 3.14e-6
psnr            0 (unchanged)          0 (unchanged)
float_moment    0 (unchanged)          0 (unchanged)
ciede           1.18e-5 -> 1.14e-5     9.71e-5 -> 4.53e-5
float_ssim      3.10e-7 -> 1.47e-7     8.31e-5 -> 8.80e-5 (scale=1 forced; formula difference, pre-existing)
ssim            1.40e-8 -> 7.40e-9     7.76e-8 -> 5.59e-8
float_ms_ssim   6.96e-8 -> 6.92e-8     2.84e-7 -> 4.69e-7 (mean 1.36e-7 -> 6.75e-8)
psnr_hvs        8.37e-5 (moves <1.3e-6) 8.42e-4 (moves <1.3e-6)
ssimulacra2     1.12e-12 (unchanged)   6.65e-12 (unchanged)
float_adm       2.50e-5 -> 2.53e-6     1.12e-6 -> 1.97e-7
cambi           0 (unchanged)          2.22e-15 (unchanged)
speed_chroma    0 (unchanged)          0 (unchanged)
speed_temporal  0 (unchanged)          0 (unchanged)
default model   1.26e-5 (unchanged)    5.56e-6 (unchanged)

scripts/ci/cross_backend_parity_gate.py --backends cpu sycl on the Netflix pair, one feature per run: the 15 runnable cells pass on both GPUs before and after. The cambi and motion cells abort on stale metric names in the gate itself; that is T-CI-PARITY-GATE-STALE-METRIC-KEYS-2026-09-29, and both twins are covered by the table above. Where the maximum grew, the residual comes from operations the flags do not control, such as sycl::log2, the summation order, or fp64 terms on the CPU; see Research-1367 §6.

Arc A380 (dg2-g11, zeus, 2026-09-30)

Setup. Kernel 7.2.8 with the A380 on the xe driver (xe.force_probe=56a5), NEO 26.35.39758.10, IGC 2.41.5, icpx 2026.0. Builds were --buildtype=release -Db_lto=false, run with ONEAPI_DEVICE_SELECTOR=level_zero:0.

meson test --suite sycl:

build                                  pass / fail   test_sycl_fp_arith_contract
master 10f27efe2, JIT                  39 / 16       (not on master)
this branch, JIT                       40 / 16       pass
this branch, AOT dg2-g11               40 / 16       pass

The 16 failures are the same tests in all three builds. They cover:

• adm (_parity, _parity_large, _tiny_frames)
• float_adm (_parity, _parity_large)
• float_vif (_parity, _parity_large)
• psnr_hvs (_parity, _parity_large, _parity_simd32)
• speed_chroma (_parity, _parity_large)
• speed_singular_parity, speed_temporal_parity_large, shared_planes and motion_tiny_frames

This change does not cause them. The xe driver returns wrong values from GPU kernels that use scratch memory:

• Two standalone SYCL kernels with no vmafx code fail on this device:

  • one indexes a private array: 16384/16384 work-items wrong;
  • one spills registers: 261888/262144 work-items wrong.

  They fail as JIT and as AOT, on Level Zero and on OpenCL, and both are correct on the OpenCL CPU device.

• IGC's zeinfo shows that each of the 16 failing tests runs a kernel with a non-zero private_size or spill_size.

• On 2026-09-25, under i915 with the same NEO and IGC, 13 of these 16 tests passed on this box. The other three (psnr_hvs_parity_simd32, shared_planes, motion_tiny_frames) did not exist yet.

Removing the scratch use from the SYCL kernels is separate work, branched from master.

On the scratch-free twins, per-frame CPU parity via scripts/dev/speed_gpu_parity.py --backend sycl looks like the B580's (max abs diff, master -> this branch):

twin          576x324 (48 frames)     3840x2160 (50 frames)
vif scale0    4.99e-8 -> 4.46e-8      7.02e-8 -> 6.99e-8
vif scale3    3.87e-7 -> 3.49e-7      1.47e-7 -> 1.73e-7
motion2/3     0 -> 0 (48/48 exact)    0 -> 0 (50/50 exact)
float_motion  3.05e-6 -> 3.09e-6      2.37e-5 -> 2.36e-5   (its `motion` output)

AOT image check. A dg2-g11-only AOT build stops at sycl_aot_image_check, on master as on this branch. With a single target, ocloc stores each TU's image as a plain zebin ELF (31/31 images here), and check_aot_image.py only accepts ar fat binaries. The AOT row above was built with ninja -k 0. A fix is in progress on a separate branch from master, fix/sycl-aot-check-single-target.

Performance (if perf or feat)

Arc B580, 3840x2160, ms per frame, before -> after. Frames were held in memory (8 frames loaded once, 50 fed through vmaf_read_pictures()), median of 5 interleaved runs, or 7 for the last three:

vif 6.28->6.42  adm 6.42->6.25  motion 6.11->5.92  motion_v2 6.64->6.70  float_psnr 7.51->7.73
float_motion 7.33->7.31  float_vif 8.19->7.98  psnr 6.90->7.01  ciede 8.56->8.55
float_ssim(scale=1) 12.07->11.70  ssim 7.47->6.94  float_ms_ssim 15.46->15.54  ssimulacra2 34.01->34.05
float_adm 9.61->8.83  cambi 8.17->7.97  speed_chroma 7.13->6.59
psnr_hvs 20.82->21.47  float_moment 5.71->5.61  speed_temporal 6.91->7.00

Every twin is within -8% to +3.1%, inside the spread of kernels whose code did not change, so no twin crossed the 10% bar. The prescribed CLI method, (t(22) - t(2)) / 20 with a median of 3, was run too. It is in the research digest, but on the shared office host it swung ±50% for unchanged twins, so it cannot resolve 10%. The A380 was not timed: under xe, half its twins compute wrong values.

Deep-dive deliverables (ADR-0108)

• Research digest — docs/research/1367-sycl-strict-fp-every-feature-tu.md
• Decision matrix — ADR-1367 ## Alternatives considered
• AGENTS.md invariant note — core/src/sycl/AGENTS.md, core/src/feature/sycl/AGENTS.md, docs/development/rebase-sensitive-invariants.md
• Reproducer / smoke-test command — below
• CHANGELOG fragment — changelog.d/fixed/sycl-strict-fp-every-feature-tu.md
• Rebase note — docs/rebase-notes.md, fix/sycl-fp-contract-all-tus entry

Reproducer

# oneAPI sourced; Intel GPUs via Level Zero. Use a multi-target list, or
# -Dsycl_icpx_aot_targets= for JIT only, until the single-target AOT check fix lands.
CC=icx CXX=icpx meson setup build core -Denable_sycl=true -Denable_cuda=false \
  -Dsycl_icpx_aot_targets=bmg-g21,adl-s --buildtype=release -Db_lto=false && ninja -C build
python3 scripts/ci/run_meson_test.py -- -C build \
  test_sycl_fp_arith_contract test_strict_fp_compiler_args test_sycl_kernel_source_contract
# Per-twin parity against the CPU extractor:
ONEAPI_DEVICE_SELECTOR=level_zero:0 python3 scripts/dev/speed_gpu_parity.py --backend sycl \
  --vmaf build/tools/vmaf --feature float_adm --feature vif --max-abs-diff 5e-5
# Arc A380, JIT build as measured above:
ONEAPI_DEVICE_SELECTOR=level_zero:0 python3 scripts/ci/run_meson_test.py -- -C build --suite sycl

Known follow-ups

• CUDA and HIP still contract in every kernel except two each (T-CUDA-FP-CONTRACT-DEFAULT-2026-09-29, T-HIP-FP-CONTRACT-DEFAULT-2026-09-29). Their division and sqrt are already IEEE. Neither is measured here.
• The Arc A380 under i915 is not measured. Under xe its scratch defect hides every kernel that uses scratch memory; see the A380 section.
• No Windows GPU run has checked the MSVC device link's images, which take the whole line (ADR-1364).
• div_rn / sqrt_rn duplicate what the line now gives on icpx. Whether plain / is as fast in the SpEED eigenvalue kernel is unmeasured.
• --feature float_moment never selects a GPU twin (T-CLI-FLOAT-MOMENT-NO-TWIN-2026-09-29).
• -Woverriding-option now appears for 60 SYCL compiles instead of 12. It is the intended override, and the icx x86 strict libraries already print it 43 times.