feat(api): import HIP device frames, dma-bufs and GL textures with event and sync_file fences (RC4 WP3, ADR-2092)

[lusoris]

Summary

RC4 work package 3, HIP lane: the VMAFx API scores frames that already live on a HIP device (device pointers, dma-bufs, HIP arrays, and OpenGL textures where the runtime can read them) with no copy through the host. On the project's pinned ROCm 10.1.0, imported frames score bit for bit as the same frames uploaded from the host for every HIP twin declared exact. The PR is stacked on the CUDA lane #2277 (rc4/api-wp3-cuda, head 15bd091e2) and reuses its GL kinds, release callback and release-event table. The lane's choices are recorded in ADR-2092, now Accepted. The measurements behind them are in Research-2159.

All evidence below comes from the pinned toolchain: build-config.env ROCM_BUILDER on master, rocm/dev-ubuntu-26.04:10.1.0-full@sha256:4f5ed1bf… (HIP runtime 7.16.26385). It ran in that image with the gfx1036 passed through. The lane was first built on the host's ROCm 7.2.4; those numbers stay as footnotes.

What it adds:

• HIP devices.
  • vmafx_device_create with VMAFX_BACKEND_HIP takes either a device index (the library creates a stream) or the caller's hipStream_t in external[0]. HIP has no context object, so external[1] must be 0.
  • vmafx_device_count, info and describe report DEVICE_POINTER, DEVICE_ARRAY, GL_TEXTURE and DMABUF memory, and NONE, HOST, HIP_EVENT, GL_SYNC and SYNC_FILE fences.
  • After vmafx_context_use_device, features registered on the context run on their HIP twins.
• One library stream per device; the twins copy on it.
  • An import is a VMAF_PICTURE_BUFFER_TYPE_HIP_DEVICE picture that carries the library stream.
  • Where a twin used to upload a host picture, it now copies the device picture device to device on that stream. Its own stream and the null stream wait for the copy on an event.
  • psnr_hvs_hip, ssimulacra2_hip and float_ms_ssim_hip level 0 used to stage planes on the host; they now copy or convert on the device.
• Import.
  • Planar device pointers are read in place, at any offset and pitch.
  • NV12, P010 and P016 are planarised on the device. The kernels come from the CUDA lane and now live in core/src/vmafx/import_convert_kernels.h, built by both nvcc and hipcc.
  • dma-bufs are imported as external memory with their own size and released behind an event. A tiled modifier is refused, and a producer size past the end of the buffer is VMAFX_E_RANGE.
  • HIP arrays are read out on the device.
• Fences.
  • A HIP_EVENT acquire fence is a wait on the library stream.
  • SYNC_FILE and GL_SYNC acquire fences are checked on the host. If one is unsignalled the import is VMAFX_E_BUSY, and the D8 rule waits and retries.
  • HOST and HIP_EVENT release fences are signalled after the last reader, through the shared release-event table (core/src/vmafx/release_events.c).
  • A SYNC_FILE release fence is refused with VMAFX_E_NOTSUP. ROCm 10.1 has no sync_file semaphore type: it refuses a DRM syncobj as an opaque descriptor (hipErrorNotSupported) and as a timeline descriptor (hipErrorInvalidValue). ROCm 7.2.4 aborts the process on the opaque descriptor.
• HIP-GL.
  • The import requires a GLX context on the device's GPU, checked before the first HIP-GL call. Without the check, ROCm 7.2 crashes after a failed setup; ROCm 10.1 refuses cleanly.
  • On ROCm 10.1 the runtime maps a GL texture but cannot read it:
    • hipMemcpy2DFromArray(Async), hipMemcpyParam2DAsync and hipMemcpy3DAsync return hipErrorInvalidValue.
    • A row copy or a texture-object read faults the GPU.
    • This happens with both Mesa versions tried: the image's 26.0.8 and the host's 26.2.4.
  • Such an import is now VMAFX_E_NOTSUP, naming desc.memory, the texture's extent and the runtime version. ROCm 7.2.4 reads these textures.
