lusoris · Jun 4, 2026 · Jun 4, 2026
@@ -0,0 +1,21 @@
+### Fixed
+
+- **CUDA VIF `filter1d.cu` 16-bit vertical rd-filter upper-bound typo** — the
+  upper-bound guard in the 16-bit vertical kernel wrote `fwidth_rd - fwidth_rd`
+  (always zero) instead of `fwidth - fwidth_rd`, causing the rd-filter window
+  to span all `fwidth` taps and index `vif_filt.filter[scale+1]` 1–4 entries
+  beyond allocation. Result: OOB reads and wrong VIF scores at scales 0–2 on
+  the CUDA backend. Fix mirrors the correct 8-bit form at line 183
+  (`fi < (fwidth - (fwidth - fwidth_rd) / 2)`).
+  File: `core/src/feature/cuda/integer_vif/filter1d.cu`.
+
+- **CUDA integer ADM CM operator-precedence bug (×2)** — two reduction loops in
+  `integer_adm/adm_cm.cu` (lines 373 and 712) computed `x_sq` as
+  `+ add_shift_sq >> shift_sq` (parsed as `+ (add_shift_sq >> shift_sq)` = `+ 0`
+  because `>>` binds tighter than `+`), silently dropping the shift normalisation.
+  `x_sq` carried an un-normalised squared value (~10^6 instead of ~0–1), causing
+  int32 overflow and wrong ADM scale-0 and AIM scores on the CUDA backend.
+  Fix adds the required parentheses: `((int64_t)accum * accum + add_shift_sq) >> shift_sq`,
+  matching the CPU reference macro `I4_ADM_CM_ACCUM_ROUND` at `integer_adm.c:743`
+  and the correct fused kernel at line 259.
+  File: `core/src/feature/cuda/integer_adm/adm_cm.cu`.
@@ -647,6 +647,24 @@ kernel variants at runtime. The current policy table is in ADR-0753.
   does NOT affect `adm_decouple.cu`.
   See [ADR-0763](../../../../docs/adr/0763-cuda-adm-decouple-ldg.md).
 
