Skip to content

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

Merged
lusoris merged 3 commits into
masterfrom
fix/cuda-vif-accum-reset-order
Oct 1, 2026
Merged

fix(cuda): reset integer_vif_cuda's accumulators on the picture stream#1750
lusoris merged 3 commits into
masterfrom
fix/cuda-vif-accum-reset-order

Conversation

@lusoris

@lusoris lusoris commented Oct 1, 2026

Copy link
Copy Markdown
Contributor

Summary

Four vmaf CUDA instances on one device return wrong vif scales in every run; one instance never does. integer_vif_cuda cleared its accumulators on a stream the kernels that add into them did not wait for. The clear now runs on the picture stream.

Closes T-UPSTREAM-1305-CUDA-VIF-ACCUM-STREAM-2026-10-01 (found and closed here). Netflix/vmaf#1305 reports the same defect on upstream (integer_vif_cuda.c:458); upstream's patch is the same one-argument change. RC3, correctness (ADR-1421).

Cause

vif_submit_plane() queued cuMemsetD8Async(accum_data, ..., s->str) and then launched the scale 0 kernels, which atomicAdd into accum_data, on the picture stream. A stream orders only its own work. With other instances loading the device, the clear can run after the first atomic adds and erase them.

I checked every other accumulator reset of the CUDA extractors for the same pattern (rg "cuMemsetD8Async" over core/src): integer_motion_sad_cuda.c (picture stream, ADR-0358), integer_adm_cuda.c (every kernel on s->str), float_adm_cuda.c, float_psnr_cuda.c, integer_psnr_cuda.c, integer_cambi_cuda.c and kernel_template.h clear on the stream of their kernels. The HIP twins queue clear and kernels on one stream (read from the source; no HIP multi-instance run).

Output on an RTX 4090

Upstream's harness repro1305 (N instances, one thread each, one shared CUcontext, Netflix 576x324 pair, 48 frames, the vmaf_v0.6.1 features, every score of every frame compared with a flushed single-instance run):

Configuration Runs Wrong values before Wrong values after
1 instance 20 0 0
2 instances, one context or two 20 0 0
4 instances, queried 0 frames behind 5 adm2 28 to 35, motion2 63 to 85, each vif scale 56 to 71 (of 192) 0
4 instances, queried 2 or 3 frames behind 10 adm2 1 to 7, motion2 1 to 20, each vif scale 2 to 14 0