• float_vif_hip registered by default. enable_float_vif_hip_autodispatch now defaults to true, per the maintainer popup of 2026-10-06 ("Default on (Recommended)").

The ABI stays at 0.1.4. --abi-check --against-ref origin/rc4/api-generation-prototype reports definition is an append-only successor of origin/rc4/api-generation-prototype (141 additions), and the same check against origin/rc4/api-wp3-cuda reports (0 additions). python3 scripts/codegen/vmafx-api.py --check: 38 generated files match core/api/vmafx.toml.

Found on the way (rows in docs/state.md; none of them is on master, so no master PR was opened):

• test_device_target_header_dependencies.py counted only feature_src_dir kernels. It failed in a CUDA build of the CUDA lane (22 fatbins against 21) and never checked the import kernels' headers. Fixed in commit 2 (T-DEVICE-HEADER-TEST-IMPORT-KERNELS-UNCOUNTED-2026-10-06).
• The WP1 symbol checker misses version nodes in the pinned image. In the pinned image's toolchain, test_vmafx_api_generator cannot see a planted wrong version node and fails. The checker's guess that nm prints no versions is wrong there. Recorded but not changed here, because those are WP1 files (T-VMAFX-SYMBOL-VERSIONS-UNSEEN-UBUNTU-2026-10-06).
• This stack's base still pins ROCm 10.0.0. build-config.env on rc4/api-wp3-cuda is behind master; the restack picks up 10.1.0.

Landing (Q-083)

Lands bottom-up after #2287 (WP4) and #2290 (incremental motion window), one API PR at a time. The PR's whole diff (old base 15bd091e2, the #2277 head it was written on, to 20a3f4d4d) was squashed into one commit, put on the restacked chain on 2026-10-07 and is now rebased onto master b037dffb8, after #2640 took #2290's change out of #2615's squash and #2290 re-landed as its own commit (e584c79a6, maintainer decision Q-307).

ABI check against master: python3 scripts/codegen/vmafx-api.py --abi-check --against-ref origin/master -> definition is an append-only successor of origin/master (0 additions). This PR changes the documentation of existing entries only, so abi_version stays 0.1.8 and no entry is renumbered.

Conflicts, resolved per hunk when the PR was first put on the chain:

