Skip to content

fix(gpu): make the integer ADM twins agree with the CPU on tiny frames and full-range content - #1476

Closed
lusoris wants to merge 22 commits into
port/upstream-2026-09from
fix/gpu-adm-tiny-frames
Closed

lusoris wants to merge 22 commits into
port/upstream-2026-09from
fix/gpu-adm-tiny-frames

Conversation

@lusoris

@lusoris lusoris commented Sep 18, 2026

Copy link
Copy Markdown
Contributor

Summary

The CUDA, HIP and SYCL integer ADM twins now agree with the CPU on tiny frames and on full-range content, and refuse the frames the CPU refuses. Stacked on #1473 (and parallel to #1474); review #1473 first.

Score fixes (all measured against the scalar CPU path):

Twin Problem Before After
CUDA, HIP Frames 17–32 px wide: the scale-0 cube shift is 0 there, and the host code set its rounding term to 1 << (0 - 1), which is 2^31 on x86 integer_adm_scale0 0.2 off; 32x32 = NaN CUDA ≤ 1.4e-7; HIP within the 1e-4 gate
CUDA, HIP Same widths: the scale-0 contrast-masking kernels read one column / one row past the band at the right and bottom edges 0.21 off at 17x17 (both twins) as above
CUDA, HIP, SYCL Accepted frames below 17x17, which the CPU and Metal refuse scored garbage -EINVAL + error naming the extractor
SYCL Kept the CPU's 16-bit scale-0 intermediates in 32/64 bits, so values that wrap on the CPU did not noise: 2.1e-4 off (4.4e-3 with Barten weights) 5e-8
SYCL Rounded two scale 1-3 terms with +2^31 where the CPU, CUDA and HIP use the wrapped -2^31 (Netflix#955, ADR-0155) ordinary video ≤ 7.8e-7 off ≤ 2.3e-7, every value moved toward the CPU

Verified on an RTX 4090 (CUDA), an Arc A380 (SYCL) and the gfx1036 iGPU of a Ryzen 9 9950X3D (HIP). The last one matters: HIP integer ADM now runs on real hardware here.

Latent bugs fixed on the way, all found by the lint cleanup below:

  • init_fex_hip() returned 0 after tearing its device state down when the feature-name dictionary could not be built; the next submit() ran against freed state.
  • init_fex_cuda() leaked the stream, events and modules when buffer allocation failed (close() never runs after a failed init()), and the 10/16-bit path swapped the reference and distorted strides (harmless while they are equal; inherited from upstream).
  • CUDA allocated two scratch buffers no kernel reads: about 50 MB less device memory at 1080p now.
  • Editing a header a CUDA or HIP kernel includes did not rebuild the kernel. The nvcc/hipcc custom targets had no depfile, so ninja tracked only the kernel source. This bit during this PR: removing two struct fields made every CUDA ADM test crash with an illegal address until the fatbin was deleted by hand.

Lint (ADR-1142 / ADR-0141): every file this PR touches is at 0 clang-tidy findings and 0 uncited NOLINTs in its lane. Before: 370 findings (adm_cm.cu 164, integer_adm_cuda.c 46, adm_cm.hip 88, integer_adm_hip.c 33, test_cuda_adm_parity.c 39) and 8 uncited NOLINTs (7 in integer_adm_sycl.cpp, 6 of which no longer suppressed anything, and 1 in integer_adm_cuda.c). GPU scores are byte-identical before and after every refactor commit (CUDA on 11 inputs, HIP on 12, SYCL on 9).

Tooling: the cuda, hip and sycl tidy lanes could not run at all, and could not see the kernel files. tidy-ratchet.py handed compiler flags to clang-tidy as clang-tidy options and used relative paths inside each TU's directory; .cu/.hip files never reach compile_commands.json. Fixed, with a new scripts/ci/gen-gpu-compile-commands.py, and the Tidy Ratchet job now self-tests the tooling on every PR (neither test file ran in CI before).

Test registrations: four HIP ADM tests were should_fail long after ADR-1211 fixed their cause; the flag also hid two tests reading adm3_score, which the HIP twin does not emit. They are plain tests now and pass on the gfx1036.

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.
  • Unit tests pass: meson test -C build.
  • If I touched any SIMD/GPU code path, I ran /cross-backend-diff and the worst ULP is ≤ 2: all three twins within the 1e-4 gate of scalar CPU on tiny frames and noise (CUDA ≤ 1.4e-7 tiny, SYCL 5e-8 noise, HIP 0 on noise); refactor commits byte-identical.
  • 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, HIP and SYCL all fixed; Metal already refused small frames.
  • 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 breaking. GPU extractors now refuse frames below 17x17 as the CPU always did.
  • 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: no ADR, bug fixes.

Bug-status hygiene (ADR-0165)

  • docs/state.md, all closed: T-GPU-ADM-TINY-FRAME-SHIFT-2026-09-18 (opened by port: reconcile with upstream Netflix/vmaf September 2026 (ADM/VIF SIMD fixes) #1473), T-HIP-ADM-TESTS-STALE-SHOULD-FAIL-2026-09-18, T-TIDY-GPU-LANES-UNRUNNABLE-2026-09-18, T-HIP-ADM-INIT-DICT-FAILURE-RETURNS-SUCCESS-2026-09-18, T-CUDA-ADM-INIT-LEAK-AND-STRIDE-SWAP-2026-09-18, T-GPU-KERNEL-HEADER-DEPS-UNTRACKED-2026-09-18, T-SYCL-ADM-INT16-SEMANTICS-2026-09-18.

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: no CPU score changes.

Deep-dive deliverables (ADR-0108)

  • Research digest — docs/research/2063-upstream-sync-2026-09-adm-vif-simd.md, zero-bit-shift section (GPU measurements and outcome).
  • Decision matrix — no alternatives: only-one-way fix (the CPU path is the reference and already defines the 17x17 minimum and the integer semantics).
  • AGENTS.md invariant note — core/src/feature/cuda/AGENTS.md, core/src/feature/hip/AGENTS.md, core/src/feature/sycl/AGENTS.md.
  • Reproducer / smoke-test command — below.
  • CHANGELOG fragment — changelog.d/fixed/gpu-adm-tiny-frames.md, sycl-adm-int16.md, tidy-gpu-lanes.md, gpu-kernel-header-deps.md, changelog.d/changed/cuda-adm-unused-accum-buffers.md.
  • Rebase note — docs/rebase-notes.md: upstream's CUDA ADM carries the rounding, border and stride defects; keep the fork forms; the SYCL narrowing and rounding invariants.

Reproducer

# CUDA (same shape for -Denable_hip=true -Denable_hipcc=true and -Denable_sycl=true)
meson setup build/cuda core -Denable_cuda=true -Denable_sycl=false -Db_lto=false --buildtype=release
ninja -C build/cuda
build/cuda/test/test_cuda_adm_tiny_frames      # rejection below 17x17, tiny-frame and noise parity
meson test -C build/cuda --suite fast

# the cuda tidy lane on this PR's files (a full-lane run still reports the
# stale baseline, see Known follow-ups)
python3 scripts/ci/gen-gpu-compile-commands.py build/cuda
python3 scripts/ci/tidy-ratchet.py --lane cuda --build-dir build/cuda \
  --extra-arg=--cuda-host-only --extra-arg=-nocudalib \
  --only core/src/feature/cuda/integer_adm_cuda.c --only core/src/feature/cuda/integer_adm/adm_cm.cu \
  --report /tmp/cuda-tidy.json   # warnings {} and nolint_uncited {} in the report

Verified locally

Check Result
test_{cuda,hip,sycl}_adm_tiny_frames pass on all three devices; fail on the old CUDA/HIP shift, the old border clamps (CUDA and HIP), and the old SYCL narrowing
Fast suite CUDA build 190/190, SYCL build 183/183, HIP build 181 + 1 existing expected fail
All CUDA / HIP / SYCL ADM parity tests pass
Refactor commits, GPU output before vs after byte-identical (CUDA 11 inputs incl. 10-bit and 120-frame 1080p, HIP 12, SYCL 9)
Tidy, per lane, every touched TU 0 findings, 0 uncited NOLINTs, 0 compile failures
ninja -t deps after the depfile change integer_adm_cuda.h / integer_adm_hip.h tracked; a header touch rebuilds the kernels
Tooling unit tests 16 + 2 pass; the two new ratchet tests fail against the old ratchet
scripts/dev/preflight.sh (with #1475's fixes) gcc, clang, m32, msvcism, tidy, cppcheck pass; sanitizers fails only on check_exported_symbols flagging ASan's __start_asan_globals / __stop_asan_globals under GNU ld, the pre-existing false positive #1475 fixes

Known follow-ups

  • The GPU tidy baselines are stale as a whole (converted on 2026-09-02 from a measurement no lane could reproduce). Scoped writes cannot lower header entries, so integer_adm_cuda.h and integer_adm_hip.h keep a baseline of 4 although both measure 0. A full per-lane re-measure should follow this PR.
  • Upstream Netflix/vmaf's CUDA ADM carries the zero-shift rounding, the border reads and the stride swap; worth reporting with the port: reconcile with upstream Netflix/vmaf September 2026 (ADM/VIF SIMD fixes) #1473 findings.
  • The HIP and SYCL twins still emit no adm3 / aim (T-GPU-ADM-AIM-DEVICE-PASS-MISSING-SYCL-HIP-2026-09-05).

…MD fixes)

Takes the September upstream fixes the fork was exposed to and records
the ones it was not:

- 03b5562c5: adm_decouple_avx2's tail bound was a multiple of 8 from
  zero instead of from `left`, so the last vector store ran up to six
  int16 past `right` in 372 of 992 band widths. Taken verbatim.
- ea012e387: adm_dwt2_8_neon's vector loop had no tail bound and wrote
  one int16 past every band, into the next band's [0][0], from
  over-read data, which made NEON scores wrong and non-deterministic at
  small widths. Bounded with a scalar tail, as upstream does.
- cba9343ed: upstream now uses the four-tap column 0 everywhere, so the
  Apple-only three-tap wrapper is retired (ADR-1257). The three akiyo
  Darwin assertions in python/test/vmafexec_test.py are left for the
  maintainer to apply by hand, as the golden-file guard requires.
- 1786bd961: its adm_buffer_alloc zeroing is ported, and the read it
  masked is fixed at the source: for n_half == 2 (frame dimensions
  17..32) the DWT index tables restarted the mirrored tail at 0 and read
  row / column -1, which made integer_adm_scale3 vary between identical
  runs.

Instead of adopting checkasm, the parity tests gain its lesson: a
guard-band helper in simd_bitexact_test.h checks that nothing outside
the compared region is written, and the ADM and VIF SIMD tests sweep
small sizes. 8f7d50d29 (the fork's bound already covers every DWT2
kernel), c023bb7cb and the CI workflow need nothing; 86da14d03 is PR
#1472.
…aths

At frame widths 17 to 32 the scale-0 horizontal and vertical cube shift is
0. The AVX2 and AVX-512 paths computed its rounding constant as
(uint32_t)pow(2, shift - 1), which converts infinity to an integer. The
ASan+UBSan lane caught it in test_integer_adm_tiny_frames.

The AVX-512 build converts with vcvttsd2usi and got 0xFFFFFFFF. That was the
whole cause of T-ADM-AVX512-SMALL-WIDTH-SCALE0-2026-09-18: scale 0 was up to
0.01 off scalar at those widths and is now bit-exact. The AVX2 build got 0,
the scalar value, by accident.

All 16 constants call adm_half_shift(), which moves from integer_adm.c into
adm_csf_fixed_point.h next to a shared adm_frame_size_check(). Scalar and
AVX2 scores are unchanged, as is every frame wider than 32 pixels. A new
sweep compares the SIMD dispatch with scalar at widths 17 to 32. It fails at
17x70 with the old AVX-512 expression and passes on NEON under QEMU.

The same investigation found the CUDA and HIP twins mis-scoring these widths
and the GPU twins accepting frames below 17x17. That is recorded as
T-GPU-ADM-TINY-FRAME-SHIFT-2026-09-18 and fixed in a separate PR.
Work in progress: functional fix, tests and HIP test registrations. The lint cleanup of the touched GPU files, docs and state follow before the PR opens.
The GPU ratchet lanes could not run. tidy-ratchet.py passed each --extra-arg
value to clang-tidy as its own argument, so the lanes' --cuda-host-only and
-x hip were rejected as unknown clang-tidy options. It also handed clang-tidy
a relative -p and a relative wrapper path while running it inside the TU's
directory, which broke the sycl lane's wrapper. Both are fixed: extra
arguments reach the compiler through clang-tidy's --extra-arg, and the
database and binary paths are made absolute.

meson compiles .cu and .hip files through nvcc and hipcc custom targets,
which leave them out of compile_commands.json. The new
gen-gpu-compile-commands.py adds a clang++ entry per kernel file, with the
rule's include paths, defines and standard, and make tidy-ratchet now runs
it (and the existing SYCL generator) before measuring.

Tidy Ratchet self-tests the tooling on every PR; neither test file ran in CI
before.
…ssions

integer_adm_sycl.cpp carried nine clang-tidy suppressions without an ADR
citation. Measured with the sycl lane, six of the single-line ones no longer
suppressed anything: the branch-clone, the four const-correctness and the
unused-parameter markers all sat on code that no longer triggers the check.
They are removed.

The remaining block suppression hid misc-use-anonymous-namespace and
misc-use-internal-linkage on every file-local function, constant and type.
Its own rationale said an anonymous namespace would work, so it was a style
preference rather than a load-bearing invariant. The file's internals now sit
in one anonymous namespace; only the extern "C" extractor keeps external
linkage.

launch_dwt_hori_pair() took an h_add rounding term it never used: the kernel
derives 1 << (h_shift - 1), which equals every table value. The parameter
and the table field are gone.

SYCL integer ADM output on the Arc A380 is byte-identical to the previous
build on src01, the 1080p checkerboard, six tiny sizes and a non-default
option set. The sycl baseline drops from 9 findings and 9 uncited NOLINTs to
0 for this file.
…tidy lanes

User docs, state rows, changelog fragments, rebase note and AGENTS.md invariants for this branch.
…a skip

The file's header comment contained `integer_adm/*.cu`, which opened a nested
comment and made every build warn (-Wcomment). The CPU and CUDA runners now
share one open / score / close sequence, which takes the three oversized
functions under the lint profile's limits. The option-table check uses two
small helpers instead of an inline search.

Two behaviour fixes in the test itself: without a CUDA device the parity
checks now set mu_skipped, so the binary exits 77 and meson reports a skip
instead of a pass; and the options dictionary is freed when a run stops
before vmaf_use_feature() takes it over, including model_opts() failing
halfway, where it used to leak.

The C nullptr suggestions are covered by the repository's ADR-1138 band, as
in the other C tests. The cuda lane measures the file at zero findings;
test_cuda_adm_parity passes 4/4 on an RTX 4090.
…rnels

The GPU ADM tiny-frame fix touches adm_cm.hip, and every touched file has
to end lint-clean (ADR-0141, ADR-1142). This removes all 88 findings the
HIP tidy lane measured for the file without changing what the kernels
compute.

The two oversized compute kernels are split into __forceinline__ device
helpers: the 3x3 neighbourhood, the threshold sum, the masked value and
the scale-0 block reduction. The helpers and the WarpShift argument struct
move into an anonymous namespace. The three kernels stay extern "C",
because the host resolves them by name from the embedded HSACO
(ADR-0372). Kernel parameters a kernel does not read keep their slot and
lose their name, so the kernel signatures and the host launch-argument
arrays are unchanged. Locals are const, if/else chains are braced, row
offsets make their widening to a pointer offset explicit, <stdint.h>
becomes <cstdint> and the typedef becomes a using alias.

The gfx1036 code object keeps its kernel symbols, kernel-argument layout
and LDS and scratch sizes, and the reduce kernel's instructions are
unchanged. HIP integer ADM scores are identical before and after on the
Netflix src01 and 1-pixel checkerboard pairs, a 10-bit src01 run, runs
with non-default options, and synthetic frames from 18x18 to 640x360.
The HIP tidy baseline for the file drops from 84 to 0.
…code

The GPU ADM tiny-frame fix touches integer_adm_hip.c, and every touched
file has to end lint-clean (ADR-0141, ADR-1142). This removes the 33
findings the HIP tidy lane measured and the compiler warnings of both HIP
build flavours, without changing what the extractor computes.

The file no longer defines the reserved __HIP_PLATFORM_AMD__ macro.
hip_runtime_dep in core/src/hip/meson.build already compiles every HIP
translation unit with -D__HIP_PLATFORM_AMD__=1, which picture_hip.c
relies on too.

write_scores, integer_compute_adm_hip, adm_cm_device_hip, init_fex_hip
and close_fex_hip exceeded the function-size limits and are split into
helpers. Device setup in init moves into self-cleaning helpers that
release in the same order as the old goto chain, and the GET_FN macro
becomes a table of kernel names. The rfactor computation that init and
the per-frame path each carried is one helper. Launch-argument arrays
cast multi-level pointers to void * explicitly, RES_BUFFER_SIZE is a
size_t, and the result accumulators are passed as const.

write_scores is live code: collect() calls it when HAVE_HIPCC is set. It
was only unused in the scaffold build without hipcc, which warned about
it and about the unused dist_pic and s. The score helpers are now built
only with HAVE_HIPCC, and the scaffold branches discard what they do not
use. The unused val_per_thread in i4_adm_cm_device_hip is removed.

The error path taken when the feature-name dictionary cannot be created
is kept as it was and marked in a comment: it releases the device state
except the luma staging buffers and still returns 0.

HIP integer ADM scores are identical before and after on the Netflix
src01 and 1-pixel checkerboard pairs, a 10-bit src01 run, runs with
non-default options (including debug output and adm_skip_scale0), and
synthetic frames from 18x18 to 640x360. The HIP tidy baseline for the
file drops from 70 to 0.
…ot be built

init_fex_hip() released most of its device state but returned 0 when vmaf_feature_name_dict_from_provided_features() failed, so the next submit() ran against destroyed state and the luma staging buffers leaked. It now frees those too and returns -ENOMEM, like the CPU extractor; close() never runs after a failed init().
…indings

adm_cm.cu carried 164 clang-tidy findings. A bug-fix PR touches the file,
and the touched-file rule (ADR-0141, ADR-1142) requires it to end clean, so
this refactors it without changing what the kernels compute.

- Locals that are never reassigned are const; if/else chains get braces.
- Device helpers, the WarpShift type and the scale-0 kernel templates move
  into anonymous namespaces. Only the extern "C" kernels keep external
  linkage, because the host resolves them by name (ADR-0747).
- The four oversized kernel bodies are split into __device__
  __forceinline__ helpers: border neighbours, the cubic accumulation, the
  DLM and AIM thresholds, and the block and warp row reductions, which the
  DLM and AIM kernels now share. The AIM_NEIGHBOR / AIM_CENTER
  function-like macros become one i4_cm_weight() helper.
- Unused kernel parameters stay in the signatures, unnamed, so the host's
  void *args[] launch layouts are unchanged. The scale-0 templates drop
  them and take the kernel parameters by reference. Passing them by value
  made nvcc copy params and ws to a 296-byte local stack.
- Row offsets are computed once as int, as before, so no index arithmetic
  is widened.
- The negated 2^31 rounding term is now an explicit (int32_t) cast with an
  ADR-0155 note, which silences nvcc warning #68-D without changing the
  value.

Verification: PTX entry signatures are identical before and after. Integer
ADM scores are byte-identical on the Netflix src01 and checkerboard pairs,
a 10-bit src01 run, a debug/adm_skip_scale0 run, a non-default CSF/EGL
run, and synthetic 18x18 to 640x360 frames. The four test_cuda_adm_*
tests pass. Register use is unchanged or within 2 per kernel, the scale-0
kernels still use no stack, and i4_adm_cm_line_kernel_fused no longer
spills its row pointers (72-byte stack before, none now).

The CUDA tidy baseline for the file is tightened from 174 to 0.
…dy findings

integer_adm_cuda.c carried 46 clang-tidy findings, one NOLINT without an
ADR citation, and eight gcc warnings. A bug-fix PR touches the file, and
the touched-file rule (ADR-0141, ADR-1142) requires it to end clean, so
this refactors it without changing behaviour.

- Oversized functions are split into named helpers.
  integer_compute_adm_cuda becomes adm_fixed_parameters, adm_dwt2_scale0,
  adm_wait_for_pictures, adm_scale0_device and adm_scale123_device.
  init_fex_cuda becomes device setup (stream/events with the ADR-1090
  graduated unwind kept label for label), module loading, per-module kernel
  lookup, buffer allocation and buffer carving. write_scores becomes the
  DLM terms, the AIM numerator and the two feature-append helpers, which
  append in the same order. close_fex_cuda and the init error path share
  one buffer-free helper. adm_cm_device and adm_cm_aim_device share the
  scale-0 region and warp-shift helpers; the AIM rows-per-thread heuristic
  and its ADR-1226 measurements move to their own function.
- The 170-line NOLINTBEGIN(performance-no-int-to-ptr) region is replaced
  by one adm_device_ptr() helper with a single cited NOLINT. Turning a
  Driver API CUdeviceptr into a typed device pointer can't be refactored
  away while the fork dispatches through the Driver API (ADR-0747).
- vmaf_fex_integer_adm_cuda keeps external linkage for the extractor
  registry, with the tree's cited NOLINTNEXTLINE (ADR-0278).
- Unused parameters are removed from the static CSF-denominator launchers.
  submit's `index` is (void)-cast. The dead buffer_stride locals are
  dropped. The const scale that gcc flagged as discarded-qualifiers is
  non-const, like its DLM twin.
- Launch arrays cast pointer-to-pointer arguments explicitly and pass
  buf/p/d_dst directly instead of &*ptr. Result-slot offsets and the result
  buffer size are computed in size_t. ceil(log2f()) becomes
  ceilf(log2f()), which gives the same value for every float. The conclude
  helpers take const accumulators.
- The data_buf vmaf_cuda_buffer_get_dptr() result is now checked like its
  tmp_res sibling. It cannot fail there because the buffer was just
  allocated.
- The ADM modules are now unloaded in the same order on the close path as
  on the init error path, which the driver does not observe.

cppcheck (repository flags) goes from 28 to 0 style findings on the file.

Verification: integer ADM scores are byte-identical on the Netflix src01
and checkerboard pairs, a 10-bit src01 run, a debug/adm_skip_scale0 run, a
non-default CSF/EGL run, and synthetic 18x18 to 640x360 frames. The four
test_cuda_adm_* tests and the fast suite (190/190) pass. 1080p throughput
is unchanged.

The CUDA tidy baseline for the file is tightened from 136 to 0 warnings
and from 1 to 0 uncited NOLINTs.
…each picture's own stride

init_fex_cuda() leaked the stream, events and kernel modules when buffer allocation failed, because close() never runs after a failed init(); it now releases them. The device step's cleanup labels also destroyed the handle whose creation had just failed; one helper now destroys only the handles that exist. The 10/16-bit path took the reference stride from the distorted picture and vice versa (as upstream does); harmless while both share a stride, now correct. CUDA ADM output is byte-identical at 8 and 10 bits.
tmp_accum (3 * w * h * 8 bytes, 49.8 MB at 1080p) and tmp_accum_h were allocated at init and passed to the CM kernels as accum_per_block, which both kernels leave unnamed and unused since the reduction was fused into them. Both buffers and their struct fields are removed; the kernels keep their signatures and receive a null device pointer. CUDA ADM output is byte-identical on src01 at 8 and 10 bits and on the 1080p checkerboard.
…hanges

The nvcc and hipcc custom targets declared no depfile, so ninja tracked only the kernel source and a header edit left stale fatbins/HSACOs. Removing two AdmBufferCuda fields made every CUDA ADM test fail with an illegal address until the fatbin was deleted by hand. nvcc now writes a depfile (-MD -MF -MT; not on Windows, where it preprocesses through cl.exe), and hipcc's device frontend does via -Xclang -dependency-file, since hipcc ignores -MD under --genco.
… the CPU

The CPU's scale-0 pipeline stores its bands as int16_t, and so does the
CUDA twin. The SYCL twin computed csf_a, csf_f and the contrast-masking
threshold in 32 or 64 bits and never wrapped. On full-range content that
changes the score: with two frames of independent 8-bit noise at
576x324, integer_adm_scale0 was 2.1e-4 away from the scalar CPU, over
the 1e-4 cross-backend tolerance.

With the default CSF weights the value that wraps is the 1/15 centre tap
of the threshold (adm_cm_thresh): once |csf_a| reaches 15360 on the
horizontal or vertical band, the CPU's int16 store turns it negative.
The kernel now narrows that tap the same way. It also narrows csf_a and
csf_f as adm_csf does, rounds the diagonal csf_a with 65535 as the CPU
does, and evaluates the contrast measure modulo 2^32 as
adm_cm_accum_round does. Those only matter with larger CSF weights: with
Barten weights of 55849/55849/56864, scale 0 was 4.4e-3 off before this
change.

After the change, scale 0 is within 6.2e-8 of the scalar CPU on noise at
34x48, 96x64, 250x130, 576x324 and 1920x1080, and within 5.3e-8 with
the Barten weights. SYCL scores for src01 (48 frames) and both 1080p
checkerboard pairs are byte-identical before and after.
…e noise

The GPU ADM parity test only fed smooth ramps. Values from a ramp never
grow large enough at scale 0 to wrap the CPU's int16 stores, so the test
could not see the SYCL divergence fixed in the previous commit.

A second parity test now scores independent full-range 8-bit noise in
the reference and the distorted picture at 96x64 and 576x324, at the
same 1e-4 tolerance. The noise is a stateless integer hash of the
sample position, seeded per picture, so every backend and every run
sees the same frames. The file header now covers both tiny frames and
full-range content.

Without the SYCL fix the test fails on the Arc A380: at 576x324,
integer_adm_scale0 is 0.45212316 on the CPU and 0.45189860 on SYCL, a
difference of 2.25e-4. With the fix it passes on SYCL. It also passes on
CUDA (RTX 4090) and HIP (gfx1036 iGPU).
The CPU stores the rounding term of its scale 1-3 ">> 32" shifts,
1u << 31, in an int32_t, where it wraps to INT32_MIN (Netflix#955). So
every csf_f value, and every 1/15 centre tap of the masking threshold,
rounds with -2^31, and the Netflix golden values encode that
(ADR-0155). The CUDA and HIP twins reproduce the wrap. The SYCL twin
added +2^31, which made each of those terms exactly one higher than the
CPU's, and put scales 1-3 up to 1.0e-6 away from the scalar CPU on
noise.

The SYCL twin now adds -2^31 in both places. On noise at 34x48, 96x64,
250x130, 576x324 and 1920x1080, scales 1-3 are now within 6.4e-8 of the
scalar CPU. As a check, a local build that also used the CPU's float
score finalisation instead of SYCL's double one matched the scalar CPU
exactly on every ADM score, on all of that noise, on src01 and on both
checkerboard pairs. The difference of about 1e-7 that remains comes
from SYCL's double-precision host finalisation.

This changes SYCL scores on ordinary content, and every changed value
moves toward the CPU. On src01, integer_adm2 and integer_adm_scale1..3
change in all 48 frames by at most 6.1e-7, and the largest difference
from the CPU drops from 7.8e-7 to 2.3e-7. On the 1-pixel checkerboard
pair they change in all three frames by at most 2.6e-7, and the largest
difference drops from 4.1e-7 to 1.6e-7. Scale 0 and the 10-pixel
checkerboard pair are unchanged.
Add T-SYCL-ADM-INT16-SEMANTICS-2026-09-18 to the recently closed rows of
docs/state.md. The row gives the before and after measurements, the CPU
narrowing points the SYCL twin now reproduces, and the stores that
cannot overflow and so need none. Add the changelog fragment and
re-render CHANGELOG.md. Note the int16 and Netflix#955 invariants for
agents in core/src/feature/sycl/AGENTS.md.
meson names a custom target's rule CUSTOM_COMMAND_DEP once it has a depfile, so gen-gpu-compile-commands.py found no kernels after the kernel targets gained depfiles. The rule pattern accepts both names; the test covers a depfile target and checks its dependency flags stay out of the analysis command.
…riants

The SYCL twin now mirrors the CPU's int16 narrowings, the 65535 diagonal rounding and the Netflix#955 -2^31 rounding term (ADR-0155, entry 0048); a sync that changes them in integer_adm.c must change the SYCL twin too.
@lusoris

lusoris commented Sep 20, 2026

Copy link
Copy Markdown
Contributor Author

Absorbed into the ADM stack train #1507, per your direction to fold this stack the way #1506 was folded.

This PR targeted the one below it in a five-deep stack, so none of the five could merge until every one below had merged and been restacked — five sequential rebase-plus-CI rounds. #1507 is one. Your work is in it unchanged; that PR's description lists the six defects the fold itself surfaced, none of which an individual PR could see, because each gate only looks at the files its own PR touches.

The branch stays on the remote.

@lusoris lusoris closed this Sep 20, 2026
lusoris added a commit that referenced this pull request Sep 22, 2026
…ny frames, HIP buffer by pointer (#1507)

Integration train for the five-deep ADM/GPU stack behind #1473, folded into one
merge per maintainer direction. Each PR previously targeted the one below it, so
none could merge until every one below had.

- #1474 AVX2 / AVX-512 contrast masking wraps like scalar on full-range noise
- #1477 CPU 16-bit vertical DWT sum formed in int64 — int32 overflows at 16 bpc
  once three samples reach 42456
- #1476 Integer-ADM twins agree with the CPU on tiny frames; SYCL wrap and
  rounding fixes; GPU tidy-lane repairs
- #1478 The same 16-bit scale-0 vertical DWT fix in the CUDA, HIP and Metal twins
- #1481 HIP ADM kernels take `AdmBufferHip` by pointer (ADR-0759) rather than
  copying 328 bytes of arguments per launch

Required Checks Aggregator green on 9bae48c with no failing check; branch level
with master at 371ff58. Merged with admin bypass because the repository has a
single collaborator who cannot self-approve (ADR-1252).

Unblocks #1518, which depends on the SIMD contrast-masking fix — see
T-CUDA-ADM-SMALL-BORDER-PARITY-2026-09-21 in docs/state.md.
lusoris added a commit that referenced this pull request Oct 1, 2026
…23e8f2 (#1761)

* docs: record the fork's check of the upstream defects verified on 6ec23e8f2

Fifteen defects reproduced on Netflix master were run against the fork:
three reproduced and are fixed (#1305, #1420 as a hang, #1613), two are
documented (#910, #755 and #1180), ten are not affected. The dated section
in known-upstream-bugs.md and the Confirmed not-affected rows of the state
ledger carry the evidence; Netflix 8e7a1ac4e (revert of #1476) needs
nothing from the fork.

* docs: regenerate the indexes and the citation map after rebasing
@lusoris
lusoris deleted the fix/gpu-adm-tiny-frames branch October 6, 2026 08:29
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

type:bug Something isn't working

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant