Fix out-of-bounds boundary reflection for tiny frames in motion feature by kjg0724 · Pull Request #1582 · Netflix/vmaf

kjg0724 · 2026-08-17T05:48:39Z

Bug

mirror(int idx, int size) reflects row/column indices at frame boundaries for the
5-tap (radius-2) motion filter. It's a single-bounce reflect-101 implementation:

static inline int mirror(int idx, int size)
{
    if (idx < 0) return -idx;
    if (idx >= size) return 2 * size - idx - 2;
    return idx;
}

For size < 3 this returns indices outside [0, size) — e.g. mirror(-2, 1) = 2,
out of range for a 1-row/1-column plane. The identical function was independently
duplicated in 4 places (integer_motion.c, x86/motion_avx2.c, x86/motion_avx512.c,
arm64/motion_neon.c), plus a differently-structured CUDA version, plus the shared
float convolution helpers used by float_motion and (via convolution_f32_avx_*)
vif_tools.c.

The float convolution path additionally has an out-of-bounds write: with radius 2,
convolution_x_c_s's borders_right = width - 3 goes negative for width < 3, so the
trailing edge loop starts before dst[0]. The equivalent bug exists in
convolution_avx.c's vertical pass (i_vec_end = height - radius goes negative,
writing tmp[-tmp_stride + j]).

This isn't only a 1×1-frame issue. vif_tools.c calls convolution_f32_avx_s with
filter widths up to 17 (radius 8), so on any AVX2 host the same OOB write triggers for
any dimension below 8, at every VIF scale — not just degenerate frames.

Fix

motion_mirror() (new, replacing 4 duplicated mirror() copies): reflects
repeatedly instead of once, so any size >= 1 is safe. size == 1/size <= 0
degenerate to 0 (the unique valid extension of a 1-sample signal — matches
OpenCV's BORDER_REFLECT_101 and scipy's mode='mirror').
Renamed from mirror to motion_mirror and consolidated into integer_motion.h
(all 4 CPU files already include it). The rename is required, not stylistic: the CUDA
translation unit also includes this header and already owns a mirror symbol with
different (intentionally unrelated, untouched) reflect-even semantics — reusing the
bare name would collide under nvcc.
convolution_mirror() (new, in convolution_internal.h): same fix, kept separate
from motion_mirror since this header serves other float features with their own
style and no existing dependency on integer_motion.h.
CUDA's own mirror(): added if (sup == 1) return 0; — one line, its sup >= 2
formula is untouched.
convolution.c / convolution_avx.c: clamped the edge-loop bounds so tiny
width/height can't write out of bounds, without changing anything for normal sizes.

Non-regression is structural, not just tested. For the real call contract
(idx in [-2, size+1], 5-tap/radius-2), a negative overshoot lands in [1,2] and a
positive one in [size-3, size-1] whenever size >= 3 — both already valid, so the
new multi-bounce loop takes exactly one iteration and computes the exact same
expression as the old code. This is proven exhaustively in tests (below), not just
argued in the abstract.

Testing

Exhaustive old-vs-new formula proofs for all three reflection implementations
(motion_mirror, convolution_mirror, and a host-side transcription of the CUDA
formula), comparing every size in [3, 8192] against every in-contract offset —
mechanically proves zero behavior change for size >= 3.
Pipeline tests at 8 tiny sizes (1×1 through 8×2) across multiple bit depths,
run through the real motion/float_motion feature extractors, asserting success,
finite scores, and — for each size — that the score obtained via forced-scalar
dispatch matches the score obtained via best-available SIMD dispatch (bit-exact for
the integer path; a documented, generously-tight tolerance for the float path, since
AVX's -mfma can contract the shared edge-filter's multiply-add into an FMA and
change the last bit or two versus the non-FMA scalar build — unrelated to this fix).
This is the only runtime coverage the AVX2/AVX512 motion kernels get from this test
suite, since no x86 hardware was available anywhere in developing this fix (see
below).
Canary-buffer bounds tests for the OOB-write fix: guard regions before/after
the real output buffer, asserted intact after the call, for the scalar convolution
path (verified with a negative control — temporarily removing the clamps reproduces
the failure) and, #if ARCH_X86-gated, the AVX path.
Live ASan red-then-green demonstration, reproduced independently on two hosts/
toolchains (macOS/Apple Clang arm64, Ubuntu 22.04/GCC 11.4.0 aarch64): built with
-Db_sanitize=address,undefined, confirmed clean on the fix, then temporarily
reverted just the two mirror function bodies to the old formula and confirmed a
real heap-buffer-overflow abort in both the integer path
(motion_score_pipeline_8) and the float path (convolution_y_c_s) — then restored
and reconfirmed clean.
Full meson test suite: 25/25 on both hosts.
Python end-to-end golden-score regression suite (quality_runner_test.py): 67/67,
run twice, on Ubuntu/GCC — no VMAF score on real video content moved.