• core/api/vmafx.toml: this PR's doc text with the parent's ABI numbers.
• core/src/picture.h: this PR's comment on VMAF_PICTURE_BUFFER_TYPE_HIP_DEVICE (master had only renumbered the ADR in the old one).
• core/test/test_device_target_header_dependencies.py: the HIP count 21 and the src_dir root of this PR on top of the CUDA lane's count 22.
• docs/api/vmafx/index.md, docs/development/rebase-sensitive-invariants.md: both sides kept (window scores of feat(api): add VMAFx window scores, the window clock and the in-flight bound (RC4 WP4, ADR-2074) #2287 and the CUDA / HIP imports; both invariant entries).
• docs/state.md: the resolver, then T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01 keeps master's label and gains this PR's ROCm 10.1 sightings; T-VMAFX-SYMBOL-VERSIONS-UNSEEN-UBUNTU-2026-10-06 stays closed, as the generator PR closed it.

Rebased onto #2290 and master since: the generated core/src/AGENTS.md index (regenerated); core/src/vmafx/fence.c, where #2637 (c67a5240b) made a host fence wait on a condition variable: master's include of <pthread.h> and its waiting paragraph are kept, together with this PR's lane paragraph, which git merged; docs/state.md by the resolver; the clang-tidy records, which took master's side and were measured again (below). The rebase note is the fragment docs/rebase-notes.d/vmafx-device-frames-hip.md (ADR-2197); the rendered files (CHANGELOG.md, ADR indexes, docs/rebase-notes.md) are master's.

Fixed while landing (the stack moved under the PR):

• core/src/meson.build appended the HIP import sources to libvmaf_sources, which the WP6 split renamed to libvmafx_sources, so a HIP build did not configure.
• core/src/hip/import_device.c called vmafx_engine_leave(previous); since WP4 it takes the context.
• core/src/vmafx/import_convert_kernels.h (shared by CUDA and HIP) had an anonymous namespace and __global__ definitions in a header: 4 clang-tidy findings in the hip lane. The header now holds static __device__ __forceinline__ bodies and import_convert.cu / import_convert.hip define the extern "C" entry points, as feature/float_moment_sum_gpu.h does.
• core/test/test_hip_shared_frame_contract.py: master's ruff (PLR2004) refused a literal 2; it is a named constant now.
• core/src/AGENTS.d/vmafx-device-frames.md named the release callback "ABI 0.1.4"; the definition says 0.1.7.
• ci(tidy): the cpu lane's baseline records 10 of this PR's units on top of feat(motion): derive motion2 and motion3 frame by frame so VMAF windows complete before the flush (ADR-2090) #2290's (the earlier record fell away in the restack); the cuda lane's records its 19.

Rebased onto master b037dffb8 after #2290 landed: the feature and docs commits applied unchanged (git range-diff: =); the cpu clang-tidy baseline, which #2290, #2646, #2647 and #2643 had moved, took master's side and the 10 cpu units were measured again (0 findings); the cuda baseline is this PR's record, master has not changed it since.

No integration-branch commit is carried.

Local gate (on 142a89892, master b037dffb8)

• CPU (-Db_lto=false, -j4, warnings as errors): build 0 warnings; --suite=fast 428 OK, 2 failed: test_vmafx_window_live and test_vmafx_window_cli, the wall-clock budget under host load (T-WINDOW-LIVE-LATENCY-UNDER-LOAD-2026-10-09, moved to the serial timing suite by fix(test): check the live window budget alone in a timing suite and its pacing on a virtual clock #2664); both pass on a re-run through the runner and 3 of 3 direct runs. test_gpu_picture_pool_uaf on its own with MALLOC_PERTURB_=0: OK; codegen tests 207 passed; make test-netflix-golden GOLDEN_NINJA_JOBS=4 280 passed, 3 skipped; preflight.sh --stage msvcism pass; affected suites: tooling 2648 passed, 6 skipped. vmafx-api.py --check: 80 generated files match; --abi-check --against-ref origin/master: 0 additions.
• HIP on the pinned toolchain: image vmafx-hip-lane:rocm10.1.0-egl (the pinned rocm/dev-ubuntu-26.04:10.1.0-full@sha256:4f5ed1bf…, HIP 7.16.26385, plus the EGL / GLES / X11 development libraries; ROCm unchanged), gfx1036, every device run flock hip-gfx1036.lock timeout 290. Build (-Denable_hip=true -Denable_hipcc=true -Dhip_gfx_targets=gfx1036, -Dwerror=true, fatal link warnings): 0 warnings. test_vmafx_import_hip 16/16; test_vmafx_import_hip_fence 8/8; test_vmafx_import_hip_gl skipped with the ROCm 10.1 refusal, as designed (feat(hip): import GL textures through EGL dma-buf export so the pinned ROCm 10.1 reads them (ADR-2132) #2360 lands the EGL route); test_vmafx_fence_kinds 5/5; test_hip_shared_frame 9/9; with feat(motion): derive motion2 and motion3 frame by frame so VMAF windows complete before the flush (ADR-2090) #2290's advance callback on master underneath: test_hip_motion_five_frame_window, test_hip_motion_parity, test_hip_motion_v2_parity pass. test_vmafx_import_hip_bitexact, one run per clip (0 to 3 and the 4K pair): 354 cells, 16098 values, 0 cells differing, 0 attempts repeated, 9042 imports, 6028 conversions, 0 host copies.
• Tidy (dev container, clang-tidy 22.1.8, scripts/dev/tidy-lane.sh --write --only ...), measured on this tree (cpu again after the rebase onto b037dffb8): cpu core/src/libvmaf.c, core/src/vmafx/{device,device_context,fence,frame_import,frame_import_admit,frame_pool,release_events,submit}.c, core/test/test_vmafx_fence_kinds.c (10 TUs) 0 findings, recorded in the cpu baseline (gl_sync.c and sync_file.c, 0 findings as well, stay read by the cuda and hip lanes only: feat(api): import SYCL device frames with event fences, dma-bufs and GL textures (RC4 WP3, ADR-2091) #2342 folds them into sync_object.c); cuda the CUDA import sources and tests with the shared units (19 TUs) 0 findings; hip the 30 units of the HIP import, the HIP twins this PR touches and the shared units, 0 findings (record unchanged).
• Train gates (deliverables, docs/state.md touch and rows, silent revert against master, praetorctl audit): pass.

The earlier device evidence on this PR (2026-10-07, before the restacks) is in the sections below and in T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01.

Type

• feat — new feature

Checklist

• Commits follow Conventional Commits (the commit-msg hook enforces this).
• make format && make lint is green locally — the commit hooks pass (clang-format, markdownlint, semgrep, source ADR citations, generated-index freshness, FFmpeg patch stack, HISS audit, assertion density).
• Unit tests pass.
  • Pinned ROCm 10.1.0 image (built there with -Denable_hip=true -Denable_hipcc=true -Dhip_gfx_targets=gfx1036): fast suite 451 OK, 1 skipped, and 1 failed (test_vmafx_api_generator, the WP1 row above). The fast suite includes the 72 fast+gpu HIP twin tests and the lane's device-free contracts. The lane programs pass as listed below.
  • Host builds:
    • CPU --suite=fast: 366 OK.
    • Device-free fast suites: HIP 380 OK, CUDA + HIP 389 OK.
    • CUDA lane on the RTX 4090: test_vmafx_import_cuda 12/12, _fence 6/6, _gl 2/2, _bitexact 236 cells with 0 differing.
    • Host ROCm 7.2.4: test_hip_exact_twin_matrix and test_hip_v1_models_no_fallback pass.
• If I touched any SIMD/GPU code path, I ran /cross-backend-diff and the worst ULP is ≤ 2. — Every exact HIP cell compares imported frames with host uploads at 0 ULP (test_vmafx_import_hip_bitexact). The twins' arithmetic is unchanged.
• If I touched a feature extractor with SIMD/GPU twins, I either updated every twin or listed the gap under "Known follow-ups" below. — Only HIP twins are touched; the CUDA lane changes are moves into shared files.
• If I added a new .c / .cpp / .cu / .h / .hpp, it has the appropriate license header (EUPL-1.2, fork-authored).
• If this is a breaking change, the commit message uses ! or BREAKING CHANGE: and the migration path is documented below. — not a breaking change: no ABI change.
• 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 — ADR-2092, number claimed with scripts/adr/next-free.sh --claim; indexes regenerated.

tidy: hip core/src/hip/{import_device,import_dmabuf,import_fence,import_frame,import_gl,picture_hip,shared_frame}.c, core/src/feature/hip/{float_vif_hip,integer_ms_ssim_hip,integer_psnr_hvs_hip,ssimulacra2_hip}.c, core/src/libvmaf.c, core/src/vmafx/{device,device_context,fence,frame_import,frame_import_admit,frame_pool,gl_sync,release_events,submit,sync_file}.c, core/test/test_hip_shared_frame.c, core/test/test_vmafx_fence_kinds.c, core/test/test_vmafx_import_hip{,_bitexact,_fence,_gl}.c (28 TUs; through them vmafx_hip{,_internal}.h, vmafx_hip_test_util.h and vmafx_device_cells.h are covered) 0 findings, 0 uncited NOLINT. After the ROCm 10.1 changes, import_frame.c and the four lane tests were re-measured: 0 findings. cpu: the touched core/src/vmafx/*.c, core/src/libvmaf.c and test_vmafx_fence_kinds.c (12 TUs) 0 findings. cuda: core/src/cuda/{import_device,import_fence,import_frame,import_gl}.c, core/test/test_vmafx_import_cuda{,_bitexact,_fence,_gl}.c and the shared core/src/vmafx/*.c (19 TUs) 0 findings. All runs used the dev container, clang-tidy 22.1.8 and scripts/dev/tidy-lane.sh --only ... <lane>. The first hip run found 18 findings and the first cuda run 1. All were fixed without NOLINT.

scripts/dev/preflight.sh --stage msvcism: pass. core/test/test_win32_pthread_shim_contract.py: pass. scripts/ci/assertion-density.sh: pass. praetorctl audit: governance gates passed, with no HISS finding in a touched file.

Bug-status hygiene (ADR-0165)

• docs/state.md updated.
  • Closed: T-HIP-NO-IMPORT-PATH-2026-10-05 (evidence on 10.1) and T-DEVICE-HEADER-TEST-IMPORT-KERNELS-UNCOUNTED-2026-10-06.
  • Opened: T-HIP-VMAFX-NO-FRAME-POOLS-2026-10-06 (RC4), T-VMAFX-SYMBOL-VERSIONS-UNSEEN-UBUNTU-2026-10-06 (RC4) and T-HIP-IMPORT-TWIN-DEVICE-COPY-2026-10-06 (RC8).
  • Three platform rows deferred on ROCm, measured on 10.1: T-HIP-ROCM-NO-SYNC-FILE-SEMAPHORE-2026-10-06, T-HIP-ROCM10-GL-TEXTURE-READ-2026-10-06 and T-HIP-GFX1036-UNFLUSHED-STREAM-ORDER-2026-10-06.
  • New 10.1 sightings on T-HIP-GFX1036-DROPPED-DISPATCHES-2026-10-01.

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 — not applicable.

Golden gate (GOLDEN_NINJA_JOBS=4 make test-netflix-golden): 280 passed, 3 skipped (core/build-golden built with gcc).

Cross-backend numerical results

Measured on ryzen-4090-arc's AMD gfx1036 (integrated GPU, Linux 7.2.9-1-cachyos), inside the pinned ROCm 10.1.0 image. The image had /dev/kfd and /dev/dri passed through, the render and video groups added, and the X display mounted. Every device run was flock ~/.cache/vmafx-locks/hip-gfx1036.lock timeout 300 docker run … timeout 290 <test>. Other lanes loaded the host throughout (load average 13 to 48).

Exit evidence (ADR-1829)	Test	ROCm 10.1.0 (pinned)
Imported frames score bit for bit as host-uploaded frames for every exact HIP twin	test_vmafx_import_hip_bitexact: every cell of scripts/ci/exact_twins.d/*.hip in three layouts — planar with odd offsets and pitches, semi-planar, and semi-planar in dma-bufs (through the D8 rule, with an exported sync_file)	Netflix 576x324: 69 cells, 11808 values. Checkerboard 1 px and 10 px: 72 cells, 765 values each. Sparks 10-bit: 69 cells, 1230 values. BBB 3840x2160 (6 frames): 72 cells, 1530 values. 354 cells, 16098 values, 0 differing, no cell repeated; 9042 imports, 6028 conversions, 0 host copies1
A skipped acquire wait gives wrong scores under load; the real wait 0 bad of N	test_vmafx_import_hip_fence	HIP_EVENT: 0 bad of 16 with the wait for psnr and vif, 16 of 16 with VMAFX_TEST_SKIP_ACQUIRE_WAIT. SYNC_FILE (a dma-buf written through GBM, with its exported sync_file as acquire fence): 32 pending at import, 32 waited for, 0 bad. With the check skipped, 32 imports took an unsignalled fence.
Release-fence canary	test_release_canary	0 bad of 16; with VMAFX_TEST_EARLY_RELEASE, 16 of 16
Host-copy counter 0	every lane program	0; with VMAFX_TEST_FORCE_HOST_COPY, 3
One import scored by two contexts	test_one_import_two_contexts (vmaf_v1.0.16_3d0h and vmaf_v0.6.1)	56 values each, 0 differing; release only after the second context
Arrays, dma-bufs	test_vmafx_import_hip	NV12, P010 (16-bit) and planar arrays, and planar and NV12 dma-bufs: 56 values each, 0 differing. 8 of 8 sync_files pending at import were VMAFX_E_BUSY.
GL textures	test_vmafx_import_hip_gl	Skipped, with the runtime's refusal printed: the HIP runtime (version 71626385) maps texture 0 (array 640x360) but refuses to read it: hipErrorInvalidValue.2

Vendor profiler. rocprofv3 --memory-copy-trace --hip-runtime-trace --kernel-trace, run in the image on the checkerboard import sessions (VMAFX_TEST_IMPORT_ONLY=1 VMAFX_TEST_CLIPS=1, 432 imports):

• The library stream ran 441 copy kernels (__amd_rocclr_copyBufferRect*, the device-to-device copies), 288 NV12 conversions and 90 float_ms_ssim_hip level-0 conversions. It ran no memory copy at all — none with a host side.
• The host-to-device copies are the test producer's uploads (12) and the twins' tables (6).
• The device-to-host copies are the twins' result read-backs, on their own streams.
• The Netflix sessions with the runtime trace abort rocprofv3 at finalisation (ring_buffer.cpp:106 mmap failed with errno 22). Their HIP API log (AMD_LOG_LEVEL=3, 6624 imports) shows 7056 device-to-device copies on the library stream and no other copy there.

Re-tests requested on 10.1:

Item	ROCm 10.1.0	ROCm 7.2.4 (host)	Decision
SYNC_FILE release (external semaphore from a DRM syncobj)	opaque: hipErrorNotSupported; timeline: hipErrorInvalidValue	opaque: process abort	Refusal kept; version recorded in ADR-2092
HIP-GL failed first setup	No crash: hipGLGetDevices() without a context or under another vendor's GLX context returns hipErrorInvalidValue, and a later GLX call works	Crash	GLX check kept (it costs nothing and guards 7.2)
HIP-GL read-out	Every read fails (hipErrorInvalidValue; row copy or texture read faults the GPU), with Mesa 26.0.8 and 26.2.4	Works	GL import refused, VMAFX_E_NOTSUP naming the runtime
hipStreamQuery() after import	Without it, the planar array case fails in 8 of 8 runs (6 values, every attempt)	5 of 6	Kept
Dropped dispatches (probe, 100000 frames)	134, 98, 179, 110 and 103 bad frames in 5 runs	0, 17 and 3 per 20000	Tests keep repeating and reporting

Gates shown failing on planted defects (measured on this branch)

Planted defect	Test	Result
acquire wait skipped (VMAFX_TEST_SKIP_ACQUIRE_WAIT)	test_vmafx_import_hip_fence	psnr and vif 16 of 16 bad on 10.1 (asserted every run)
sync_file check skipped	test_vmafx_import_hip_fence	32 imports took an unsignalled sync_file (asserted)
release recorded at submit (VMAFX_TEST_EARLY_RELEASE)	test_vmafx_import_hip_fence	16 of 16 canaries on 10.1 (asserted)
planar import staged through the host (VMAFX_TEST_FORCE_HOST_COPY)	test_vmafx_import_hip	host-copy counter 3 (asserted)
hipStreamQuery() after the import removed	test_vmafx_import_hip	planar arrays fail in 8 of 8 runs on 10.1
unreadable GL texture reported as a device failure (the code before the refusal)	test_vmafx_import_hip_gl	fails on 10.1 with VMAFX_E_DEVICE; the contract test refuses the source
null stream not waiting for the copies	test_vmafx_import_hip_bitexact	adm frame 0: 7 values differ in every attempt (7.2.4)
float_vif_hip without the HIP flag	test_vmafx_import_hip_bitexact	the float_vif session is refused by admission
release event never recorded	test_vmafx_import_hip	test_release_fence_kinds: not signalled after the release
release callback before the recording	test_vmafx_import_hip	test_release_callback fails (the 5 s wait runs out)
HIP pool refused with the generic host-frame message	test_vmafx_import_hip	test_no_frame_pools fails
second twin reading the shared frame without the wait	test_hip_shared_frame (device-free fakes)	fails
eleven source regressions — see list below	test_vmafx_import_hip_contract.py	each refused (the planted cases run every time)
cells table missing, adding or mis-aliasing a declared exact HIP twin	test_vmafx_import_cuda_cells_contract.py	each refused
vmafx/import_convert_kernels.h missing from a shared-header list	test_device_target_header_dependencies.py	fails naming the header

The eleven source regressions the contract test refuses:

• host staging of a device picture in psnr_hvs_hip, ssimulacra2_hip and float_ms_ssim_hip level 0
• upload of a device picture as host memory
• missing null-stream wait
• unsubmitted import
• HIP-GL called before the GLX check
• dma-buf imported with the producer's size
• a size past the end of the dma-buf accepted
• an unreadable GL texture not refused
• a faked SYNC_FILE release

Performance (if perf or feat)

No timing claim (RC8). An import costs one device-to-device copy per plane into the twins' buffers, on the library stream. The shared frame makes that copy once per frame for every twin that reads the plane. Reading producer memory in place is the RC8 row T-HIP-IMPORT-TWIN-DEVICE-COPY-2026-10-06.

Deep-dive deliverables (ADR-0108)

• Research digest — docs/research/2159-vmafx-hip-device-frames.md: runtime behaviour on ROCm 10.1 (with 7.2.4 for comparison) — external memory and semaphores, events, hipFree, pools, GBM dma-bufs, unsubmitted work, HIP-GL setup and read-out, dropped commands — plus the exit evidence.
• Decision matrix — docs/adr/2092-vmafx-hip-device-frames.md ## Alternatives considered.
• AGENTS.md invariant note — new core/src/hip/AGENTS.d/vmafx-device-frames.md; core/src/hip/AGENTS.d/host-picture-staging.md, core/src/feature/hip/AGENTS.d/picture-upload-sync.md and core/src/AGENTS.d/vmafx-device-frames.md updated; docs/development/rebase-sensitive-invariants.md entry "VMAFx HIP frames are read on one stream per device".
• Reproducer / smoke-test command — below.
• CHANGELOG fragment — changelog.d/added/api-hip-device-frames.md, changelog.d/changed/float-vif-hip-autodispatch-default.md, changelog.d/changed/hip-twins-device-level0.md.
• Rebase note — fragment docs/rebase-notes.d/vmafx-device-frames-hip.md (ADR-2197), "VMAFx device frames on HIP (RC4 WP3 HIP lane)".

User documentation:

• docs/api/vmafx/index.md gains a "HIP devices" section (descriptors, memory kinds, the ROCm 10.1 GL refusal, fences, a dma-buf example, ordering, profiling).
• docs/backends/hip/overview.md, uploads.md, build-flags.md, vif.md and features.md are updated to match.

Reproducer

The image is the pinned ROCM_BUILDER, plus meson, nasm, zimg, xxd, gbm, drm, X11 and Mesa GLX for the build and the tests:

FROM rocm/dev-ubuntu-26.04@sha256:4f5ed1bf6532a4b9920400401b0ad0f706af356dae61e40afefbfd6c8073f0ce
RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
      meson ninja-build nasm pkg-config xxd libzimg-dev libgbm-dev libdrm-dev libx11-dev \
      libgl-dev libgl1-mesa-dri libglx-mesa0 python3-pytest python3-numpy git ca-certificates

docker build -t vmafx-hip-lane:rocm10.1.0 -f Containerfile .
R="docker run --rm --init --device /dev/kfd --device /dev/dri \
  --group-add $(stat -c %g /dev/kfd) --group-add $(stat -c %g /dev/dri/card0) \
  --user $(id -u):$(id -g) -e HOME=/tmp -e DISPLAY -e XAUTHORITY -v /tmp/.X11-unix:/tmp/.X11-unix:ro \
  -v $XAUTHORITY:$XAUTHORITY:ro \
  -v $PWD:$PWD -w $PWD vmafx-hip-lane:rocm10.1.0"
$R bash -c 'meson setup build-rocm10 core -Denable_hip=true -Denable_hipcc=true \
  -Dhip_gfx_targets=gfx1036 -Denable_cuda=false -Denable_sycl=false -Db_lto=false \
  -Denable_dnn=disabled && ninja -C build-rocm10 -j6'
for t in test_vmafx_import_hip test_vmafx_import_hip_bitexact test_vmafx_import_hip_fence \
         test_vmafx_import_hip_gl; do
  flock ~/.cache/vmafx-locks/hip-gfx1036.lock timeout 300 $R timeout 290 ./build-rocm10/test/$t
done
python3 core/test/test_vmafx_import_hip_contract.py
VMAFX_TEST_IMPORT_ONLY=1 VMAFX_TEST_CLIPS=1 flock ~/.cache/vmafx-locks/hip-gfx1036.lock timeout 300 \
  $R rocprofv3 --memory-copy-trace --hip-runtime-trace --kernel-trace -f rocpd -d /tmp/rp -o imp \
  -- ./build-rocm10/test/test_vmafx_import_hip_bitexact

The fixtures are python/test/resource/yuv and testdata/bbb (4K). The dma-buf cases need libgbm and the device's render node, found through sysfs when /dev/dri/by-path is missing, as in a container. The GL test needs an X display and a GLX driver for the HIP device's GPU, and sets DRI_PRIME itself.

Known follow-ups

• Restack onto master together with feat(api): import device frames with fences, pools and the import rule on the CPU device (RC4 WP3, ADR-1929) #2213 and feat(api): import CUDA device frames with event fences and GL textures (RC4 WP3, ADR-2023) #2277; this also brings build-config.env to ROCm 10.1.0. Until then check-silent-revert.py runs against origin/rc4/api-wp3-cuda.
• HIP-GL on the pinned toolchain waits for a ROCm that reads mapped GL textures (T-HIP-ROCM10-GL-TEXTURE-READ-2026-10-06).
• HIP frame pools are refused (T-HIP-VMAFX-NO-FRAME-POOLS-2026-10-06).
• The twins copy imported frames once more on the device (RC8, T-HIP-IMPORT-TWIN-DEVICE-COPY-2026-10-06).
• The WP1 symbol checker in the pinned image (T-VMAFX-SYMBOL-VERSIONS-UNSEEN-UBUNTU-2026-10-06) belongs to the WP1 owner.
• The hipStreamQuery() workaround stays until a ROCm or driver update passes 8 runs without it.
• WP9 (filters): VAAPI and DRM frames can come in through the dma-buf path with their exported sync_file.

Footnotes

1. On the host's ROCm 7.2.4 the same cells had 0 differing in every attempt, with 2 cells repeated once (1 value each, the platform defect). ↩

2. On ROCm 7.2.4: 42 values, 0 differing. ↩