Accumulator bounds: CUDA and HIP¶

Appendix of integer accumulator bounds: every integer accumulator, size product and offset of the CUDA and HIP feature twins, read on master 571565a47 (before the fixes the main page lists). Rows marked OVERFLOW or DEPENDS name their state row on the main page.

• Source: master 571565a47.
• Mode: read only. Nothing was edited, built or run on a device. The only commands were grep, sed and cat, plus a few python3 bound computations. The ADM replay is scripts/dev/adm_cm_row_bound.py.
• Scope:
• every file under core/src/feature/cuda/ and core/src/feature/hip/ (130 files);
• the shared core/src/feature/*.h headers those files include for GPU-side integer math;
• core/src/cuda/*.c, core/src/hip/*.c and core/src/cuda/cuda_helper.cuh, for size products only.
• Envelope:
• 16K: W ≤ 15360, H ≤ 8640, N = 132,710,400.
• 8K DCI: 8192x4320.
• Cap: W, H ≤ 32768, N ≤ \(2^{30}\).
• Samples up to 16 bit, 4:4:4.
• Worst-case content: max-difference frames or constructed extremal patterns.

How the audit was done¶

The audit was split into eight feature groups. Each group read every file of its group in full and derived every bound from the code it read. Each group reports its own constants and derivations. Every row that is not SAFE was then re-checked every row that is not SAFE against the source, and spot-checked a set of SAFE rows (listed at the end of this section).

• Feature sections: every feature has CUDA and HIP subsections.
• Integer ADM is the exception. Its HIP kernels are line-for-line ports of the CUDA kernels, so each ADM row cites both files and is tagged "CUDA …; HIP …".
• Rows on shared headers sit in their own subsection of the feature that uses them.

Verdict totals (743 rows)¶

Two verdicts changed from the first draft, both in integer ADM: - Scale-0 CM row total: DEPENDS → OVERFLOW@16K. The envelope covers every W ≤ 15360, and W = 63–64 wraps at default options. - Scale-0 CM frame accumulator: DEPENDS → SAFE. The frame sum cannot wrap for any int64 row value.

Every row that is not SAFE¶

OVERFLOW@16K¶

1. **Integer ADM decouple, scales 1-3: double → int32 out of range (CUDA and HIP).**
2. Location: cuda/integer_adm/adm_decouple_inline.cuh:160, hip/integer_adm/adm_decouple_inline.hip:163.
3. Expression: int32_t rst = (int32_t)(((k * o_val) + 16384) >> 15) * adm_enhn_gain_limit;
4. The (int32_t) cast binds to the shifted product only. The multiplication by the double gain yields a double, which is implicitly converted to int32_t *before* min(rst, t_val) / max(rst, t_val).
5. Bound: |r| ≤ |o| ≤ 1.449e9 at scale 1 and ≤ 7.5e8 at scales 2-3. Those are the DWT band maxima; they agree with adm_csf_fixed_point.h:88-90.
6. The default gain is 100 (adm_options.h:40). With angle_flag set, |r| > 21,474,837 produces a double outside the int32 range, which is UB under C++ [conv.fpint].
7. This does not depend on frame size. It needs only a band value above about 1.5 % of the scale-1 maximum and a dis band within 1° of the ref band, which is ordinary content on high-quality encodes.
8. The CPU (integer_adm_kernels.h:380-385) takes MIN(rst * gain, t) in double first, which is defined.
9. The GPU scores match the CPU only because the hardware conversions saturate: INT32_MAX gives min(INT32_MAX, t) = t. This is consistent with the existing bit-exact parity results, but it is not guaranteed under LLVM (fptosi out of range is poison).
10. Fix: rst = (int32_t)fmin((double)r * gain, (double)t) and the mirrored fmax, as the CPU does.
11. **Integer ADM scale-0 contrast-masking row total, int64 (CUDA, HIP, and the CPU reference).**
12. Locations:
    • CUDA adm_cm.cu:561, 894 (accum_row), :515 (warp_reduce), :519.
    • HIP adm_cm.hip:440, 719, :457, :465.
    • CPU integer_adm_kernels.h:1017-1019, 1079-1099.
13. Term: ((x² + 2^28) >> 29) · x >> (ceil(log2 Wb) − 4), where Wb = W/2. The shift uses the band width (integer_adm.c:711 passes w2; integer_adm_kernels.h:989).
14. Worst case: ref == dis, so a = 0 and thr = 0, and x = |band| · weight. With 4 pixel rows (0, 0, max, 0), the exact column-sign maximum comes from the DP in scripts/dev/adm_cm_row_bound.py. It was re-run for this page.

15. Row maximum as a fraction of INT64_MAX at the default Watson weights:

    W	row max / INT64_MAX
    48	0.807
    56	0.878
    **64**	**1.021 (wraps)**
    128	0.973
    1080p (1920)	0.857
    8K DCI	0.912
    15360	0.855
    16384	0.912
    cap	0.912

16. The 8-bit maximum at W = 64 is 1.009 and the 10-bit maximum is 1.018.

17. W ≥ 17 is accepted (ADM_MIN_FRAME_DIM, adm_csf_fixed_point.h:277), so W = 63–64 runs.
18. With non-default CSF weights the row also overflows at 16K and at the cap: h/v ≥ 38,406 or d ≥ 60,965 at 16K, h/v ≥ 37,588 at the cap. Those weights are reachable through adm_csf_scale / adm_csf_diag_scale / the viewing-geometry options.
19. The bound is not monotonic in W, because the normalising shift is ceil(log2 Wb).
20. The CPU has the same signed int64 sum, so all backends share the defect: signed-overflow UB on the CPU, and a two's-complement wrap on the GPUs.

OVERFLOW@CAP-ONLY¶

1. **HIP PSNR-HVS prefix scan silently stops at 32,768 chunks (hip/integer_psnr_hvs/psnr_hvs_score.hip:462-471).**
2. Code: limit = num_chunks < 32768u ? num_chunks : 32768u.
3. num_chunks = ceil(total_blocks / 256), and total_blocks = Σ planes ((w−8)/7+1)·((h−8)/7+1) (integer_psnr_hvs_hip.c:389-393).
4. 16K 4:4:4: 3 · 2194 · 1234 = 8,122,188 blocks, so 31,728 chunks. That is under the cap, with a 3.2 % margin.
5. The cap is passed above 8,388,608 blocks, for example:
   • 16384x8640 4:4:4: 8,662,680 blocks.
   • luma-only frames above about 20.3K x 20.3K.
   • cap 4:4:4: 256,779 chunks.
6. Above that, chunk_offsets[c ≥ 32768] are never written. d_scratch is a hipMalloc buffer that is never cleared, so hvs_compact_hip writes packed_terms + chunk_offsets[chunk] + intra from stale or garbage offsets. The results are out-of-bounds device writes, a short total_terms, and wrong scores.
7. The CUDA twin (psnr_hvs_score.cu:443) scans every chunk.
8. The uint32_t sum itself is safe: ≤ 4,207,058,112 at the cap.
9. **CUDA SpEED covariance divisor rounded to fp32 (cuda/speed/speed_score.cu:1219).**
10. Code: static_cast<float>(g.sub_w * g.sub_h).
11. This is not a wrap: the uint32_t product is at most 67,010,596. The problem is the fp32 exact-integer range, \(2^{24}\).
12. The count is exact at 16K for every option value: the largest is 3836 · 2156 = 8,270,416 at speed_prescale 4.
13. Above 16K with a prescale above about 2, the count can be inexact. For example, W = H = 32740 at prescale 4 gives sub 8181, and 8181² = 66,928,761 is odd and above \(2^{24}\).
14. The CPU (speed.c:847) divides the double sum by the exact double count. The twin is then no longer bit-exact.
15. The mean divisor (:1183) rounds the count the same way the CPU's compute_mean (speed.c:775) does, so it is SAFE.
16. **HIP SpEED covariance divisor (hip/speed/speed_hip_device.h:624).** Same defect as row 4.

DEPENDS¶

1. **CUDA integer PSNR clip-level apsnr_sse[p] += sse (cuda/integer_psnr_cuda.c:394), uint64_t.**
2. One plane's frame SSE ≤ N · 65535² per frame.

3. Frames to wrap at 16-bit max-diff:

   resolution	frames to wrap
   1080p	2,072
   8K DCI	122
   16K	33
   cap	5

4. These come from \(2^{64}\) / (N · 4,294,836,225) = 2071.4, 121.4, 32.4 and 4.0.

5. At 12 bit the counts are 530,502 / 31,085 / 8,290 / 1,025.
6. At 10 bit they are 8.5M / 498,076 / 132,821 / 16,417.
7. At 8 bit they are 136.8M / 8.0M / 2.1M / 264,205.
8. The wrap is silent, and apsnr comes out too high. CPU integer_psnr.c uses the same uint64_t.
9. **HIP integer PSNR apsnr_sse[p] += sse (hip/integer_psnr_hip.c:476).** Same as row 6.
10. **HIP float PSNR 16bpc kernel at bpc 10/12 (hip/float_psnr/float_psnr_score.hip:151,161,166), uint32_t block sum of 256 float squares.**
11. In-range samples give 256 · 4095² = 4,292,870,400, which is 2,096,895 below UINT32_MAX.
12. uint16 codewords above \(2^{\mathrm{bpc}} - 1\) can wrap the sum. Nothing in picture.c / libvmaf.c checks samples against bpc.
13. At 16 bit the kernel splits each square into 16-bit halves, so that case is safe.
14. **HIP integer VIF horizontal accum_mu1/mu2 += coeff * tmp.mu1[...], uint32_t (hip/integer_vif/vif_statistics.hip:499-500, 530-534).**
15. The HIP vertical pass stores tmp.mu1 without the (uint16_t) truncation that the CPU (integer_vif.c:450-451) and CUDA (filter1d.cu:603-604) apply.
16. For bpc 9-15 with samples above \(2^{\mathrm{bpc}} - 1\), the stored value reaches 8,388,480, and the 17-tap sum reaches 5.5e11. That wraps, and HIP then differs from the CPU.
17. In-range samples are SAFE: ≤ 4,294,901,760.
18. **HIP integer VIF vif_downsample_store accum_ref_rd (hip/integer_vif/vif_statistics.hip:506-507, 329-336).** Same cause as row 9: ref_convol is not truncated.
19. **CUDA motion batch-boundary frame counter (cuda/integer_motion_cuda.c:830, :533-534, :545).**
    • Code: last_batch_boundary = (int)index.
    • Index \(2^{31} - 1\) ≡ 7 (mod 8) is a boundary, so last_batch_boundary + 1 is signed overflow (UB) at the flush.
    • Beyond \(2^{31}\), (int)index is an implementation-defined narrowing.
    • It needs ≥ 2,147,483,647 frames: 414 days at 60 fps.
20. **Integer ADM flt = (4369·|i16| + 2048) >> 12 narrowed to int16 (CUDA and HIP).**
    • Locations: CUDA adm_csf.cu:171, adm_cm.cu:844-847; HIP adm_csf.hip:193, adm_cm.hip:658.
    • It wraps negative once |i16| ≥ 30,720, which needs an h/v CSF weight ≥ 43,900. The default is 36,453, which gives at most 27,212.
    • The CPU (integer_adm_kernels.h:488-489) narrows the same way.
    • It depends on options, not on frame size.
21. **Integer ADM scale-0 x_sq = (int32_t)((x*x + 2^28) >> 29) (CUDA adm_cm.cu:218, HIP adm_cm.hip:150).**
    • SAFE while thr ≥ 0.
    • It wraps only when row 12's negative int16 makes thr negative, so it also depends on options.

Asides (not integer overflow; found during the reading, worth a ticket)¶

• A1 (CUDA float_motion, out-of-bounds read): verified.
• fm_mirror() (cuda/float_motion/float_motion_score.cu:43-50) reflects once without a clamp.
• fm_load_tile() (:85-90) loads the full 20x20 tile for every block. For W or H in {3..9, 17}, an index goes negative (down to −13) and the load reads before the plane.
• The HIP twin and the CUDA motion SAD kernels clamp through vmaf_*_tile_index().
• A2 (CUDA integer VIF, full-mask shuffle inside a divergent branch).
• vif_hori_flush_accums() → warp_reduce() runs __shfl_down_sync(0xffffffff, …) under if (y < h && x_start < w) (cuda/integer_vif/filter1d.cu, the kernel body around :513-535).
• Lanes past the plane edge do not take part, which is undefined for a full mask.
• The code shape is verified. The effect on hardware is not.
• **A3 (core/src/cuda/cuda_helper.cuh:129-130).** (x >> 32) << 32 on a negative long long is UB before C++20.
• A4 (stale comments). hip/float_psnr/float_psnr_score.hip:107-111,139-141 describe a partials[2*block] layout, but the code writes one u64 per block.
• A5 (outside scope, CPU).
• third_party/xiph/psnr_hvs.c:311-317: the int index reaches exactly INT_MAX for a >8-bit 32768² plane.
• integer_motion.c:193/243: the uint32_t row_sad wraps for out-of-range samples (bpc 9, W ≥ 513).
• sycl/speed_sycl_pipeline.cpp:808: the same fp32 covariance divisor as rows 4-5.
• A6 (memory footprint only, no wrap).
• HIP integer ADM tmp_accum takes 24 B/pixel: 3.2 GB at 16K and 25.8 GB at the cap.
• CUDA vif_cuda allocates 7 planes that no kernel reads: 3.2e9 B at 16K.
• The integer SSIM twins need 7 · \(2^{33}\) B at the cap.
• SpEED at prescale 4 needs about 34 GB at 16K.
• Every one of these is a size_t product, so allocation fails cleanly.
• Common root of rows 8-10. No libvmaf entry point checks that samples are ≤ \(2^{\mathrm{bpc}}\) − 1. The CPU truncates or uses wide types in places where HIP does not.

Verification log¶

• Row 1:
• Read adm_decouple_inline.cuh:125-170 and adm_decouple_inline.hip:150-170.
• Read the CPU integer_adm_kernels.h:355-385.
• Confirmed the gain default 100.0 at adm_options.h:40.
• Row 2:
• Re-ran scripts/dev/adm_cm_row_bound.py; the output is the table in row 2.
• Confirmed the term and shifts: adm_cm.cu:216-220, :500-561; integer_adm_kernels.h:989-1003.
• Confirmed the band width w2 (integer_adm.c:711) and the minimum frame size 17 (adm_csf_fixed_point.h:277).
• For the frame accumulator: region rows ≤ 0.875 · \(2^{\lceil \log_2 H_b \rceil}\), with the maximum at Hb = 16 or 32. So |Σ (row >> s)| < 0.875 · \(2^{63}\) for any int64 row value.
• Row 3:
• Read psnr_hvs_score.hip:430-526 and integer_psnr_hvs_hip.c:169-189, :375-394.
• Recomputed the 16K block and chunk counts.
• Rows 4-5:
• Read speed_score.cu:1175-1222 and speed_hip_device.h:545-628.
• Read the CPU speed.c:766-776, :835-850, :1285-1317, and the option range speed.c:1518-1523 (speed_prescale 0.1..4.0).
• Recomputed sub at 16K at prescale 4: 3836 x 2156.
• Found that 32768 at prescale 4 gives 8186² = 67,010,596, which is a multiple of 4 and so exact in fp32. Only some sizes above \(2^{24}\) are inexact, for example 32740.
• Rows 6-8:
• Read integer_psnr_cuda.c:380-400, the apsnr_sse sites in both twins, and float_psnr_score.hip:1-181.
• Recomputed the wrap frames.
• Rows 9-10: read vif_statistics.hip:455-540 against filter1d.cu:585-612.
• Row 11: read integer_motion_cuda.c:525-550, :822-831, and grepped both trees for other (int)index narrowings (none).
• Spot-checked SAFE rows, each read in the source:
• CUDA psnr_score.cu (uint64 SSE, int64 diff, one atomic per block).
• CUDA motion_v2_score.cu: |h| ≤ 65535; SAD ≤ N · 65535 = \(2^{46}\) at the cap.
• CUDA moment_score.cu: u64 sums; ref2 ≤ \(2^{62}\) at the cap.
• SSIMULACRA 2 CUDA unsigned index 2u * pixels + i: ≤ 3 · \(2^{30} - 1\) < \(2^{32}\).
• HIP SpEED (uint32_t)plane_bytes: \(2^{31}\) at the cap, which fits.
• CAMBI histogram cells: ≤ 65 · 65.
• The runtime picture allocators: all size_t.

Section index¶

G1: integer ADM, CUDA and HIP twins (integer-overflow audit)¶

Repo: master at 571565a47. All paths are under core/src/feature/. Read-only. Nothing was built or run on a device. The bounds come from the code, the composite-filter bounds per DWT scale of core/src/feature/adm_csf_fixed_point.h (held by test_integer_adm_cm_budget), and scripts/dev/adm_cm_row_bound.py (the exact maximum of a scale-0 CM row over every column sign pattern, found by dynamic programming over the overlapping 4-tap windows).

Constants used by every row:

• DWT taps. lo = {15826, 27411, 7345, -4240} and hi = {-4240, -7345, 27411, -15826}. Each has an absolute sum of 54822. The lo taps sum to 46342 (integer_adm.h:187-191; the GPU copies are in adm_fixed_parameters() and adm_hip_fixed_params()).
• Band maxima (half the absolute sum of the composite filter; they agree with adm_csf_fixed_point.h:88-90):
• scale 0: every band ≤ 22,929.8.
• scale 1: band_a ≤ 1,491,390,732 and h/v/d ≤ 1,448,979,042.
• scale 2: ≤ 751,510,749.
• scale 3: ≤ 748,642,077.
• Every maximum is below INT32_MAX. They are per sample and do not depend on frame size.
• CSF weight budgets (ADR-1472, adm_csf_fixed_point.h:118-127):
• scale 0 h/v: < 46,603.4.
• scale 0 d: < 65,536 (the storage limit).
• scale 1: < 279,958,309; scale 2: < 539,893,111; scale 3: < 546,406,567.
• Defaults (Watson97, 3H, 1080): 36,453 for h/v and 49,417 for d (tabulated).
• The options adm_csf_mode and adm_csf_scale / adm_csf_diag_scale (continuous doubles, 0..50, applied multiplicatively in barten_csf_tools.h:159-160) can place a weight anywhere below its budget. So can a non-default adm_norm_view_dist or adm_ref_display_height.
• Region. cols = Wb - 2*int(0.1*Wb - 0.5), where Wb is the band width (W/2 at scale 0).
• 16K: Wb = 7680, cols = 6146.
• 8K DCI: Wb = 4096, cols = 3278.
• cap: Wb = 16384, cols = 13110.
• Scale-0 rows at 16K: Hb = 4320, so 3458 rows.

integer_adm: DWT (scale 0 fused kernel, scales 1-3 combined kernels)¶

integer_adm: decouple and CSF (inline in the CSF / CM kernels)¶

integer_adm: CSF denominator (adm_csf_den)¶

integer_adm: contrast masking (DLM and AIM, scale 0 and scales 1-3)¶

integer_adm: host buffers, strides and launch geometry¶

shared headers (GPU-relevant parts)¶

Coverage¶

Row counts by verdict, 42 rows in total:

G2 integer-overflow audit: float ADM, float VIF, integer VIF (CUDA + HIP)¶

Repo: master at 571565a47. Read-only. Paths below are relative to core/src/feature/.

Envelope: 16K N = 132,710,400 (\(2^{26.98}\)), cap N = \(2^{30}\), W,H <= 32768. Samples <= 65535 (bpc 16) and in range (<= \(2^{\mathrm{bpc}} - 1\)) unless a row says otherwise.

Basis for integer VIF, checked against the CPU reference (integer_vif.h:39-46, integer_vif.c:175-228, 254-341, 406-456):

• Every row of vif_filter1d_table sums to exactly 65536. Widths are {17, 9, 5, 3}, and the largest tap is 43,728 (scale 3 centre). So any tap sum is at most 65536 * (max input), whatever W and H are.
• Scale-0 vertical shifts at bpc b: shift_VP = b, shift_VP_sq = 2(b-8). Scales 1-3 shift by 16. The horizontal shift is always 16. The GPU twins copy these values: integer_vif_cuda.c:572-589 and integer_vif_hip.c:334-341.
• Inputs to scales 1-3 are the uint16 decimated planes: at most 65,535 for 16-bit input, at most 65,280 for 8-bit input.
• The int64 accumulators hold the sums of one frame and one scale. They are cleared every frame (integer_vif_cuda.c:811, integer_vif_hip.c:689), so the frame count does not matter.

integer VIF: CUDA (vif_cuda)¶

Asides (CUDA integer VIF, no overflow, no verdict):

• Possible undefined-behaviour shuffle, unverified. filter1d.cu:513/535: vif_hori_flush_accums() (warp_reduce, __shfl_down_sync(0xffffffff, …)) runs only inside if (y < h && x_start < w). When the scale width is not a multiple of 64, lanes of the edge warp skip the shuffle that the full mask names. The CUDA guide calls this undefined, and the values read from those lanes are undefined. This is not an overflow. test_integer_vif_cpu_cuda_parity (256x144, so a 32-pixel width at scale 3 with 16 active lanes) passes places=4, so current hardware returns benign values. It is not checked on a device here.
• Idle horizontal blocks. filter1d_16_grid (integer_vif_cuda.c:620) sets grid_hori_x = ceil(w/128), but a block covers 256 pixels (vpt = 2), so about half the x-blocks do nothing. There is no wrap.
• Unused buffer memory. data_sz also carves 2 stride_16 planes and 5 stride_32 planes (mu1, mu2, mu1_32, mu2_32, ref_sq, dis_sq, ref_dis) that no kernel reads. That is about 3.18e9 B at 16K and 2.58e10 B at cap of device memory, which matters for memory, not for overflow.

integer VIF: HIP (vif_hip)¶

integer VIF: shared headers¶

float VIF: CUDA (float_vif_cuda) and HIP (float_vif_hip)¶

The accumulators are fp32 throughout: fvif_row_sum and fvif_sum_rows, as in the CPU. They are never converted to an integer, so they are out of scope. Only index, size and float-to-int math is listed. The only conversion is (float)((int32_t)exponent - 127).

float ADM: CUDA (float_adm_cuda) and HIP (float_adm_hip)¶

The accumulators are fp32: the per-row sum fadm_row_sum and the frame fold fadm_fold_rows. They are never converted to an integer. The device code has no float-to-int conversion; fadm_abs is a bit mask. Integer math is index and size arithmetic only. Band planes at scale 0 are half_w0 = (W+1)/2 by half_h0 = (H+1)/2, with buf_stride = (half_w0 + 3) & ~3.

Coverage¶

All 27 files in group G2 were read in full or every grep hit was read in context. Rows per file (111 rows in total: 109 SAFE, 2 DEPENDS, 0 OVERFLOW@16K, 0 OVERFLOW@CAP-ONLY):

• cuda/float_adm_cuda.c: 8
• cuda/float_adm_cuda.h: none (declares the PTX symbol only)
• cuda/float_adm/float_adm_device.h: none (macros and typedefs only)
• cuda/float_adm/float_adm_score.cu: 7
• cuda/float_vif_cuda.c: 3
• cuda/float_vif_cuda.h: none (declares the PTX symbol only)
• cuda/float_vif/float_vif_device.h: none (macros and typedefs only)
• cuda/float_vif/float_vif_score.cu: 6
• cuda/integer_vif_cuda.c: 7
• cuda/integer_vif_cuda.h: none (struct types only; vif_accums is 7 int64_t, covered under the .cuh rows)
• cuda/integer_vif/filter1d.cu: 20 (includes the warp_reduce/atomicAdd_int64 rows)
• cuda/integer_vif/vif_statistics.cuh: 12
• hip/float_adm_hip.c: 4
• hip/float_adm_hip.h: none
• hip/float_adm/float_adm_hip_math.h: none (macros only)
• hip/float_adm/float_adm_score.hip: 5
• hip/float_vif_hip.c: 2
• hip/float_vif_hip.h: none
• hip/float_vif/float_vif_score.hip: 2
• hip/integer_vif_hip.c: 5
• hip/integer_vif_hip.h: none (vif_accums_hip is 7 int64_t, covered under the .hip rows)
• hip/integer_vif/vif_statistics.hip: 22 (2 DEPENDS)
• float_adm_gpu_common.h: 3
• float_vif_gpu_common.h: 3
• integer_vif_sv_sq.h: 1
• vif_log2_table.h: 1
• adm_float_reference.h: integer parts are adm_border_s and adm_pool_bands_s (whose body, get_noise_constant, uses w * h). Both are covered by the float ADM rows; 0 rows of its own.

Auxiliary files read to bound the terms:

• integer_vif.c and integer_vif.h: the CPU reference shifts, casts, filter table and log2_32/64.
• cuda/cuda_helper.cuh:125-139: warp_reduce and atomicAdd_int64.
• cuda/cuda_tile_index.h: the reflect/clamp index helpers.
• mem.h:26-31: ALIGN_CEIL, MAX_ALIGN = 32.
• cuda/common.h:54-55: warp size 32, cache line 128.
• adm_tools.c:66-99: get_noise_constant and adm_border_s.
• adm_options.h:24: ADM_BORDER_FACTOR = 0.1.
• src/meson.build:1553: CUDA builds with --std c++20.
• test_integer_vif_cpu_cuda_parity.c: the fixture is 256x144.

G3a: motion family, CUDA + HIP twins (integer-overflow audit)¶

Repo: master at 571565a47. Read-only. Paths are relative to core/src/feature/ unless they start with core/.

Envelope: 16K N = 132,710,400 (W 15360, H 8640); 8K DCI N = 35,389,440; 1080p N = 2,073,600; cap W,H <= 32768, N <= \(2^{30}\) (core/src/picture.c:46); bpc 8..16 (core/src/picture.c:37-38).

Shared derivation: per-pixel motion SAD term¶

Taps {3571, 16004, 26386, 16004, 3571}, sum 65536 = \(2^{16}\) (cuda/integer_motion_v2/motion_v2_score.cu:39, hip/integer_motion_v2/motion_v2_score.hip:67, identical to CPU integer_motion.h:26). CPU reference integer_motion.c:164-257 (pipeline_8 / pipeline_16): vertical shift bpc, round \(2^{\mathrm{bpc}-1}\); horizontal shift 16, round \(2^{15}\). Twins copy it exactly.

• diff d = prev - cur: |d| <= \(2^{\mathrm{bpc}} - 1\) for in-range samples, <= 65535 for any uint16 sample.
• vertical v = (sum f_k d_k + 2^(bpc-1)) >> bpc: |sum| <= 65536 * |d|max. In range: |v| <= floor(65536 (\(2^{\mathrm{bpc}} - 1\)) / \(2^{\mathrm{bpc}}\) + 0.5) <= 65535 (bpc 16: 65535; bpc 8: 65280). Out-of-range samples (uint16 value > \(2^{\mathrm{bpc}} - 1\), worst bpc = 9): |v| <= (65536 * 65535 + 256) >> 9 = 8,388,480 (~\(2^{23}\)).
• horizontal h = (sum f_k v_k + 2^15) >> 16: |h| <= |v|max. In range |h| <= 65535 (reached by a uniform max-difference frame: ref 65535 / dis 0). Out-of-range worst |h| <= 8,388,480.
• Frame SAD = sum over N pixels of |h|. In range: 16K 8,697,176,064,000 (\(2^{42.98}\)); 8K 2,319,246,950,400 (\(2^{41.08}\)); 1080p 135,893,376,000 (\(2^{36.98}\)); cap 70,367,670,435,840 (\(2^{46.0}\)). Out-of-range worst (bpc 9): 16K 1.113e15 (\(2^{49.98}\)); cap 9,007,061,815,787,520 (< \(2^{53}\) = 9,007,199,254,740,992).
• The SAD accumulator is zeroed before every frame (CUDA cuda/integer_motion_sad_cuda.c:177, HIP hip/integer_motion_sad_hip.c:172), so there is no clip-level (cross-frame) integer accumulator in G3a; frames-to-overflow: not applicable.

motion SAD kernel (motion_cuda + motion_v2_cuda) - CUDA¶

motion_v2_cuda host - CUDA¶

motion_cuda host - CUDA¶

float_motion_cuda - CUDA¶

Float blur and fp32 row sums are not integer accumulators and are never converted to an integer (skipped per brief).

motion SAD kernel (motion_hip + motion_v2_hip) - HIP¶

motion_hip / motion_v2_hip host - HIP¶

float_motion_hip - HIP¶

Float blur, fp32 transposed differences and fp32 row sums are not integer accumulators (skipped per brief).

Shared headers¶

Aside A: CUDA float_motion tile index can go negative (out-of-bounds read, not an integer overflow)¶

cuda/float_motion/float_motion_score.cu:43-50 (fm_mirror) reflects once and does not clamp, and fm_load_tile (:83-90) loads a fixed 20x20 tile for every block, padding threads included. The largest tile index on an axis of extent sup is 16 * (ceil(sup / 16) - 1) + 17; it maps below zero when that exceeds 2 * (sup - 1), i.e. for sup in {3..9, 17}. Width or height 3 gives index -13 (13 rows before the ref_in allocation when it is the row axis), 17 gives -1. Those samples feed no valid output (the consumed taps stay in range from 3x3 up), but the load reads outside the buffer. The HIP twin clamps (hip/float_motion/float_motion_score.hip:101-104, ADR-1381), and the CUDA motion SAD kernel clamps (cuda/cuda_tile_index.h). core/test/test_cuda_motion_tiny_frames.c:67 covers 3x3 and 17x17 for motion_cuda and motion_v2_cuda only, not float_motion_cuda. The init guard only refuses below 3x3 (cuda/float_motion_cuda.c:262).

Aside B: CPU reference row accumulator (outside G3a, for the CPU group)¶

integer_motion.c:193 and :243 hold a row's SAD in uint32_t row_sad. In range it is <= 32768 * 65535 = 2,147,450,880 (< \(2^{32}\)), so it is safe up to the cap. If a 16-bit container carries samples above \(2^{\mathrm{bpc}} - 1\) (bpc 9..15; sample range against bpc was not checked in this audit), abs(h) reaches 8,388,480 at bpc 9 and the row sum wraps from W >= 513 (bpc 10: abs(h) <= 4,194,240, W >= 1025). The GPU twins add in 64 bits and would then differ from the CPU.

Coverage¶

All paths under core/src/feature/. Every file was read in full, or grepped with the brief's pattern list and every hit read in context (for the 600-900-line host files).

G3b: PSNR (integer + float), moment, CIEDE, tile-index headers (CUDA + HIP)¶

Repo: master @ 571565a47. Paths are relative to core/src/feature/. Envelope: 16K N = 132,710,400 (\(2^{26.98}\)); cap N = \(2^{30}\). Max 16-bit term \(65535^{2}\) = 4,294,836,225. Its fp32 square is 4,294,836,224 (the twins' float-square terms). Both values are below \(2^{32}\). Row W <= 32768 = \(2^{15}\).

float_psnr¶

CUDA (cuda/float_psnr/float_psnr_score.cu, cuda/float_psnr_cuda.c, float_psnr_rows.h)¶

Grid aside: gridDim.y = H (FPSNR_BY = 1), <= 32768, below the 65535 y-dimension limit.

HIP (hip/float_psnr/float_psnr_score.hip, hip/float_psnr_hip.c)¶

Note: the kernel doc comments (float_psnr_score.hip:107-111, 139-141) describe a partials[2 * block] / [2 * block + 1] layout. The code writes one u64 per block (partials[block_idx]), and the host reads one u64 per block. The comments are out of date. This does not cause an overflow.

psnr (integer)¶

CUDA (cuda/integer_psnr/psnr_score.cu, cuda/integer_psnr_cuda.c, psnr_score.h)¶

vmaf_psnr_aggregate (psnr_score.h:89-97) and the mse at :397 convert uint64_t to double. They have no integer math.

HIP (hip/integer_psnr/psnr_score.hip, hip/integer_psnr_hip.c)¶

float_moment¶

CUDA (cuda/integer_moment/moment_score.cu, cuda/integer_moment_cuda.c)¶

(double)sums_host[k] at integer_moment_cuda.c:329-332 converts uint64 to double. It is not integer math. Above \(2^{53}\), exactness is a parity question that ADR-1497 handles; it does not cause a wrap.

HIP (hip/float_moment/moment_score.hip, hip/float_moment_hip.c)¶

Shared (float_moment_sum.h, float_moment_sum_gpu.h; ordered_sum.h read for the units helpers). Used by the CUDA and HIP twins¶

ciede¶

CUDA (cuda/integer_ciede/ciede_score.cu, ciede_device.h, cuda/integer_ciede_cuda.c)¶

The per-pixel float plane has no integer accumulator. Only index and size math is listed.

HIP (hip/integer_ciede/ciede_score.hip, hip/ciede_hip.c)¶

tile-index headers¶

Coverage¶

Verdict totals: SAFE 80, DEPENDS 3, OVERFLOW@16K 0, OVERFLOW@CAP-ONLY 0.

G4a: integer-overflow audit of SSIM (integer and float) and MS-SSIM, CUDA and HIP twins¶

Repo: master at 571565a47. This audit was read-only. Every path below is relative to core/src/feature/. Envelope: 16K N = 132,710,400; cap N = \(2^{30}\) (W, H <= 32768); 16-bit samples, 4:4:4.

Derived term maxima (shared by every integer-SSIM row)¶

• Integer Gaussian gaussian_filter_init(1.5, 5) (integer_ssim.c:60-96) has kernel_len = floor(1.5sqrt(-2 ln(sqrt(pi/2)1.5/256))) = floor(4.70) = 4, so it has 9 taps [2,9,28,55,68,55,28,9,2] with a sum of exactly 256. The CUDA kernel (ISSIM_KERNEL, integer_ssim_score.cu:70) and the HIP kernel (integer_ssim_score.hip:84) copy it.
• Horizontal moments (a window of at most 9 taps, at most 65535 per sample):
• mux and muy: at most 256*65535 = 16,776,960 (\(2^{24}\)).
• x2, xy and y2: at most 256*\(65535^{2}\) = 1,099,478,073,600 (< \(2^{40}\)).
• w: at most 256.
• Vertical moments (9 taps over the horizontal moments):
• mux and muy: at most 65536*65535 = 4,294,901,760 (< \(2^{32}\)).
• x2, xy and y2: at most 65536*\(65535^{2}\) = 281,466,386,841,600 (< \(2^{48}\)).
• w: at most 65,536 (\(2^{16}\)).
• The bounds above are local to one window and do not depend on the frame size. Frame-level weight sums are Σ m.w ≤ 65536*N: 8,697,308,774,400 (\(2^{42.98}\)) at 16K and \(2^{46}\) at the cap.
• The float_ssim decimation is a fixed-point sum in units of \(2^{-52}\) over a scale x scale window, with tap = fl(1/\(\mathrm{scale}^{2}\)) and sample ≤ 65535/256 = 255.996. A window therefore sums to at most about 256, which is ≤ 255.996*\(2^{52}\) ≈ \(2^{59.99998}\) < \(2^{60}\). This bound holds for every scale. The cap of 128 on the scale (SSIM_MAX_EXACT_SCALE / VMAF_HIP_SSIM_MAX_EXACT_SCALE) is there for exactness, not for overflow.
• No accumulator in G4a spans more than one frame (clip level). Every sum is a local in collect() or a per-frame device buffer, so no frames-to-overflow figure applies.

integer SSIM (ssim)¶

CPU reference (integer_ssim.c), used only to derive the term maxima¶

CUDA (cuda/integer_ssim/integer_ssim_score.cu, cuda/ssim_cuda.c)¶

HIP (hip/integer_ssim/integer_ssim_score.hip, hip/integer_ssim_hip.c)¶

float SSIM (float_ssim)¶

CUDA (cuda/integer_ssim/ssim_score.cu, cuda/integer_ssim_cuda.c)¶

HIP (hip/float_ssim/ssim_decimate.h, hip/float_ssim/ssim_score.hip, hip/float_ssim_hip.c)¶

float MS-SSIM (float_ms_ssim)¶

CUDA (cuda/integer_ms_ssim/ms_ssim_score.cu, cuda/integer_ms_ssim_cuda.c)¶

HIP (hip/integer_ms_ssim/ms_ssim_arith.h, hip/integer_ms_ssim/ms_ssim_score.hip, hip/integer_ms_ssim_hip.c)¶

Shared integer helpers¶

Asides (verdicts unchanged)¶

• Device memory at the cap: the CUDA and HIP integer SSIM twins allocate 6 int64 moment planes plus 1 double term plane, 7 x \(2^{33}\) B = 56 GiB. Every size is computed in size_t and passed to size_t allocator parameters (cuda/common.h:148, 273; cuda/kernel_template.h:187-188; hip/kernel_template.h:125). At that size the allocation fails cleanly; no size wraps.
• **warp_reduce(int64_t) (cuda_helper.cuh:129-130, outside this group):** the helper left-shifts the shuffled signed high half (<< 32), which is UB for negative values before C++20. That is unreachable here: the weights are ≥ 0 and below \(2^{23}\), so the high half is 0.
• Grid dimensions: the largest gridDim.y in the group is ceil(32768/8) = 4096, well under 65535. No grid expression wraps.
• **picture_copy() (picture_copy.cpp, outside this group):** the CUDA and HIP MS-SSIM hosts call it. It walks rows with ptrdiff_t pointer increments and unsigned loop counters, and has no integer product.
• Tightest margin in the group: the signed int index at cuda/integer_ms_ssim/ms_ssim_score.cu:154 (\(2^{30} - 1\) against INT32_MAX). Next is peak * peak in unsigned at cuda/integer_ms_ssim_cuda.c:306 and hip/integer_ms_ssim_hip.c:620 (131,070 below UINT32_MAX). Both are safe under the envelope.

Coverage¶

Verdict totals (126 rows): SAFE 126; OVERFLOW@16K 0; OVERFLOW@CAP-ONLY 0; DEPENDS 0.

G4b: SSIMULACRA 2, CUDA + HIP twins, ordered-sum header (integer-overflow audit)¶

Repo: master @ 571565a47. Read-only; no build, no device run. All paths below are relative to core/src/feature/.

Envelope inputs taken from the code: - VMAF_PIC_DIM_MAX 32768u (core/src/picture.c:46, checked at :185), VMAF_PIC_BPC_MIN/MAX 8u/16u (picture.c:37-38, checked at :181). - Both twins refuse w or h < 8 (cuda/ssimulacra2_cuda.c:788, hip/ssimulacra2_hip.c:931). That makes pixels >= 64 and chunks >= 1. - Chunk = 256 lanes x 4 px = 1024 px (SS2C_CHUNK_PIXELS, SS2H_CHUNK_PIXELS). A batch is 1024 chunks (SS2C_BATCH, SS2H_BATCH). - Scale pyramid: scale_w[i] = (scale_w[i-1]+1)/2. The pyramid stops before the first scale with a side below 8. Chunks are counted per plane (one channel); ceil applies.

Device buffers are sized for scale 0. That scale has the most chunks.

The brief assumed fp64 term planes (6 terms x 3 channels x N x 8 B). These do not exist. Every kernel recomputes the six fp64 terms from the five blurred fp32 planes and the two XYB fp32 planes (ss2c_terms / ss2h_terms). Per scale, the code stores only three per-(channel, sum, chunk) arrays: - chunk_sums: double, 8 B - plan: int16, 2 B - units: 2 x int64, 16 B

Neither twin has a clip-level (cross-frame) integer accumulator. Each frame's score is appended to the collector on its own, so no row is DEPENDS.

ordered_sum (shared, ordered_sum.h, ADR-1433). Used by the CUDA and HIP twins¶

The increment of one term in binade e is the term divided by u = \(2^{e-52}\), rounded. In integers this is mantissa >> (e - ex): - mantissa is in [\(2^{52}\), \(2^{53}\)). - shift = 0 gives an increment of at most \(2^{53}\) - 1. - shift >= 1 gives at most \(2^{52}\) after rounding up. - shift > 53 gives 0. - A term that does not fit (higher binade, negative, inf or NaN) gives VMAF_ORDSUM_UNFIT = \(2^{54}\).

Every increment is >= 0 (round_shifted returns whole + {0,1}).

The composition vmaf_ordsum_then saturates at \(2^{54}\). The int64 range therefore does not depend on the term count, the chunk size or N. For comparison, an uncapped sum would wrap after 512 UNFIT terms (512 x \(2^{54}\) = \(2^{63}\)). With 1024 terms at the maximum increment it would reach 1024 x (\(2^{53} - 1\)) = \(2^{63}\) - 1024, which still fits.

The cap never changes a result. A capped value (>= \(2^{54}\)) always fails total > VMAF_ORDSUM_BINADE_END in add_chunk, and that chunk is then added term by term.

SSIMULACRA 2: CUDA twin (cuda/ssimulacra2_cuda.c, cuda/ssimulacra2_cuda.h, cuda/ssimulacra2/ssimulacra2_device.cu, cuda/ssimulacra2/ssimulacra2_blur.cu)¶

There are no integer accumulators besides the ordered-sum units. The rows below are size products, index and offset math, counters and grid dims.

Input pictures are cuMemAllocPitch planes. pitch is carried as size_t (ssimulacra2_cuda.h:71, set at ssimulacra2_cuda.c:441 from (size_t)stride) and multiplied in size_t (device.cu:82).

SSIMULACRA 2: HIP twin (hip/ssimulacra2_hip.c, hip/ssimulacra2_hip.h, hip/ssimulacra2/ssimulacra2_device.hip)¶

The HIP twin does not read the picture's device pitch. submit() packs every plane into pinned staging with row_bytes = plane_w x bytes_per_sample (ss2h_stage_plane, :895-907) and copies it into a flat hipMalloc per plane. The kernel therefore indexes sy * plane_w + sx in size_t.