fix(cuda): reset integer_vif_cuda's accumulators on the picture stream

CHANGELOG.md

@@ -1118,6 +1118,18 @@
   places or better.
 
 
+- **Several CUDA instances on one device no longer get wrong `vif` scores.**
+  `integer_vif_cuda` cleared its accumulators on its private stream while the
+  scale 0 kernels that add into them ran on the picture stream, with nothing
+  ordering the two. One instance never lost the race; with four instances on
+  one `CUcontext` a late clear erased the first adds, and every run returned
+  wrong `vif` scales (and, through the model, wrong VMAF) for tens of frames.
+  The clear now runs on the picture stream. Measured on an RTX 4090 with four
+  instances on one context, 48 frames of the Netflix 576x324 pair: 28 to 85
+  wrong frames of 192 per feature in every run before, none in 105 runs after
+  (Netflix/vmaf#1305).
+
+
 - **Omit `-march=native` from the reference binary build in `vmaf-dev-mcp`.**
   Intel oneAPI `icx` contracts multiply-accumulate operations in unvectorized CPU
   extractor scalar loops when `-march=native` exposes FMA target capabilities,
@@ -1187,6 +1199,23 @@
   [HIP backend](docs/backends/hip/overview.md#adm_hip-returns-the-cpus-values-bit-for-bit-2026-10-01)).
 
 
+- **`float_moment_hip` and `vif_hip` no longer return a wrong first frame in
+  a later, larger context of one process.** Both cleared their device
+  accumulators ahead of the frame's plane upload. On a gfx1036 such a clear
+  has no effect in the first context of a process that needs larger planes
+  than the contexts before it, so the frame's sums were added onto the sums
+  the earlier context had left in recycled device memory: after a 640x360
+  context, the first frame of a 3840x2160 context had `float_moment_ref1st`
+  130.53 where the CPU has 127.00, and `vif_hip` scale 0 at 0.6748 where the
+  CPU has 0.6934. Every HIP extractor now uploads, then clears, then launches
+  its kernels (`float_psnr_hip` is reordered too; its scores were right).
+  The `vmaf` tool creates one context per process and was not affected;
+  programs that score several clips through the library were. A device test
+  per extractor and a source check over every HIP file hold the order
+  ([ADR-1427](docs/adr/1427-hip-clear-after-upload.md),
+  [HIP backend](docs/backends/hip/overview.md#a-frame-clears-its-accumulators-after-its-upload-adr-1427)).
+
+
 - **`float_motion` on the HIP backend is bit-identical to the CPU, with every
   option.** The CPU extractor adds the absolute differences of a row into one
   `float`, the row sums into a second one, and divides in `float`; those
@@ -1556,6 +1585,30 @@
   `float_ssim_l/c/s` means the same way; linear scores move by less than 6e-8.
 
 
+- **`float_vif_sycl` returns the CPU's scores bit for bit.** The SYCL twin
+  filtered with a table of Gaussian taps the CPU extractor stopped using
+  (it derives them at start-up with `vif_get_filter()`), called the device
+  `log2` where the CPU evaluates a polynomial, took `vif_sigma_nsq` as a
+  `float` where the CPU keeps it in `double`, and reduced per sub-group and
+  per block where the CPU keeps one `float` running sum per row and one over
+  the rows. Measured on an Arc A380 that left no frame of the Netflix
+  576x324 pair identical, up to 3.8e-5 away, and 7.0e-6 at 3840x2160. The
+  twin now takes the CPU's taps, evaluates the CPU's per-pixel statistic
+  without a 64-bit floating-point type (pairs of floats, and the CPU's
+  `double` operations replayed in integers next to a rounding boundary) and
+  adds in the CPU's order
+  ([ADR-1422](docs/adr/1422-sycl-float-vif-cpu-arithmetic.md), after
+  ADR-1412 for CUDA). At `--precision max` every output of every frame is
+  identical on the Netflix pair at 8, 10, 12 and 16 bits, both 1080p
+  checkerboard pairs and 200 frames of BBB 3840x2160, with `debug=true` and
+  with non-default options. `float_vif_sycl` also accepts the CPU's
+  `vif_scale1_min_val`, `vif_scale2_min_val` and `vif_scale3_min_val` now. A
+  3840x2160 frame takes 23.95 ms on the A380, 20.54 ms before, and 100 MB
+  more device memory. The parity gate compares this twin with tolerance 0.
+  Stored `float_vif_sycl` outputs change by up to 3.8e-5. The HIP and Metal
+  twins still agree with the CPU to four decimal places.
+
+
 - **SYCL and HIP `psnr_hvs` return the CPU extractor's scores bit for bit**
   (`T-SYCL-PSNR-HVS-EXACT-SUM-2026-10-01`, `T-HIP-PSNR-HVS-EXACT-SUM-2026-10-01`,
   [ADR-1401](docs/adr/1401-psnr-hvs-sycl-hip-exact-twins.md)). Like the CUDA

changelog.d/fixed/cuda-vif-accum-reset-order.md

@@ -0,0 +1,10 @@
+- **Several CUDA instances on one device no longer get wrong `vif` scores.**
+  `integer_vif_cuda` cleared its accumulators on its private stream while the
+  scale 0 kernels that add into them ran on the picture stream, with nothing
+  ordering the two. One instance never lost the race; with four instances on
+  one `CUcontext` a late clear erased the first adds, and every run returned
+  wrong `vif` scales (and, through the model, wrong VMAF) for tens of frames.
+  The clear now runs on the picture stream. Measured on an RTX 4090 with four
+  instances on one context, 48 frames of the Netflix 576x324 pair: 28 to 85
+  wrong frames of 192 per feature in every run before, none in 105 runs after
+  (Netflix/vmaf#1305).

changelog.d/fixed/hip-first-frame-accumulator-clear.md

@@ -0,0 +1,15 @@
+- **`float_moment_hip` and `vif_hip` no longer return a wrong first frame in
+  a later, larger context of one process.** Both cleared their device
+  accumulators ahead of the frame's plane upload. On a gfx1036 such a clear
+  has no effect in the first context of a process that needs larger planes
+  than the contexts before it, so the frame's sums were added onto the sums
+  the earlier context had left in recycled device memory: after a 640x360
+  context, the first frame of a 3840x2160 context had `float_moment_ref1st`
+  130.53 where the CPU has 127.00, and `vif_hip` scale 0 at 0.6748 where the
+  CPU has 0.6934. Every HIP extractor now uploads, then clears, then launches
+  its kernels (`float_psnr_hip` is reordered too; its scores were right).
+  The `vmaf` tool creates one context per process and was not affected;
+  programs that score several clips through the library were. A device test
+  per extractor and a source check over every HIP file hold the order
+  ([ADR-1427](docs/adr/1427-hip-clear-after-upload.md),
+  [HIP backend](docs/backends/hip/overview.md#a-frame-clears-its-accumulators-after-its-upload-adr-1427)).

changelog.d/fixed/sycl-float-vif-cpu-arithmetic.md

@@ -0,0 +1,22 @@
+- **`float_vif_sycl` returns the CPU's scores bit for bit.** The SYCL twin
+  filtered with a table of Gaussian taps the CPU extractor stopped using
+  (it derives them at start-up with `vif_get_filter()`), called the device
+  `log2` where the CPU evaluates a polynomial, took `vif_sigma_nsq` as a
+  `float` where the CPU keeps it in `double`, and reduced per sub-group and
+  per block where the CPU keeps one `float` running sum per row and one over
+  the rows. Measured on an Arc A380 that left no frame of the Netflix
+  576x324 pair identical, up to 3.8e-5 away, and 7.0e-6 at 3840x2160. The
+  twin now takes the CPU's taps, evaluates the CPU's per-pixel statistic
+  without a 64-bit floating-point type (pairs of floats, and the CPU's
+  `double` operations replayed in integers next to a rounding boundary) and
+  adds in the CPU's order
+  ([ADR-1422](docs/adr/1422-sycl-float-vif-cpu-arithmetic.md), after
+  ADR-1412 for CUDA). At `--precision max` every output of every frame is
+  identical on the Netflix pair at 8, 10, 12 and 16 bits, both 1080p
+  checkerboard pairs and 200 frames of BBB 3840x2160, with `debug=true` and
+  with non-default options. `float_vif_sycl` also accepts the CPU's
+  `vif_scale1_min_val`, `vif_scale2_min_val` and `vif_scale3_min_val` now. A
+  3840x2160 frame takes 23.95 ms on the A380, 20.54 ms before, and 100 MB
+  more device memory. The parity gate compares this twin with tolerance 0.
+  Stored `float_vif_sycl` outputs change by up to 3.8e-5. The HIP and Metal
+  twins still agree with the CPU to four decimal places.

core/src/feature/AGENTS.md

@@ -995,14 +995,16 @@ feature/
   fp32 sum, fp32 divide: NOT the old table's decimals (26 of 34 taps
   differ, centre tap of scale 0 by 14 ulp). GPU twin must take taps from
   `vif_get_filter()` on host, never a literal table: stale table =
-  3.8e-5 on Netflix pair (ADR-1412; CUDA fixed, SYCL / HIP / Metal open,
-  `T-GPU-FLOAT-VIF-CPU-ARITHMETIC-2026-10-01`). Also load-bearing for
+  3.8e-5 on Netflix pair (ADR-1412 CUDA fixed, ADR-1422 SYCL fixed; HIP /
+  Metal open, `T-GPU-FLOAT-VIF-CPU-ARITHMETIC-2026-10-01`). Also load-bearing for
   twins: `VIF_OPT_FAST_LOG2` (`vif_options.h`) makes CPU `log2f` the
   polynomial `log2f_approx()`, no libm; `vif_pixel_statistic_s()` keeps
   `vif_sigma_nsq` in `double`; `vif_statistic_s()` sums row by row in
   fp32. Change any of these -> change
-  `cuda/float_vif/float_vif_device.h` same PR
-  (`test_float_vif_device_math` fails until it follows).
+  `cuda/float_vif/float_vif_device.h` and
+  `sycl/sycl_float_vif_math.h` same PR
+  (`test_float_vif_device_math`, `test_sycl_float_vif_math` fail until
+  they follow).
   [Research-0024](../../../docs/research/0024-vif-upstream-divergence.md)
   = history of the table era.
 

core/src/feature/cuda/AGENTS.md

@@ -1150,6 +1150,22 @@ with a new ADR and measurements, never by reviving ADR-0753 text.
 - Guard: `test_gpu_adm_bright_16bit_parity` in `test_gpu_adm_tiny_frames.c`
   (parity only; device wrap hides the UB itself).
 
+## Accumulator reset runs on the stream of the kernels that add into it (T-UPSTREAM-1305-CUDA-VIF-ACCUM-STREAM-2026-10-01)
+
+- Rule: `cuMemsetD8Async` of an accumulator + the kernels that `atomicAdd` into it
+  = same stream. Streams order only within themselves.
+- `integer_vif_cuda.c`: reset on the picture stream (scale 0 kernels run there),
+  scales 1-3 run on `s->str` after an event recorded behind reset + scale 0,
+  DtoH on `s->str`. Reset on `s->str` raced scale 0 under GPU contention: late
+  reset erased the first adds, vif scales wrong with >= 2 instances on one
+  device (Netflix/vmaf#1305).
+- Audited same rule, no change needed: motion SAD (`integer_motion_sad_cuda.c`),
+  `integer_adm_cuda.c` (all on `s->str`), `float_adm_cuda.c`,
+  `float_psnr_cuda.c`, `integer_psnr_cuda.c`, `integer_cambi_cuda.c`,
+  `kernel_template.h`.
+- Guard: `test_cuda_multi_instance` (4 instances, shared `CUcontext`, scores `==`
+  a single instance; fails without the fix on every run).
+
 ## Device-resident CAMBI and SpEED (ADR-1379, ADR-1380)
 
 - **One readback, one wait per frame.** `cambi_cuda`: 88-byte

core/src/feature/cuda/integer_vif_cuda.c

@@ -731,14 +731,20 @@ static int vif_submit_scales(VifStateCuda *s, CudaFunctions *cu_f, VmafPicture *
 /* vif_submit_plane - one plane's four scales and its accumulator readback.
  *
  * HISS-04: the loop body of submit_fex_cuda, moved whole: the same memset,
- * waits, dispatches, event and readback, in the same order and on the same
- * streams.
+ * waits, dispatches, event and readback, in the same order.
+ *
+ * The accumulator reset runs on the picture stream, the stream of the scale 0
+ * kernels that add into it (Netflix/vmaf#1305). On the private
+ * stream it was ordered against nothing the scale 0 kernels wait for, so
+ * under GPU contention it could land after their first atomic adds and erase
+ * them. Scales 1 to 3 run on the private stream after an event recorded behind
+ * the reset and scale 0, and the readback runs there too.
  */
 static int vif_submit_plane(VifStateCuda *s, CudaFunctions *cu_f, VmafPicture *ref_pic,
                             VmafPicture *dist_pic, unsigned plane)
 {
-    CHECK_CUDA_RETURN(cu_f,
-                      cuMemsetD8Async(s->buf.accum_data->data, 0, sizeof(vif_accums) * 4, s->str));
+    CHECK_CUDA_RETURN(cu_f, cuMemsetD8Async(s->buf.accum_data->data, 0, sizeof(vif_accums) * 4,
+                                            vmaf_cuda_picture_get_stream(ref_pic)));
     CHECK_CUDA_RETURN(cu_f, cuStreamWaitEvent(vmaf_cuda_picture_get_stream(ref_pic),
                                               vmaf_cuda_picture_get_ready_event(dist_pic),
                                               CU_EVENT_WAIT_DEFAULT));

core/src/feature/hip/AGENTS.md

@@ -502,6 +502,29 @@ Rebase-sensitive:
   `_large` (device, `==`, 48 outputs), `test_hip_kernel_source_contract.py`
   (10 planted regressions).
 
+## Frame order: upload, clear, kernels (ADR-1427)
+
+Every HIP twin, every frame: `vmaf_hip_plane_source_acquire*()` /
+`vmaf_hip_picture_upload*()` FIRST, then `hipMemsetAsync` of accumulators,
+then kernels. Rebase-sensitive:
+
+- Clear queued ahead of the upload = lost on gfx1036 in the first context of
+  a process needing larger planes than earlier contexts. Kernels add onto
+  the earlier context's sums in recycled device memory. Was: `adm_hip` (run
+  fails), `float_moment_hip` (moments +3 %), `vif_hip` (scores -0.02).
+  `vmaf` CLI = one context per process, fresh memory zero, never shows.
+- Clear at allocation = no fix. `hipMemset` on device memory is async on the
+  null stream (hipamd `ihipMemset()`, ROCm 7.2.4); `float_moment_hip` stays
+  wrong with it.
+- Rule holds through helpers: a helper that clears counts as a clear at its
+  call site, a helper that uploads as an upload.
+- New twin: add a row to `cases[]` in `core/test/test_hip_first_frame_clear.c`
+  AND its name to `hip_first_frame_twins` in `core/test/meson.build` (one
+  binary per twin; only the first larger context of a process is exposed).
+- Guards: `test_hip_first_frame_clear_<twin>` (device),
+  `test_hip_clear_after_upload_contract.py` (device-free, every `*.c` under
+  `core/src/feature/hip/` and `core/src/hip/`, six planted regressions).
+
 ## adm_hip = CPU bits (ADR-1423, `EXACT_TWINS`)
 
 Scores = `integer_adm.c`'s bits (gfx1036: 21 pairs, 6192 values with
@@ -531,9 +554,8 @@ Scores = `integer_adm.c`'s bits (gfx1036: 21 pairs, 6192 values with
   upload = lost in the first context needing larger planes than earlier
   contexts of the process (recycled device memory -> NaN numerator, run
   fails). `hipMemset` at allocation = no fix: async on the null stream for
-  device memory (hipamd `ihipMemset()`, ROCm 7.2). Same defect in
-  `float_moment_hip`, `vif_hip`:
-  `T-HIP-FIRST-FRAME-ASYNC-CLEAR-OTHER-TWINS-2026-10-01`.
+  device memory (hipamd `ihipMemset()`, ROCm 7.2). Rule for every twin:
+  "Frame order: upload, clear, kernels" above (ADR-1427).
 - Guards: `test_hip_adm_exact` (device, nine contexts in one process, two
   frames each, `==`), `test_hip_adm_exact_contract.py` (device-free, ten
   planted regressions).

core/src/feature/hip/float_moment_hip.c

@@ -162,8 +162,8 @@ static int moment_hip_launch_kernel(MomentStateHip *s, ptrdiff_t row_w, hipStrea
 }
 
 /*
- * Per-frame submit body: zero the four uint64 accumulators, HtoD copies
- * of both luma planes, kernel launch, DtoH readback.
+ * Per-frame submit body: HtoD copies of both luma planes, zero the four
+ * uint64 accumulators, kernel launch, DtoH readback.
  */
 static int moment_hip_launch(MomentStateHip *s, VmafPicture *ref_pic, VmafPicture *dist_pic)
 {
@@ -172,23 +172,24 @@ static int moment_hip_launch(MomentStateHip *s, VmafPicture *ref_pic, VmafPictur
     hipStream_t str = vmaf_hip_stream_of(s->lc.str);
     const size_t sums_bytes = (size_t)MOMENT_HIP_COUNTERS * sizeof(uint64_t);
 
-    /* Zero the four uint64 atomic accumulators before dispatch. */
-    hipError_t hip_rc = hipMemsetAsync(s->rb.device, 0, sums_bytes, str);
-    if (hip_rc != hipSuccess)
-        return moment_hip_rc(hip_rc);
-
     /* Returns once both pictures are read: the caller recycles them when
      * submit() returns (T-HIP-PAGEABLE-UPLOAD-RACE-2026-09-18). */
     int err = vmaf_hip_plane_source_acquire_luma(&s->planes, s->hip_frame, ref_pic, dist_pic,
                                                  s->lc.str, &s->ref_in, &s->dis_in);
+    /* Zero the four uint64 accumulators after the upload, directly ahead of
+     * the kernel that adds into them (ADR-1427): a clear queued ahead of the
+     * upload is lost in the first context of a process that needs larger
+     * planes than the contexts before it. */
+    if (err == 0)
+        err = moment_hip_rc(hipMemsetAsync(s->rb.device, 0, sums_bytes, str));
     if (err == 0)
         err = moment_hip_launch_kernel(s, row_w, str);
     if (err != 0)
         return err;
 
     /* Record submit event, DtoH copy of the four uint64 accumulators,
      * record finished event. */
-    hip_rc = hipEventRecord(vmaf_hip_event_of(s->lc.submit), str);
+    hipError_t hip_rc = hipEventRecord(vmaf_hip_event_of(s->lc.submit), str);
     if (hip_rc == hipSuccess) {
         hip_rc =
             hipMemcpyAsync(s->rb.host_pinned, s->rb.device, sums_bytes, hipMemcpyDeviceToHost, str);

core/src/feature/hip/float_psnr_hip.c

@@ -198,7 +198,7 @@ static int float_psnr_hip_launch_kernel(FloatPsnrStateHip *s, ptrdiff_t plane_pi
 }
 
 /*
- * Per-frame submit body: zero accumulator, HtoD copies, kernel launch,
+ * Per-frame submit body: HtoD copies, zero the partials, kernel launch,
  * submit-event record, DtoH copy.
  */
 static int float_psnr_hip_launch(FloatPsnrStateHip *s, VmafPicture *ref_pic, VmafPicture *dist_pic)
@@ -207,21 +207,23 @@ static int float_psnr_hip_launch(FloatPsnrStateHip *s, VmafPicture *ref_pic, Vma
     const ptrdiff_t plane_pitch = (ptrdiff_t)(s->frame_w * bpp);
     hipStream_t str = vmaf_hip_stream_of(s->lc.str);
 
-    hipError_t hip_rc = hipMemsetAsync(s->rb.device, 0, (size_t)s->wg_count * sizeof(float), str);
-    if (hip_rc != hipSuccess)
-        return hip_err(hip_rc);
-
     /* Returns once both pictures are read: the caller recycles them when
      * submit() returns (T-HIP-PAGEABLE-UPLOAD-RACE-2026-09-18). */
     int err = vmaf_hip_plane_source_acquire_luma(&s->planes, s->hip_frame, ref_pic, dist_pic,
                                                  s->lc.str, &s->ref_in, &s->dis_in);
+    /* The partials are cleared after the upload, like every HIP twin's
+     * buffers (ADR-1427). The kernel writes every partial, so the frame
+     * does not depend on this clear. */
+    if (err == 0) {
+        err = hip_err(hipMemsetAsync(s->rb.device, 0, (size_t)s->wg_count * sizeof(float), str));
+    }
     if (err == 0)
         err = float_psnr_hip_launch_kernel(s, plane_pitch, str);
     if (err != 0)
         return err;
 
     /* Record submit event, DtoH copy of partials, record finished event. */
-    hip_rc = hipEventRecord(vmaf_hip_event_of(s->lc.submit), str);
+    hipError_t hip_rc = hipEventRecord(vmaf_hip_event_of(s->lc.submit), str);
     if (hip_rc == hipSuccess) {
         hip_rc = hipMemcpyAsync(s->rb.host_pinned, s->rb.device,
                                 (size_t)s->wg_count * sizeof(float), hipMemcpyDeviceToHost, str);

core/src/feature/hip/integer_vif_hip.c

@@ -643,23 +643,25 @@ static int submit_fex_hip(VmafFeatureExtractor *fex, VmafPicture *ref_pic, VmafP
 #else
     VifStateHip *s = fex->priv;
 
-    hipError_t rc = hipMemsetAsync(s->accum_dev, 0, sizeof(vif_accums_hip) * 4u, s->str);
-    if (rc != hipSuccess)
-        return vif_hip_err(rc);
-
     /* ADR-0537: the host Y planes in device memory. Returns once both
      * pictures are read: the caller recycles them when submit() returns
      * (T-HIP-PAGEABLE-UPLOAD-RACE-2026-09-18). */
     int err = vmaf_hip_plane_source_acquire_luma(&s->planes, fex->hip_frame, ref_pic, dist_pic,
                                                  vmaf_hip_stream_bits(s->str), &s->ref_in_dev,
                                                  &s->dis_in_dev);
+    /* The accumulators are cleared after the upload, directly ahead of the
+     * kernels that add into them (ADR-1427): a clear queued ahead of the
+     * upload is lost in the first context of a process that needs larger
+     * planes than the contexts before it. */
+    if (err == 0)
+        err = vif_hip_err(hipMemsetAsync(s->accum_dev, 0, sizeof(vif_accums_hip) * 4u, s->str));
     if (err == 0)
         err = vif_hip_launch_scales(s, ref_pic->w[0], ref_pic->h[0], ref_pic->bpc);
     if (err != 0)
         return err;
 
-    rc = hipMemcpyAsync(s->accum_host, s->accum_dev, sizeof(vif_accums_hip) * 4u,
-                        hipMemcpyDeviceToHost, s->str);
+    hipError_t rc = hipMemcpyAsync(s->accum_host, s->accum_dev, sizeof(vif_accums_hip) * 4u,
+                                   hipMemcpyDeviceToHost, s->str);
     if (rc != hipSuccess)
         return vif_hip_err(rc);