Repository navigation
perf(cuda): ciede 8/16bpc — __ldg() read-only cache routing (F3 fix, ADR-0762) - #102
Merged
Merged
Conversation
lusoris
marked this pull request as ready for review
May 29, 2026 09:37
…ADR-0762) F3: extract const uint8_t *__restrict__ (8bpc) and uint16_t *__restrict__ (16bpc) channel pointers from VmafPicture struct args before per-pixel body of calculate_ciede_kernel_8bpc and calculate_ciede_kernel_16bpc. Replace all 6 indexed channel reads with __ldg(&ptr[idx]) to route through L1 read-only texture cache (L2-pressure reduction at 1080p and above). Add __launch_bounds__(BLOCK_X * BLOCK_Y) register-budget hint to both kernels. Mirrors the F3 pattern from ADR-0754 (PR #93, SSIM vert_combine). VmafPicture passes void *data[3] by value, hiding the alias-free invariant from the compiler; extracting typed __restrict__ pointers makes it visible. Resolve pre-existing merge-conflict stub in integer_vif_cuda.c (inherited from commit 24bb5da): HEAD side retained (ADR-0743 comment block preserved). Correctness: CUDA vs CPU places=4 PASS, max ciede2000 diff = 0.0 on Netflix 576x324 reference pair. Build: meson + ninja (CUDA 13.3); meson test --suite=fast 55/55 OK. Deliverables: ADR-0762, AGENTS.md invariant, changelog.d/perf/cuda-ciede-ldg.md, rebase-notes entry, state.md row. no digest needed: F3 fix mirroring PR #93 pattern no per-surface docs needed: internal CUDA kernel optimization, no user-discoverable API change no ffmpeg-patches impact: no public C API change Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris
force-pushed
the
perf/cuda-ciede-ldg-20260529
branch
from
May 29, 2026 09:40
daff121 to
2fce9ab
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
State row T-HIP-ADM-ADR0759-REVERTED-2026-09-18 closed (#102 as the reverting merge), changelog fragment, and a rebase note on keeping the pointer form and checking for it after a merge. The HIP AGENTS.md now describes the code as it is: the stale "P1 known issue" section that said the struct was still passed by value is gone, and the ADR-0759 invariant names buf_dev, adm_hip_upload_buf() and the launch argument, states the no-write-after-init precondition, and records the measured effect. Research-0759 gets a dated addendum: the revert, the real struct size (328 bytes), that the CUDA twin passes AdmBufferCuda by value, the gfx1036 runtime verification and the scratch, VGPR and fps measurements. ADR-0759's body is unchanged; it is accurate again.
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
calculate_ciede_kernel_8bpcandcalculate_ciede_kernel_16bpcincore/src/feature/cuda/integer_ciede/ciede_score.cu.const uint8_t *__restrict__(8bpc) andconst uint16_t *__restrict__(16bpc) channel pointers fromVmafPicturestruct args before the per-pixel body; replace all 6 indexed channel reads with__ldg(&ptr[idx])to route through the L1 read-only texture cache.__launch_bounds__(BLOCK_X * BLOCK_Y)register-budget hint to both kernels.integer_vif_cuda.c(inherited from commit24bb5daf89): HEAD side (ADR-0743 comment block) retained.Correctness
ciede2000at places=4: PASS (max diff = 0.0 on Netflix 576×324).meson test --suite=fast: 55/55 OK (CUDA 13.3, RTX 4090).Reproducer
Deliverables checklist
## Alternatives consideredcore/src/feature/cuda/AGENTS.md—__ldg()pattern for VmafPicture channel readschangelog.d/perf/cuda-ciede-ldg.mddocs/rebase-notes.md— ADR-0762 entrydocs/state.md: T-CUDA-CIEDE-LDG-F3-20260529 row added🤖 Generated with Claude Code