What's not verified at runtime

No x86 or CUDA hardware was available anywhere in developing this fix.

x86 (AVX2/AVX512): verified only via cross-target clang -target x86_64-... -mavx2/-mavx512f -fsyntax-only/-c checks (proves it compiles and type-checks, not
that it's correct at runtime) plus the exhaustive formula proofs and mechanical
rename review. This PR's own test suite gives these kernels real coverage the moment
CI runs it on an x86 runner.
CUDA: the one-line guard is exhaustively proven inert for sup >= 2 via a
host-side transcription, but the actual .cu file has never been compiled. This
repo's CI doesn't build the CUDA path, so a maintainer with CUDA access verifying the
guard would be appreciated.

The 5-tap radius-2 boundary reflection helper used by the integer motion feature was independently duplicated across four CPU implementations (scalar, AVX2, AVX512, NEON), all sharing the same single-bounce formula that reads out-of-bounds indices when width or height is less than 3 (Netflix#1580). Replace the four local copies with one shared, loop-based reflect-101 implementation in integer_motion.h, which all four TUs already include. Named motion_mirror rather than mirror to avoid colliding with the unrelated __device__ mirror() in cuda/motion_score.cu, which also includes this header transitively and implements different boundary semantics that are out of scope here. For size >= 3 with the offsets these callers ever pass (radius 2), the new implementation takes at most one reflection and is bit-identical to the old formula; only sizes below 3 change behavior, from out-of-bounds reads to well-defined reflect-101 indices. Also removes edge_16(), an unused dead-code helper in the same header that shared the old broken formula and had no callers.

Exhaustively checks that motion_mirror matches the old single-bounce formula for every size in [3, 8192] and every offset in the 5-tap radius-2 call contract, proving the consolidation changed no output for normal frame sizes. Also pins down the exact reflect-101 values for size 1 and 2, the cases the fix specifically targets.

The CUDA device-side mirror() in motion_score.cu has the same class of out-of-bounds bug as the CPU implementation for degenerate 1-pixel dimensions, but with a different (reflect-even) formula. Add a minimal guard so sup==1 maps every index to 0, without altering the existing reflect-even behavior for sup>=2. Also extends test_motion_mirror.c with a host-side transcription of both the pre- and post-fix CUDA formulas, proving the fix is a no-op for sup>=2 and degenerates safely for sup==1. This does not compile or link any CUDA code.

…ilter The float convolution path shared by ADM, VIF and float motion assumed width and height are at least filter_width, and misbehaved below that. convolution_internal.h reflected boundary taps with a single bounce, which still lands outside [0, size) once size is smaller than the filter radius. Replace it with convolution_mirror(), a repeated whole-sample symmetric reflection that is safe for any size >= 1 and bit-identical to the previous formula for size >= radius + 1. convolution.c and convolution_avx.c additionally computed negative edge-loop bounds for such sizes, so the trailing edge loops started at a negative index and wrote before the start of dst/tmp. Clamp those bounds. In convolution_avx.c the clamped horizontal trailing start is kept in a separate variable so that the j_end argument of convolution_f32_avx_s_1d_h_scanline stays unraised; raising it would push the scanline's own vector stores past the end of dst. All clamps are no-ops for sizes at or above the filter width, and for the sizes in between they only drop redundant writes of identical values, so the computed output is unchanged.

Exhaustively checks that convolution_mirror matches the single-bounce reflection it replaces for every size in [3, 8192] and every offset in the 5-tap radius-2 call contract, proving the float convolution path produces identical output for normal frame sizes.

The comment claimed the vertical trailing loop's start is i_vec_end's only other use. It is not: i_vec_end is also the end bound of the vector loop. State the actual reason the in-place raise is safe, so the comment still guards against a maintainer collapsing the clamp into a bug. Comment-only change; no behavior change.

The sizes below the 5-tap filter's radius were excluded because the single-bounce mirror() went out of bounds there, making both the scalar and the SIMD score meaningless. Now that the reflection is safe for any size >= 1, those sizes must agree bit-exactly like every other size.

Drive both the integer and the float motion feature extractor end to end at 1x1 through 8x2, for 8, 10, 12 and 16 bpc and for both the scalar and the SIMD dispatch. There is no meaningful reference score at these sizes, so what is asserted is that extraction succeeds, the score is finite and two identical runs agree bit for bit; the memory-safety half of the check comes from running the same binary under AddressSanitizer. The float sub-test skips itself when the library is built with -Denable_float=false, in which case the extractor is not registered.

convolution_f32_c_s() delegates to the AVX kernel and returns as soon as AVX2 is available, so driving the float motion extractor with the CPU flags left at their default only reached convolution_f32_avx_s() on an x86 host. The scalar convolution_y_c_s()/convolution_x_c_s() pair went unexercised there, which is precisely where the boundary reflection had to be corrected. Run the float sizes under both a zeroed and an unrestricted CPU mask, as the integer test already did, and share the mask table between the two.

motion_mirror() and convolution_mirror() special-cased size == 1, where the reflection period 2*(size-1) is zero and the loop would not terminate. Widen the guard to size <= 1 so the functions are total rather than merely correct over their call contract. Both helpers are only ever reached from loops bounded by the same size (for i < h / for j < w in integer_motion.c and the AVX2/AVX-512 motion kernels, and the clamped border loops in convolution.c / convolution_avx.c), so size is always >= 1 in practice and no reachable result changes.

The tiny-size pipeline tests ran the extractor twice per CPU mask and only asserted the two runs agreed, which proves determinism but never compares one dispatch against another. The x86 AVX2/AVX-512 motion kernels therefore had no runtime verification at all; upstream CI runs meson test on an x86_64 runner, so a cross-mask comparison is the one place those kernels get exercised. Capture the cpu_masks[0] (mask 0, forced scalar) score for each (size, bpc) and assert every other mask reproduces it. The integer path is compared exactly, as test_motion_neon.c already does for NEON. The float path is compared with a relative tolerance instead: convolution_avx.c is built with -mfma while convolution.c is not, so the compiler contracts the shared convolution_edge_s() accumulation into fused multiply-adds in the AVX translation unit only and the two results legitimately differ in the last bits. A boundary regression moves the score by O(1) and is still caught.

The border clamps in convolution.c and convolution_avx.c fix an out-of-bounds write, not just an out-of-bounds read: below the filter width the unclamped borders_left/borders_right (and the AVX i_vec_end/j_vec_end) went negative and the trailing edge loop stored outside dst. Only a sanitizer would have caught a regression there, and no sanitizer job exists in CI, so a later "simplification" back to the buggy form would leave CI green. Surround every destination buffer with guard elements holding a sentinel and assert the sentinel survives. The scalar test covers convolution_x_c_s() and convolution_y_c_s() over dimensions 1..4 and runs on every host. The AVX test covers convolution_f32_avx_s(), _sq_s() and _xy_s(), guarding both dst and the caller-supplied tmp scratch, and is compiled only under ARCH_X86 so it is inert elsewhere. Verified against the pre-fix code by removing each clamp in turn: the scalar test reports "convolution_x_c_s wrote before dst" and "convolution_y_c_s wrote before dst" respectively.

…holes Three HIGH findings from an independent adversarial review of this harvest. Each was reproduced before being fixed, and the regression test fails on the pre-fix tree. 1. The Netflix/vmaf#1582 border clamp landed only on the scalar path. convolution_f32_c_s dispatches straight into convolution_f32_avx_s whenever AVX2 is present — every CI runner and the dev workstation — so the clamp this PR added was dead code on x86. The AVX2 and AVX-512 twins derive the same `height - radius` split at three sites each and kept it unclamped. For a plane shorter than the radius that value is negative, so the trailing border loop starts at a negative row and the leading one runs past the end. Both are heap WRITES, not reads. All six sites now share the scalar clamp, which moved into convolution_internal.h as a static inline. 2. motion_filter_size=1 bypassed the minimum-dimension guard entirely. motion_check_min_dim gated its check on `effective_filter_size > 1`, but motion_blur_plane keeps filter_size = 5 for that value and only swaps in the FILTER_5_NO_OP_s coefficients, so the radius is still 2. A 1-row plane therefore reached the convolution in (1) through a documented public option with range 0..9. The guard now mirrors motion_blur_plane exactly: 3 taps only for motion_filter_size == 3, otherwise 5. 3. Odd-height 4:2:0 chroma planes were under-allocated by one row. motion_chroma_heights used `h / 2` while picture.c and the guard both use the ceiling `(h + 1) >> 1`, so motion_copy_and_blur overran ref, tmp and every MOTION_BLUR_RING blur buffer for both U and V. Even heights were unaffected, which is why neither golden fixture caught it. Also removes a stray `} // namespace` inside the _WIN32 block of core/tools/vmaf.cpp that closed a namespace never opened. It broke every Windows build and was invisible on Linux, where the preprocessor drops the block. This PR is a draft and drafts run no CI here, so nothing had compiled it. The file now has exactly one namespace opener and one closer, neither inside any conditional. New test core/test/test_motion_convolution_oob.c drives float_motion through the public vmaf_read_pictures entry point, because no existing test could reach the dispatched SIMD path: test_motion_min_dim only calls init(), and test_convolution_edge_small calls the scalar kernels directly. Verified both ways. The new test fails on the pre-fix tree and passes after. Under -Db_sanitize=address the pre-fix tree reports "heap-buffer-overflow ... WRITE of size 4" in convolution_f32_avx_s reached from vmaf_read_pictures, "0 bytes after 32-byte region" — the single-row buffer. Post-fix: zero sanitizer reports, the guard returns -EINVAL, and the odd-height case scores cleanly. meson test --suite=fast: 111 Ok, 0 Fail. Netflix golden gate: 271 passed, 12 skipped, 0 failed. The twelve MEDIUM and LOW findings from the same review are not addressed here and remain open on the PR. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

…ADR-1166) Verify a batch of long-open Netflix/vmaf issues against this tree — the fork diverged far enough (ADR-0700's `libvmaf/` -> `core/` rename, several C-to-C++ conversions, four fork-added GPU backends) that an upstream report is neither automatically applicable nor automatically stale — and fix the subset that still bites. The full triage table, including the ALREADY-FIXED and NOT-APPLICABLE verdicts, is in docs/research/1166-upstream-issue-harvest-2026-09-03.md. Memory safety, all reachable from the public C API today: * reported upstream as Netflix/vmaf#1582 (mirror half also Netflix/vmaf#1581): the reflect-101 fold in convolution_edge_s / _sq_s / _xy_s bounced an out-of-range tap exactly once, which only lands in range for size >= radius + 1; and convolution_x_c_s / convolution_y_c_s derived the trailing border bound as dim - (filter_width - radius), which goes negative for a plane narrower than the filter and starts the trailing loop at a negative index (heap underflow write). Two live paths reached those sizes: `--feature float_vif` on 9..15px frames (the guard admitted >= 9 but the four-scale ladder needs >= 16 — the binding constraint is scale 3), and `--feature float_motion` with motion_add_uv on a 4x4 YUV420P frame (the guard validated luma only while the blur runs at the 2x2 chroma dimensions). The fold is now iterative and bit-identical to the single bounce for every in-contract size; the borders are clamped; float_vif derives its floor from vif_get_min_dim(kernelscale); float_motion validates every plane it convolves. * reported upstream as Netflix/vmaf#1580: the three fork-added Metal motion extractors were written after the Research-0094 sweep and never got the min-dim guard, so a 1- or 2-pixel-tall frame read out of bounds on device. Correctness and contracts: * reported upstream as Netflix/vmaf#1242: vmaf_model_feature_overload() leaked the caller's dictionary on the -ENOMEM path, vmaf_model_collection_feature_overload() swallowed the copy error and dereferenced *model_collection unchecked, and feature.h / model.h documented opposite ownership rules — one of the two readings a latent double free. All three public headers now state the implemented contract identically. Supersedes ADR-0806. * reported upstream as Netflix/vmaf#1551, which retracts Netflix/vmaf#1422: the MSVC __builtin_clz shim used __lzcnt, which emits LZCNT with no runtime gate and silently retires as BSR on any x86-64 without ABM — wrong VIF and ADM log2 shifts, no fault, and invisible to CI because every hosted Windows runner has LZCNT. Now _BitScanReverse, with an architecture guard so an MSVC ARM64 leg compiles. User-visible surfaces: * reported upstream as Netflix/vmaf#743: the CLI wrote UTF-8 braille and a CSI erase to a Windows console it never configured, so the progress line was mojibake under every default code page. The console is switched to UTF-8 + VT for the run and restored on exit, with an ASCII fallback. * reported upstream as Netflix/vmaf#1178: libvmaf.pc omitted the C++ runtime, so `pkg-config --static --libs libvmaf` produced a link line that fails with hundreds of undefined references — the reason ADR-0198's static FFmpeg reproducer had to add -lstdc++ by hand. * reported upstream as Netflix/vmaf#1573: the nvcc fatbin include list used relative paths that stopped resolving at ADR-0700, and three shell-driven tool tests declared no `depends`, so a subset run built nothing and exited 127. Behaviour changes: float_vif now rejects frames below 16px in either dimension, and float_motion with motion_add_uv rejects sub-minimum chroma planes. Both convert previously undefined behaviour into a documented -EINVAL. Regression tests: core/test/test_convolution_edge_small.c (NaN-poisoned guard buffers; fails pre-fix), core/test/test_compat_clz.c, core/test/test_model_feature_overload_ownership.c, core/test/test_spinner.cpp, scripts/ci/check-msvc-clz-shim.sh (fails pre-fix), plus extended cases in test_motion_min_dim.c and test_float_vif_min_dim.c, and a real static link in the libvmaf-build-matrix pkg-config step. Netflix golden scores unchanged: 76.66744 / 35.070245 / 7.985956 (271 passed, 12 skipped). Confirmed but not batched, one docs/state.md row each: Netflix/vmaf#1564, #930, off-by-one found while triaging Netflix/vmaf#1580. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>