+- **`integer_adm/adm_cm.cu` `x_sq` reduction requires explicit parentheses around
+  `add_shift_sq` before the right-shift (r6-cuda-kernel / 2026-06-04).** The expression
+  `(int64_t)accum * accum + add_shift_sq >> shift_sq` is parsed by C++ as
+  `+ (add_shift_sq >> shift_sq)` = `+ 0` because `>>` binds tighter than `+`. The
+  correct form is `((int64_t)accum * accum + add_shift_sq) >> shift_sq`, which matches
+  the CPU reference macro `I4_ADM_CM_ACCUM_ROUND` in `integer_adm.c:743` and the fused
+  kernel at `adm_cm.cu:259`. This defect affected two reduction loops in the file
+  (lines 373 and 712). On rebase: if either loop is modified, verify the parenthesisation
+  of the `x_sq` computation before pushing.
+
+- **`integer_vif/filter1d.cu` 16-bit rd-filter upper-bound guard must use
+  `(fwidth - fwidth_rd)`, not `(fwidth_rd - fwidth_rd)` (r6-cuda-kernel / 2026-06-04).**
+  The correct guard is `fi < (fwidth - (fwidth - fwidth_rd) / 2)`, matching the 8-bit
+  form at line 183. Writing `(fwidth_rd - fwidth_rd)` (always zero) widens the tap
+  window to all `fwidth` taps and causes OOB reads into `vif_filt.filter[scale+1]`.
+  On rebase: if the vertical-pass loop in the 16-bit path is modified, verify the
+  upper-bound guard expression before pushing.
+
 - **`integer_adm/adm_csf.cu` and `integer_adm/adm_cm.cu` carry the F3 `__ldg()` fix on the
   active path (ADR-0773).** The six inline `__device__` helpers in `adm_cm.cu`
   (`inline_i4_csf_a`, `inline_i4_decouple_r`, `inline_s0_csf_a`, `inline_s0_decouple_r`,

@@ -370,7 +370,7 @@ adm_cm_line_kernel(AdmBufferCuda buf, int h, int w, int top, int bottom, int lef
     for (int row = 0; row < rows_per_thread; ++row) {
         int32_t accum_thread = accum_thread_reg[row];
         const int32_t x_sq =
-            (int32_t)((((int64_t)accum_thread * accum_thread) + add_shift_sq >> shift_sq));
+            (int32_t)(((((int64_t)accum_thread * accum_thread) + add_shift_sq) >> shift_sq));
         accum += (((int64_t)x_sq * accum_thread) + add_shift_cub) >> shift_cub;
     }
 
@@ -708,8 +708,8 @@ adm_cm_aim_line_kernel(AdmBufferCuda buf, int h, int w, int top, int bottom, int
 
     for (int row = 0; row < rows_per_thread; ++row) {
         int32_t accum_thread_val = accum_thread_reg[row];
-        const int32_t x_sq =
-            (int32_t)((((int64_t)accum_thread_val * accum_thread_val) + add_shift_sq >> shift_sq));
+        const int32_t x_sq = (int32_t)((
+            (((int64_t)accum_thread_val * accum_thread_val) + add_shift_sq) >> shift_sq));
         accum += (((int64_t)x_sq * accum_thread_val) + add_shift_cub) >> shift_cub;
     }
 

@@ -554,7 +554,7 @@ filter1d_16_vertical_kernel(VifBufferCuda buf, uint16_t *ref_in, uint16_t *dis_i
                     accum_dis[off] += img_coeff_dis * (uint64_t)imgcoeff_dis;
                     accum_ref_dis[off] += img_coeff_ref * (uint64_t)imgcoeff_dis;
                     if (fi >= (fwidth - fwidth_rd) / 2 &&
-                        fi < (fwidth - (fwidth_rd - fwidth_rd) / 2) && fwidth_rd > 0) {
+                        fi < (fwidth - (fwidth - fwidth_rd) / 2) && fwidth_rd > 0) {
                         const uint16_t fcoeff_rd =
                             vif_filt.filter[scale + 1][fi - ((fwidth - fwidth_rd) / 2)];
                         accum_ref_rd[off] += fcoeff_rd * imgcoeff_ref;

@@ -59,6 +59,21 @@ upstream. No Netflix golden assertions or upstream-mirrored Python are touched.
 
 ---
 
+## fix/cuda-vif-filter1d-adm-cm-opprec (2026-06-04)
+
+**Files touched:**
+`core/src/feature/cuda/integer_vif/filter1d.cu`,
+`core/src/feature/cuda/integer_adm/adm_cm.cu`
+
+no rebase impact: pure kernel arithmetic fixes. No public C API header, no
+meson build option, no FFmpeg patch surface, and no upstream-mirrored Python
+file is touched. The fixes correct two silent arithmetic defects (a typo in
+the rd-filter upper-bound guard in `filter1d.cu` and a missing parenthesis
+pair in two `x_sq` reduction loops in `adm_cm.cu`). Cross-backend SYCL/HIP/
+Vulkan ADM and VIF twins do not carry the same expressions and are unaffected.
+
+---
+
 ## docs/vulkan-overview-mark-removed-adr0726 (2026-06-04)
 
 **Files touched:**

@@ -2,6 +2,7 @@
 _Updated: 2026-06-04 (T-SIMD-PSNR-16BIT-SCALAR-TAIL-OVERFLOW-2026-06-04 closed — scalar-tail loops in `psnr_sse_line_16_avx2`, `psnr_sse_line_16_avx512`, and `psnr_sse_line_16_neon` computed squared error as `(int32_t)e * (int32_t)e` where `e` can reach ±65535 — signed-integer overflow UB (65535² > INT32_MAX, UBSan-flagged). Fix: replace with `(uint64_t)(uint32_t)abs((int32_t)ref[j] - (int32_t)dis[j]) * e` unsigned-multiply pattern, mirroring `sse_line_16_c` in `integer_psnr.c`. No ADR per CLAUDE §12 r8. Row added to Recently closed.)_
 _Updated: 2026-06-04 (T-CPU-SCORING-NAN-UB-GUARDS-2026-06-04 opened — Round-6 audit surfaced 11 CPU-path scoring edge-case bugs across APSNR (log10(0)), MS-SSIM (pow(neg,frac) + size overflow), float SSIM/MS-SSIM (convert_to_db NaN), iqa/ssim_tools.c (assert abort), ADM (score_aim uninit + harmonic-mean NaN + skip_scale0 sentinel), MOTION (wrong bilinear stride + OOM crash), and CAMBI (v_band_size uint16_t overflow). All fixed in one cluster PR. DRAFT PR: fix/r6-cpu-scoring-nan-ub-guards. ADR-1033. Row will be moved to Recently closed when PR merges.)_
 _Updated: 2026-06-04 (T-VMAF-INIT-DOUBLE-INIT-GUARD-2026-06-04 closed — three HIGH-severity API-contract bugs fixed: (1) vmaf_init unconditionally overwrote *vmaf, leaking the old context on double-init; fix: return -EINVAL when *vmaf is non-NULL. (2) vmaf_close left the caller pointer dangling; fix: document pointer-invalidity contract and null-after-close pattern in the public header. (3) DNN vmaf_dnn_session_open returned error on missing int8 sidecar instead of falling through to fp32; fix: rc=0 + fall-through. ADR-1032. Rows added to Recently closed.)_
+_Updated: 2026-06-04 (T-R6-CUDA-VIF-FILTER1D-WIDTH-GUARD-2026-06-04 and T-R6-CUDA-ADM-CM-OPERATOR-PRECEDENCE-2026-06-04 closed — two HIGH-severity CUDA kernel arithmetic bugs fixed: (1) `filter1d.cu` 16-bit vertical rd-filter upper-bound guard typo `fwidth_rd - fwidth_rd` → `fwidth - fwidth_rd`, preventing OOB reads into `vif_filt.filter[scale+1]`; (2) `adm_cm.cu` lines 373+712 `x_sq` operator-precedence defect `+ add_shift_sq >> shift_sq` → `+ add_shift_sq) >> shift_sq` matching CPU reference macro. PR: fix/cuda-vif-filter1d-adm-cm-opprec.)_
 _Updated: 2026-06-04 (T-R5-MEMORY-ORDERING-2026-06-04 closed — three HIGH-severity concurrency bugs: acq_rel ref-count ordering, mutex-destroy-after-unlock in feature_collector, picture_pool slot copy before unlock. ADR-1020. Row added to Recently closed.)_
 _Updated: 2026-06-04 (T-Y4M-DST-BUF-READ-SZ-OVERFLOW-2026-06-04 closed — five `dst_buf_read_sz` arithmetic expressions in `core/tools/y4m_input.c` used bare `int * int` multiplication for chroma branches 420/420jpeg/420mpeg2, 420p10, 420p12, 422p10, and 422p12. Signed overflow underallocated `dst_buf_read_sz` vs `dst_buf_sz`; subsequent `fread` at line 931 could overrun the heap buffer. Fix: `(size_t)` casts on all five expressions. ADR-1022. Row added to Recently closed.)_
 _Updated: 2026-06-04 (T-CPP-STD-C23-BUMP-INCONSISTENCY-2026-06-04 closed — cpp_std=c++11 project default inconsistent with C++23 wave files; test_feature_collector_coverage link-fail fixed. ADR-1003. Row added to Recently closed.)_
@@ -261,6 +262,10 @@ landed fix yet._
 
 | **T-SIMD-PSNR-16BIT-SCALAR-TAIL-OVERFLOW-2026-06-04** | Scalar-tail loops in `psnr_sse_line_16_avx2` (`core/src/feature/x86/psnr_avx2.c:105`), `psnr_sse_line_16_avx512` (`core/src/feature/x86/psnr_avx512.c:109`), and `psnr_sse_line_16_neon` (`core/src/feature/arm64/psnr_neon.c:96`) computed squared error as `(int32_t)e * (int32_t)e` where `e` can reach ±65535 for 16-bit input. The product 65535² = 4 294 836 225 exceeds INT32_MAX (2 147 483 647), invoking signed-integer overflow — undefined behaviour under C99/C11 that UBSan flags and that can produce wrong SSE values on optimising compilers. Fix: use `const uint32_t e = (uint32_t)abs((int32_t)ref[j] - (int32_t)dis[j])` and accumulate as `(uint64_t)e * e`, mirroring the `sse_line_16_c` reference in `integer_psnr.c`. `<stdlib.h>` added for `abs()` in all three files. No ADR per CLAUDE §12 r8. | no ADR | fix/simd-psnr-16bit-scalar-tail-overflow | `meson test -C build --suite=fast` — all fast tests pass; UBSan build of the scalar tail with input (ref=65535, dis=0) produces `4294836225` (correct) not UB. | (2026-06-04) |
 
+| **T-R6-CUDA-VIF-FILTER1D-WIDTH-GUARD-2026-06-04** | `filter1d.cu` 16-bit vertical kernel had a typo in the rd-filter upper-bound guard: `fwidth_rd - fwidth_rd` (always 0) instead of `fwidth - fwidth_rd`. This widened the rd-filter tap window to cover all `fwidth` taps and indexed `vif_filt.filter[scale+1]` 1–4 entries past its allocation, causing OOB reads and wrong VIF scores at scales 0–2 on the CUDA backend. Fix: change the guard to `fi < (fwidth - (fwidth - fwidth_rd) / 2)`, matching the correct 8-bit form at line 183. | no ADR: bug fix per CLAUDE §12 r8 | fix/cuda-vif-filter1d-adm-cm-opprec | `meson test -C build-cuda --suite=fast` + `scripts/dev/cross_backend_diff.py` CPU vs CUDA for VIF convergence within existing tolerance. | (2026-06-04) |
+
+| **T-R6-CUDA-ADM-CM-OPERATOR-PRECEDENCE-2026-06-04** | `adm_cm.cu` lines 373 and 712 computed `x_sq` as `(int64_t)accum * accum + add_shift_sq >> shift_sq`, parsed by C++ as `+ (add_shift_sq >> shift_sq)` = `+ 0` because `>>` binds tighter than `+`. The normalisation shift was silently dropped, leaving `x_sq` carrying an un-normalised squared value (~10^6 vs correct ~0–1) and causing int32 overflow in the cubic accumulator — producing wrong ADM scale-0 and AIM scores on the CUDA backend. Fix: wrap as `((int64_t)accum * accum + add_shift_sq) >> shift_sq`, matching CPU reference macro `I4_ADM_CM_ACCUM_ROUND` at `integer_adm.c:743` and the correct fused kernel at line 259. | no ADR: bug fix per CLAUDE §12 r8 | fix/cuda-vif-filter1d-adm-cm-opprec | `meson test -C build-cuda --suite=fast` + `scripts/dev/cross_backend_diff.py` CPU vs CUDA for ADM convergence within existing tolerance. | (2026-06-04) |
+
 | **T-R5-MEMORY-ORDERING-2026-06-04** | Three HIGH-severity concurrency bugs found and fixed in the r5 audit. (1) `vmaf_ref_fetch_decrement` in `ref.c` and `ref.cpp` used implicit seq_cst; replaced with `memory_order_acq_rel` to match the canonical C11/C++ last-decrementer pattern and make the acquire-on-zero-transition intent explicit. Header `ref.h` gained `using std::memory_order_*` declarations for the C++ branch. (2) `vmaf_feature_collector_destroy` (feature_collector.cpp) unlocked then immediately destroyed the mutex, leaving a window where a concurrent locker acquired a destroyed mutex (UB). Fix: `bool destroyed` field added to `VmafFeatureCollector`; set under lock before the final unlock; all five public entry points test the flag after locking and return `-ENODEV` if set. (3) `vmaf_picture_pool_fetch` (picture_pool.c) unlocked before reading `pool->pictures[idx]`; `vmaf_picture_pool_close` could free the array in that window. Fix: copy the slot to a stack-local before unlocking. | [ADR-1020](adr/1020-r5-memory-ordering.md) | fix/r5-memory-ordering | `meson test -C build --suite=fast` + TSan build optional. | (2026-06-04) |
 
 | **T-Y4M-DST-BUF-READ-SZ-OVERFLOW-2026-06-04** | `core/tools/y4m_input.c::y4m_input_open_impl()` computed `dst_buf_read_sz` for five chroma branches (420/420jpeg/420mpeg2, 420p10, 420p12, 422p10, 422p12) using bare `int * int` arithmetic. A crafted Y4M header with large `pic_w`/`pic_h` values causes signed-integer overflow, underallocating `dst_buf_read_sz` relative to the malloc'd `dst_buf_sz`. The subsequent `fread(_y4m->dst_buf, 1, _y4m->dst_buf_read_sz, _fin)` at line 931 could then request a read of a negative-wrapped (huge) `size_t` number of bytes, overflowing the heap buffer. Fix: add `(size_t)` casts to all five arithmetic expressions, mirroring the pattern already applied to `dst_buf_sz` and all other `dst_buf_read_sz` branches. SEI CERT C INT30-C / INT32-C. | [ADR-1022](adr/1022-y4m-dst-buf-read-sz-overflow.md) | fix/y4m-dst-buf-read-sz-overflow | `python3 -c "import ctypes; print((65536*65536 + 2*(65536//2)*(65536//2)) < 0)"` → `True` before fix (pre-fix int overflow); `(size_t)65536*65536` → `4294967296` (no overflow). | (2026-06-04) |