After the fix: 0 wrong values in 105 runs of the seven configurations. A scratch run of four instances on every registered CUDA extractor (vif, adm, motion, motion_v2, float_adm, float_motion, float_psnr, float_ssim, float_vif, float_ms_ssim, float_moment, psnr, psnr_hvs, ssim, ciede, cambi, speed_chroma, speed_temporal, ssimulacra2; 24 frames; the CSV output of each instance compared byte for byte with a single instance's) gives identical output for all.

Tests

  • test_cuda_multi_instance (new, fast + gpu, skips without a device): four threads, each with its own VmafContext and VmafCudaState on one CUcontext, score 48 frames with the features of vmaf_v0.6.1; every vif scale, adm2 and motion2 is compared with == against a single instance. Without the fix it fails on every run (3 of 3: an instance's frame turns non-finite and the call returns -EINVAL); with it 5 of 5 pass.

Type

  • fix — bug fix

Deep-dive deliverables (ADR-0108)

  • Research digest — no digest needed: a one-argument stream fix; the audit of the other resets is in the state row.
  • Decision matrix — no alternatives: only-one-way fix (the reset must share the kernels' stream).
  • AGENTS.md invariant note — core/src/feature/cuda/AGENTS.md, "Accumulator reset runs on the stream of the kernels that add into it".
  • Reproducer / smoke-test command — under "Reproducer" below.
  • CHANGELOG fragment — changelog.d/fixed/cuda-vif-accum-reset-order.md.
  • Rebase note — docs/rebase-notes.md, "integer_vif_cuda resets its accumulators on the picture stream".

Reproducer

On a CUDA host, in a build with -Denable_cuda=true:

flock ~/.cache/vmafx-locks/cuda-4090.lock timeout 120 build/test/test_cuda_multi_instance

It prints pass with the fix and an instance failed without it. Upstream's harness is /home/kilian/.cache/vmafx-upstream-rebase/evidence/c1305/repro1305.c.

Notes

  • Touches core/test/meson.build (one block, next to test_cuda_vif_min_dim); no open PR touches integer_vif_cuda.c.
  • No public API, ABI or FFmpeg patch impact. No Netflix golden assertion changed.

lusoris added 3 commits October 1, 2026 21:12
@lusoris
fix(hip): queue each accumulator clear after the frame's upload so a …
9c46de3 
…later, larger context scores its first frame (#1749)

* fix(hip): queue each accumulator clear after the frame's upload so a later, larger context scores its first frame

float_moment_hip and vif_hip cleared their device accumulators with
hipMemsetAsync ahead of the frame's plane upload. On a gfx1036 a clear
queued there has no effect in the first context of a process that needs
larger planes than the contexts before it, and the kernels add onto the
sums the earlier context left in recycled device memory (ADR-1427).
adm_hip had the same order and failed the run; ADR-1423 moved its clear.

One frame in a 640x360 context, then one in a 3840x2160 context, each
twin in its own process, on master e2954fc, every run:
float_moment_ref1st 130.53 where the CPU has 127.00, vif_hip scale 0 to
2 at 0.6748 / 0.8040 / 0.8749 where the CPU has 0.6934 / 0.8260 /
0.8988. The vmaf tool has one context per process, where fresh device
memory is zero, and never showed it.

float_moment_hip, vif_hip and float_psnr_hip now upload first and queue
the clear directly ahead of their kernels, the order psnr_hip,
float_vif_hip, float_adm_hip and cambi_hip already had. float_psnr_hip's
scores were right (its kernel writes every partial); it is reordered so
that one rule holds for every twin. After the change float_moment_hip
and adm_hip equal the CPU on that frame and vif_hip is within 1.2e-7.

What is measured about the cause: with the clear after the upload the
frame is correct in 23 of 23 runs, also with an allocation, a second
upload, an upload of a new host buffer or an event wait between the
clear and the kernel; a hipStreamSynchronize() after a clear ahead of
the upload cures it; a hipMemset at allocation does not, since on device
memory that call is queued on the null stream without waiting (hipamd
ihipMemset(), rocm-7.2.4); blocking and non-blocking streams both fail.
Why the clear is lost is the runtime's or the driver's and not
established.

Tests: test_hip_first_frame_clear_<twin>, one binary per twin for
fourteen twins, because only the first larger context of a process is
exposed; three fail on master and all pass here, three of three runs
each. test_hip_clear_after_upload_contract.py reads every HIP source and
reports a function that clears and uploads afterwards, helpers included;
it reports the four pre-fix functions on master and has six planted
regressions.

Time per frame is unchanged within the samples: float_moment_hip 1.95
and 1.97 ms at 1080p, 7.22 and 7.14 at 4K; vif_hip 46.86 and 47.84,
207.30 and 202.10.

docs/state.md: T-HIP-FIRST-FRAME-ASYNC-CLEAR-OTHER-TWINS-2026-10-01
closed.
@lusoris
fix(sycl): compute float_vif in the CPU's arithmetic without fp64 so …
7c0fa70 
…the twin is bit-identical (#1736)

* fix(sycl): compute float_vif in the CPU's arithmetic without fp64 so the twin is bit-identical

float_vif_sycl matched the CPU extractor on no frame of the Netflix pair
and was up to 3.8e-5 from it (the SYCL part of
T-GPU-FLOAT-VIF-CPU-ARITHMETIC-2026-10-01, which ADR-1412 opened when it
fixed the CUDA twin). It had the same four differences. Each put back
alone into the corrected twin, on an Arc A380, largest difference on the
Netflix pair and at 3840x2160:

- a table of Gaussian taps the CPU dropped in #758, when float_vif.c
  began to compute its filters with vif_get_filter(): 3.83e-5, 7.5e-6;
- the device log2, where VIF_OPT_FAST_LOG2 makes the CPU's log2f the
  polynomial log2f_approx(): 2.1e-7, 1.2e-7;
- vif_sigma_nsq as a float, where vif_pixel_statistic_s() keeps it in
  double: 1.3e-7, 7.5e-8;
- sub-group and block reductions, where vif_statistic_s() adds a row
  into one float and the rows into another: 5.0e-7, 5.9e-6.

ADR-1422:

- float_vif_sycl.cpp takes each scale's taps from vif_get_filter() and
  passes them to the kernels by value.
- sycl/sycl_float_vif_math.h holds vif_pixel_statistic_s() and
  log2f_approx() operation for operation. A SYCL kernel has no fp64 type
  (ADR-0220), so the reference's two fp64 expressions are evaluated as
  exact fp32 pairs and, within 2^-12 of an fp32 step of a rounding
  boundary (one sample in 1650), by replaying the fp64 add, divide and
  conversion in 64-bit integers. The pair alone rounds wrongly on 85 of
  8.4e9 random quotients; the selected value on none.
- Per scale the filter kernel stores the variances, a statistic kernel
  turns them into the two terms, and a row kernel adds each row in one
  work-item; the host adds the rows in fp32. The statistic has its own
  kernel because it spilled registers inside the filter kernel, and no
  kernel of the twin uses scratch memory (ADR-1395).
- The twin gains the CPU's vif_scale1..3_min_val floors.
- EXACT_TWINS lists float_vif / sycl: the gate cell is an equality.

Measured on the A380 at --precision max, identical frames on every
scale before and after: Netflix 576x324 0/48 and 48/48, checkerboard
1 px 0/3 and 3/3, 10 px 1/3 and 3/3, BBB 3840x2160 0/20 and 200/200;
also identical at 10, 12 and 16 bits, with debug=true and with
non-default options. A 3840x2160 frame takes 23.95 ms, 20.54 before
(T-SYCL-FLOAT-VIF-EXACT-THROUGHPUT-2026-10-01); 576x324: 0.93 and 0.73.
The HIP and Metal twins still hold the tap table.

Tests: test_sycl_float_vif_parity asserts equality over seven cases
(float_vif_twin_parity.h) and fails on the old twin;
test_sycl_float_vif_math compares the header with vif_statistic_s() on
the host and in a device kernel, and with the compiler's fp64 on twelve
operands where the pair alone is wrong;
test_sycl_float_vif_exact_contract.py pins the design.

* docs: regenerate the indexes and the citation map after rebasing
@lusoris
fix(cuda): reset integer_vif_cuda's accumulators on the picture stream (
80c5a03 
#1750)

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

The accumulator reset was queued on the private stream while the scale 0
kernels that add into it run on the picture stream, with nothing ordering
the two. Under contention from other instances on one device a late reset
erased the first adds: four instances on one CUcontext returned wrong vif
scales in every run (Netflix/vmaf#1305). The reset now takes the picture
stream.

test_cuda_multi_instance runs four instances on one context and compares
every vif scale, adm2 and motion2 with a single instance; it fails on every
run without the fix.
@lusoris lusoris mentioned this pull request Oct 1, 2026
Merged
5 of 7 tasks
@lusoris
lusoris force-pushed the fix/cuda-vif-accum-reset-order branch from 30c5786 to 80c5a03 Compare October 1, 2026 19:22
@lusoris
lusoris merged commit 80c5a03 into master Oct 1, 2026
66 of 78 checks passed
@lusoris
lusoris deleted the fix/cuda-vif-accum-reset-order branch October 1, 2026 19:22
This was referenced Oct 1, 2026
Closed
Merged
Merged
@github-actions github-actions Bot added the type:bug Something isn't working label Oct 2, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

type:bug Something isn't working

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

1 participant

@lusoris