perf(hip): pass AdmBufferHip to the ADM kernels by pointer again (ADR-0759)

[lusoris]

Summary

Re-applies ADR-0759 (Accepted): the four HIP ADM kernels take const AdmBufferHip * __restrict__ again, instead of the 328-byte struct by value. 31a51afb2 (#101) implemented it, and the next merge, 92ea978a4 (#102, an unrelated CUDA ciede PR built from an older base), reverted every HIP hunk the same day. The HIP AGENTS.md note and an unreleased changelog fragment kept describing the by-pointer form ever since. Stacked on #1478; review #1476 and #1478 first.

Precondition checked. One device copy uploaded at init is correct only if the struct never changes per launch. In today's integer_adm_hip.c, s->buf is written only during init (adm_hip_alloc_buffers(), adm_hip_slice_bands(), adm_hip_slice_results()) and on teardown. submit_fex_hip() passes &s->buf unchanged to all four launch helpers, and no launch builds a modified copy.

The change.

• adm_hip_upload_buf() copies the struct to the device once, at the end of adm_hip_init_device().
• The four launches pass (void *)&s->buf_dev.
• The copy is freed in close and on both init failure paths.

Measured honestly (gfx1036, from the HSACO metadata):

Kernel	Kernel arguments (bytes)	Scratch	VGPR
adm_csf_kernel_1_4	856 → 536	0 → 0	44 → 38
i4_adm_csf_kernel_1_4	856 → 536	0 → 0	53 → 45
i4_adm_cm_line_kernel	904 → 584	0 → 0	59 → 59
adm_cm_line_kernel_8	968 → 648	936 → 936	128 → 128

• The VGPR drops come from the lint restructure of adm_csf.hip, not the pointer. Pointer-only and lint-only variants were built separately to tell the two apart.
• The 936 bytes of scratch are register spills in adm_cm_line_kernel_8 (239 spilled at the 128-register cap), not a copy of the struct. That is a separate follow-up.
• No throughput change: 1080p 10-bit, ten alternating runs, median 29.54 fps before vs 30.02 after, with runs spread from 27.4 to 31.1. The decision is restored because it is Accepted and cuts 320 bytes of launch arguments per kernel, not for speed.

Lint. adm_csf.hip had 35 clang-tidy findings; touching it meant taking it to 0 (baseline tightened). integer_adm_hip.c and adm_cm.hip stay at 0.

no docs needed: no user-visible change; scores are byte-identical and the extractor's options, output and supported inputs are unchanged.

Type

• feat — new feature
• fix — bug fix
• perf — performance improvement
• refactor — no behavior change
• docs — documentation only
• test — test-only
• build / ci — tooling / infra
• port — cherry-pick from upstream Netflix/vmaf
• sycl / cuda / simd — backend-specific

Checklist

• Commits follow Conventional Commits (the commit-msg hook enforces this).
• make format && make lint is green locally: 0 clang-tidy findings and 0 uncited NOLINTs on the three touched sources (hip lane); cppcheck and pre-commit pass.
• Unit tests pass: meson test -C build: HIP fast suite 181 OK + 1 existing expected fail; all eight HIP ADM tests pass, including test_hip_adm_tiny_frames and the large parity test.
• If I touched any SIMD/GPU code path, I ran /cross-backend-diff and the worst ULP is ≤ 2: 0; HIP output byte-identical before and after on 14 cases (src01 at 8/10/12/16-bit, odd 575x323 and 197x101 at 8/10/16-bit, bright 16-bit noise, both 1080p checkerboard pairs, a 60-frame 1080p 10-bit clip).
• If I touched a feature extractor with SIMD/GPU twins, I either updated every twin or listed the gap under "Known follow-ups" below: HIP only. The CUDA twin passes by value by decision (ADR-0756 chose __ldg() extraction and put adm_cm out of scope).
• If I added a new .c / .cpp / .cu / .h / .hpp, it has the appropriate license header (see CONTRIBUTING.md): no new files.
• If this is a breaking change, the commit message uses ! or BREAKING CHANGE: and the migration path is documented below: not breaking.
• If this PR adds an ADR, the ADR row lives in docs/adr/_index_fragments/<NNNN-slug>.md and the slug is appended to docs/adr/_index_fragments/_order.txt: no new ADR; ADR-0759 is Accepted and becomes true again.

Bug-status hygiene (ADR-0165)

• docs/state.md: T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed.

Netflix golden-data gate (ADR-0024)

• I did not modify any assertAlmostEqual(...) score in the Netflix golden Python tests.
• If I believe a golden value must change, I have explained why below AND pinged @lusoris for a CODEOWNERS exception: no score changes.

Deep-dive deliverables (ADR-0108)

• Research digest — no digest needed: re-applies an existing decision; its digest docs/research/research-0759-hip-adm-buffer-by-pointer.md gains a dated addendum correcting its size and CUDA-mirroring claims.
• Decision matrix — no alternatives: only-one-way fix (ADR-0759's Accepted decision, re-applied).
• AGENTS.md invariant note — core/src/feature/hip/AGENTS.md.
• Reproducer / smoke-test command — below.
• CHANGELOG fragment — changelog.d/fixed/hip-adm-buffer-by-pointer-reapplied.md.
• Rebase note — docs/rebase-notes.md.

Reproducer

meson setup build/hip core -Denable_hip=true -Denable_hipcc=true -Denable_cuda=false \
  -Denable_sycl=false -Db_lto=false --buildtype=release
ninja -C build/hip && meson test -C build/hip --suite fast
build/hip/tools/vmaf -r ref.yuv -d dis.yuv -w 1920 -h 1080 -p 420 -b 10 \
  --backend hip --feature adm_hip --no_prediction --precision max -o out.json --json

Known follow-ups

• adm_cm_line_kernel_8 spills 239 registers (936 bytes of scratch per thread) at the 128-VGPR cap; __launch_bounds__(128) alone still spills 112. A register-pressure pass is its own change.
• AdmFixedParametersHip (248 bytes) is still passed by value; ADR-0759 did not cover it.
• The HIP tidy baseline still lists adm_decouple.hip (49 findings), a file deleted under ADR-1154; a full re-measure (L-41) will drop it.
• Two silent reverts of this kind (perf(cuda): ciede 8/16bpc — __ldg() read-only cache routing (F3 fix, ADR-0762) #102 over perf(hip): pass AdmBufferHip by pointer in ADM kernels (F3 fix, ADR-0759) #101, and the rename merge that brought back the i686 and MinGW64 lanes) prompted a history audit, which is running.

[lusoris]

Absorbed into the ADM stack train #1507, per your direction to fold this stack the way #1506 was folded.

This PR targeted the one below it in a five-deep stack, so none of the five could merge until every one below had merged and been restacked — five sequential rebase-plus-CI rounds. #1507 is one. Your work is in it unchanged; that PR's description lists the six defects the fold itself surfaced, none of which an individual PR could see, because each gate only looks at the files its own PR touches.

The branch stays on the remote.