Repository navigation
perf(hip): pass AdmBufferHip by pointer in ADM kernels (F3 fix, ADR-0759) - #101
Merged
Merged
Conversation
lusoris
marked this pull request as ready for review
May 29, 2026 09:37
…759) Research-0755 / PR #95 HIP audit identified AdmBufferHip (~272 bytes) being passed by value in four __global__ kernel signatures: - adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4 (adm_csf.hip) - i4_adm_cm_line_kernel, adm_cm_line_kernel_8 (adm_cm.hip) This PR replaces `AdmBufferHip buf` with `const AdmBufferHip * __restrict__ buf_ptr` in all four kernels. A device-side copy of the struct is allocated once in init_fex_hip via hipMalloc + hipMemcpy(hipMemcpyHostToDevice) and stored as AdmStateHip::buf_dev. Host dispatch helpers pass &buf_dev in args[]. The struct contains only stable device pointers set at init time; no per-frame update is required. Change is numerically transparent (pointer indirection only; all load addresses are identical). Matches CUDA F3 fix pattern from PR #93. Build verified: meson -Denable_hipcc=true -Denable_cuda=false succeeds (35 targets, zero errors). Runtime verification pending — no AMD GPU on audit host. Deliverables: ADR-0759, Research-0759, AGENTS.md invariant, changelog fragment, rebase-notes entry, state.md row. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris
force-pushed
the
perf/hip-adm-buffer-by-pointer-20260529
branch
from
May 29, 2026 09:39
e25cff3 to
7caac96
Compare
5 tasks
14 of 26 tasks
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 20, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy, in 31a51af (#101). The next merge, 92ea978 (#102, a CUDA ciede change cut from an older base), put the by-value signatures back, and every launch since has copied the 328-byte struct into the kernel arguments. This re-implements the decision on today's structure instead of reverting the revert. The design needs the struct to be the same on every launch. It is: s->buf is written only by adm_hip_alloc_buffers(), adm_hip_slice_bands() and adm_hip_slice_results() during init and by adm_hip_free_buffers() on teardown. submit() passes &s->buf to every launch helper, and no launch passes a modified copy; the scale 1-3 CSF and CM kernels get the scale as a separate argument and use the same i4 band pointers at every scale. adm_hip_upload_buf() allocates and uploads the copy at the end of adm_hip_init_device(), after the band and result slices are final, and the four launches pass &s->buf_dev. The copy is freed in close() and on both init failure paths. HIP output is byte-identical at %.17g before and after, with debug features, on 8, 10, 12 and 16-bit input, odd 575x323 and 197x101 frames, bright 16-bit noise, the 1080p 8-bit checkerboard pairs and a 60-frame 1080p 10-bit clip. The kernel argument segment of each kernel shrinks by 320 bytes. Per-thread scratch and VGPRs do not change because of the pointer: the 936 bytes on adm_cm_line_kernel_8 are VGPR spills (239 at the 128-register cap), not a copy of the struct. End-to-end 1080p 10-bit throughput is unchanged within run-to-run noise. adm_csf.hip had not been lint-cleaned yet, so this also takes it from 35 clang-tidy findings to 0: the device helpers move into an anonymous namespace, locals become const, and the band select becomes a helper. That restructure, not the pointer, lowers the two CSF kernels from 44 to 38 and 53 to 45 VGPRs. HIP tidy baseline tightened.
lusoris
added a commit
that referenced
this pull request
Sep 22, 2026
ADR-0759 moved the four HIP integer ADM kernels that read AdmBufferHip (adm_csf_kernel_1_4, i4_adm_csf_kernel_1_4, i4_adm_cm_line_kernel, adm_cm_line_kernel_8) from the struct by value to a pointer to a device copy. It landed in 31a51af (#101); the next merge, 92ea978 (#102, an unrelated CUDA ciede change squashed from an older base), put the by-value signatures back the same day. The invariant note in core/src/feature/hip/AGENTS.md survived the revert, so the file has since carried both a "P1 known issue: struct-by-value" section and a "MUST be passed by pointer - Resolved" section, with the code matching the first. Documentation that contradicts the code is what the next rebase trusts, which is worse than either half alone. This re-implements the decision on today's structure rather than reverting the revert, and rewrites the AGENTS.md section to say what is true: the pointer convention, the revert history so it is recognised if it happens again, the host-side contract, and the fact that AdmFixedParametersHip (248 bytes) is still by value by decision, not by oversight. The design needs s->buf to be identical on every launch. It is: the pointer fields are written only by the allocation and the two slicing blocks in init_fex_hip() and cleared in close_fex_hip(), and submit() passes &s->buf to the launch helpers without modifying it. AdmStateHip gains a void *buf_dev device copy, uploaded at the end of init_fex_hip() once every pointer inside s->buf is final, passed as args[0] by address, and freed in close_fex_hip() and on the new fail_buf_dev init path. The upload shares the failure exit the feature-name dictionary already used, so the file gains no goto and no oversized block: praetorctl reports 1063 total HISS infractions before and after. Measured with hipcc --genco (ROCm 7.2.53211), reading each kernel's metadata with clang-offload-bundler --unbundle and llvm-readelf --notes. Identical on gfx1036, gfx1100 and gfx90a: adm_csf_kernel_1_4 kernarg 856 -> 536 i4_adm_csf_kernel_1_4 kernarg 856 -> 536 i4_adm_cm_line_kernel kernarg 904 -> 584 adm_cm_line_kernel_8 kernarg 968 -> 648 adm_cm_reduce_line_kernel_4 kernarg 296 -> 296 (does not read the struct) 320 bytes = sizeof(AdmBufferHip) - sizeof(void *), off every launch. The by-value copy was also spilling: adm_cm_line_kernel_8's per-thread scratch drops 920 -> 608 bytes on gfx1036, 932 -> 616 on gfx90a and 664 -> 352 on gfx1100, and its SGPR count 86 -> 78 on gfx1036. ADR-0759 and the earlier notes quote the struct as ~272 bytes and the parameter struct as ~244; a harness compiled against the real header prints 328 and 248. The frozen ADR body keeps its numbers; the correction is recorded in Research-0759. Verified on the host's real AMD device (gfx1036, the Raphael iGPU, ROCm 7.2): a differential harness loaded both HSACO variants into one process, gave them byte-identical randomised band contents and parameters, and compared every output buffer - csf_f bands 1-3, i4_csf_f bands 1-3, the tmp_accum per-thread scratch and the three int64 adm_cm accumulators. All four kernels are bit-identical, over three shape/seed configurations (98x50 stride 128, 33x17 stride 36, 160x90 stride 160), with non-zero output throughout. Not run: an end-to-end vmaf --backend hip score. The workstation was saturated by sibling agents and no full enable_hipcc build was made. Refs: ADR-0759, Research-0759.
lusoris
added a commit
that referenced
this pull request
Sep 22, 2026
A pull request can remove work from master with nothing in its review diff looking like a removal of somebody else's commit. `92ea978a4` (#102, a CUDA ciede change) reset three HIP ADM files to the exact blobs they held before `31a51afb2` (#101, ADR-0759), merged 47 seconds earlier: the post-#102 blob of `core/src/feature/hip/integer_adm_hip.c` is `baf2b339`, byte-identical to the pre-#101 blob, and the 50 lines #101 added are the 50 lines #102 removed. Every check stayed green. Because #101, its ADR, its AGENTS.md note and its state row all stayed in tree, four months of audits read the change as shipped; it was restored only by PR #1481, and the read-only audit of master's first-parent history found dozens more in the same class. `scripts/ci/check-silent-revert.py` measures the real merge result (`git merge-tree --write-tree`, not the branch tree in isolation) and fails when it removes target work the branch never set out to touch: a file rewound to an older blob it held on the target's history, a target commit undone hunk-for-hunk, or lines dropped or resurrected by a conflict resolution rather than by any non-merge commit of the branch. The first two cover the rebased-branch shape that produced #102, where intent and effect are the same diff and only history can show the content is old; the last two lift the principle of the existing post-restack `check-resurrected.py` to a base/head pair. It runs as the required context `Silent-Revert Guard` on every non-draft PR, resolving the live target tip rather than the PR's recorded base because the defect class is master moving after the branch was cut, and locally as `make silent-revert-check`. A deliberate revert is declared where a reviewer sees it (`revert:` title, `reverts: #N`, `intentional revert: <reason>`); there is no in-tree suppression. The gate fails closed on an unresolvable ref, unrelated histories, a conflicting merge, or a git below 2.38. Replayed over master's last 30 first-parent commits it fires twice, both correctly: `5a467e629` re-lands a migration the `c2a3c7e0f` mega-squash had reverted, and `1beb3b8a9` removes an earlier fix along with its changelog fragment. Both need a one-line declaration, which is the intended workflow. A third case, a comment rewrite in `core/tools/vmaf.cpp`, was a false positive and tightened the predicate out of existence. `dropped` and `resurrected` check the surviving tree before reporting, so a line still present in the merged file was relocated, not lost. Dogfooding the gate against a 48-commit stack showed why: without that check both fire on a diff-alignment artefact, where inserting text above a line makes the cumulative diff re-pair it as a delete plus an add although no commit's own diff shows the pair. 12 fixture tests cover the reproduced revert, the merge-commit resolution case, four must-stay-clean shapes, the declaration path, the rejected template placeholder, every fail-closed exit, and a replay of the real `31a51afb2` / `92ea978a4` pair that skips rather than passes when those commits are not in the clone.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
AdmBufferHip(~272 bytes) was passed by value in four__global__kernel signatures acrossadm_csf.hipandadm_cm.hipAdmBufferHip bufwithconst AdmBufferHip * __restrict__ buf_ptrin all four kernels; a device-side copy of the struct is allocated once at init viahipMalloc+hipMemcpy, stored asAdmStateHip::buf_dev, freed inclose_fex_hipAffected kernels
adm_csf.hipadm_csf_kernel_1_4AdmBufferHip buf→ pointeradm_csf.hipi4_adm_csf_kernel_1_4AdmBufferHip buf→ pointeradm_cm.hipi4_adm_cm_line_kernelAdmBufferHip buf→ pointeradm_cm.hipadm_cm_line_kernel_8AdmBufferHip buf→ pointerBuild verification
Result: 35 targets, 0 errors (host:
/opt/rocm/bin/hipcc, ROCm on CachyOS).Runtime verification status
PENDING — requires AMD GPU with ROCm runtime. Smoke command when available:
Expected: ADR-0214 places=4 parity vs CPU. Change is numerically transparent (pointer indirection only; all kernel arithmetic and load addresses are identical).
Deliverables checklist
docs/research/research-0759-hip-adm-buffer-by-pointer.md## Alternatives considered(in-kernel ldg,__constant__, host-side pointer (chosen), unmodified)AGENTS.mdinvariant:core/src/feature/hip/AGENTS.md— AdmBufferHip pass-by-pointer mandatory; future struct args follow same rulechangelog.d/perf/hip-adm-buffer-by-pointer.mddocs/rebase-notes.md— no rebase impact (all files are fork-added)docs/state.md: T-HIP-ADM-BUFFER-BY-POINTER-20260529 row addeddocs/adr/README.md: ADR-0759 row addedNo auto-merge
DRAFT — blocked on AMD hardware runtime verification before ready.
🤖 Generated with Claude Code