Skip to content

docs(cuda): ADR-0756 F3 struct-by-value audit — 20 kernels, top-5 dispatch - #96

Merged
lusoris merged 1 commit into
masterfrom
research/cuda-f3-struct-by-value-audit-20260529
May 29, 2026
Merged

docs(cuda): ADR-0756 F3 struct-by-value audit — 20 kernels, top-5 dispatch#96
lusoris merged 1 commit into
masterfrom
research/cuda-f3-struct-by-value-audit-20260529

Conversation

@lusoris

@lusoris lusoris commented May 29, 2026

Copy link
Copy Markdown
Contributor

Summary

Six deep-dive deliverables

  • Research digest: docs/research/research-0756-cuda-f3-struct-by-value-audit.md
  • Decision matrix in ADR-0756 ## Alternatives considered (in-kernel __ldg vs host-side extraction vs skip)
  • AGENTS.md invariant note: F3 pattern section added to core/src/feature/cuda/AGENTS.md
  • Reproducer: ncu -k "regex:ms_ssim_vert_lcs" --set basic --launch-count 10 ... (in research digest)
  • Changelog fragment: changelog.d/perf/cuda-f3-struct-by-value-audit.md
  • Rebase-notes: docs/rebase-notes.md — no rebase impact: documentation-only PR

Per-surface docs

research only — no user-discoverable CLI/API surface changed by this PR.

state.md

Row added to Recently closed: T-CUDA-F3-STRUCT-BY-VALUE-AUDIT-2026-05-29.

Smoke test

# Confirm no CUDA source files were modified (audit-only)
git diff origin/master -- '*.cu' '*.cuh' '*.c' '*.h'
# Expected: empty diff

Affected kernel inventory (summary)

Severity Kernel File
HIGH ms_ssim_vert_lcs integer_ms_ssim/ms_ssim_score.cu:136
HIGH ms_ssim_horiz integer_ms_ssim/ms_ssim_score.cu:100
HIGH ms_ssim_decimate integer_ms_ssim/ms_ssim_score.cu:70
MEDIUM-HIGH calculate_ciede_kernel_8bpc/16bpc integer_ciede/ciede_score.cu
MEDIUM adm_decouple_kernel + adm_decouple_s123_kernel integer_adm/adm_decouple.cu
MEDIUM psnr_hvs integer_psnr_hvs/psnr_hvs_score.cu
LOW calculate_psnr_kernel_*, calculate_moment_kernel_* psnr_score.cu, moment_score.cu
ALREADY FIXED calculate_ssim_vert_combine PR #93 / ADR-0754

🤖 Generated with Claude Code

Drafted by routines

@lusoris
lusoris enabled auto-merge (squash) May 29, 2026 09:17
@lusoris lusoris mentioned this pull request May 29, 2026
Merged
6 tasks done
@lusoris @claude
docs(cuda): ADR-0756 F3 struct-by-value kernel audit — 20 kernels, to…
2e4dac5 
…p-5 dispatch

Fork-wide inventory of every __global__ kernel accepting VmafCudaBuffer,
VmafPicture, or AdmBufferCuda by value (the F3 pattern deferred in PR #93).
The hidden CUdeviceptr prevents ptxas from emitting ld.global.nc / __restrict__
alias-free analysis on inner-loop reads.

Results: 20 kernel variants across 8 metric families affected; 5 high-severity
(hot inner loop, no smem tile, no existing pointer extraction); 1 already fixed
(calculate_ssim_vert_combine, PR #93 / ADR-0754, measured -4.2% at 1080p).

Top-5 dispatch PRs defined in Research-0756, ranked by measured DRAM throughput
from PR #77 ncu profiles:
  1. ms_ssim_vert_lcs  — DRAM pattern identical to ssim_vert_combine (~55% est)
  2. ms_ssim_horiz     — 7 VmafCudaBuffer args, K=11 horiz loop
  3. ciede 8/16bpc     — 6 channel reads/pixel from VmafPicture.data[]
  4. adm_decouple      — AdmBufferCuda 6 band pointers, 7.8% DRAM @576p
  5. psnr_hvs          — VmafCudaBuffer tile load only, smem-resident post-load

Per-surface docs: research only (no user-discoverable CLI/API surface changed).
No rebase impact: documentation-only PR.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris
lusoris force-pushed the research/cuda-f3-struct-by-value-audit-20260529 branch from b7b3619 to 2e4dac5 Compare May 29, 2026 09:31
@lusoris
lusoris merged commit 950d2af into master May 29, 2026
11 of 19 checks passed
@lusoris
lusoris deleted the research/cuda-f3-struct-by-value-audit-20260529 branch May 29, 2026 09:31
@lusoris lusoris mentioned this pull request May 29, 2026
Merged
9 tasks done
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
perf(cuda): ms_ssim_vert_lcs + ms_ssim_horiz — __ldg + __launch_bound…
898b8e1 
…s (F3 fix #2, ADR-0757)

Apply the F3 __ldg() / __restrict__ pointer-extraction pattern (first used in
ADR-0754 for calculate_ssim_vert_combine, PR #93) to the two top ms_ssim audit
candidates identified in PR #96:

- ms_ssim_vert_lcs: extract 5 const float *__restrict__ pointers before K=11
  loop; use __ldg() on all 5x11 = 55 inner-loop loads.
- ms_ssim_horiz: extract 2 const float *__restrict__ pointers before K=11
  loop; use __ldg() on all 2x11 = 22 inner-loop loads.
- Both kernels: add __launch_bounds__(128) (actual launch is 16x8=128 threads).

LDG.E.CONSTANT confirmed in sm_89 SASS via cuobjdump. Build: clean nvcc
compile (zero warnings). Predicted -4 to -6% kernel duration at 1080p
(memory-bound regime where combined intermediate footprint exceeds L2 capacity).

Six deep-dive deliverables: ADR-0757, AGENTS.md invariant extended,
changelog.d/perf/cuda-ms-ssim-vert-lcs-horiz-ldg.md, rebase-notes.md entry,
state.md row, no user-discoverable surface change (perf-only).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
perf(cuda): ms_ssim_vert_lcs + ms_ssim_horiz — __ldg + __launch_bound…
f2e2c03 
…s (F3 fix #2, ADR-0757) (#99)

Apply the F3 __ldg() / __restrict__ pointer-extraction pattern (first used in
ADR-0754 for calculate_ssim_vert_combine, PR #93) to the two top ms_ssim audit
candidates identified in PR #96:

- ms_ssim_vert_lcs: extract 5 const float *__restrict__ pointers before K=11
  loop; use __ldg() on all 5x11 = 55 inner-loop loads.
- ms_ssim_horiz: extract 2 const float *__restrict__ pointers before K=11
  loop; use __ldg() on all 2x11 = 22 inner-loop loads.
- Both kernels: add __launch_bounds__(128) (actual launch is 16x8=128 threads).

LDG.E.CONSTANT confirmed in sm_89 SASS via cuobjdump. Build: clean nvcc
compile (zero warnings). Predicted -4 to -6% kernel duration at 1080p
(memory-bound regime where combined intermediate footprint exceeds L2 capacity).

Six deep-dive deliverables: ADR-0757, AGENTS.md invariant extended,
changelog.d/perf/cuda-ms-ssim-vert-lcs-horiz-ldg.md, rebase-notes.md entry,
state.md row, no user-discoverable surface change (perf-only).

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris lusoris mentioned this pull request May 29, 2026
Closed
6 of 9 tasks
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
perf(cuda): psnr_hvs — __ldg() + __restrict__ + __launch_bounds__(64)…
28a9eb3 
… (ADR-0764)

Apply the F3 struct-by-value fix (PR #96 candidate #5) to the psnr_hvs
CUDA kernel in core/src/feature/cuda/integer_psnr_hvs/psnr_hvs_score.cu:

- Extract `const float *__restrict__ ref_buf` and `dist_buf` from the
  VmafCudaBuffer struct args before the cooperative 64-thread tile load.
- Apply `__ldg()` to both per-thread element reads, routing them through
  the L1 read-only texture cache (LDG.E.CONSTANT in SASS).
- Add `__launch_bounds__(64)` matching the actual 8×8 block dispatch.

No arithmetic change; bit-identical scores (ADR-0214 places=4). Predicted
-3 to -5% kernel duration at >=1080p (mirrors ADR-0754 1080p -4.2% result).

Six deep-dive deliverables: ADR-0764, Research-0764, changelog.d/perf,
AGENTS.md invariant extended, rebase-notes.md entry, state.md row.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris lusoris mentioned this pull request May 29, 2026
Closed
6 of 11 tasks
lusoris added a commit that referenced this pull request Jun 3, 2026
@lusoris @claude
docs(process): per-surface doc compliance audit — last 30 PRs (ADR-0848)
f8e157e 
Audits 30 most recent merged commits (PRs #96–#174) for CLAUDE §12 r10
compliance: user-discoverable surface change must ship human-readable docs
in the same PR.

Score: 22/30 N/A (no surface change), 5/8 with surface changes compliant
(62.5 %). Three confirmed gaps:

- Gap A (HIGH): PR #47 (ADR-0726, Vulkan drop) — docs/backends/vulkan/overview.md,
  docs/metrics/features.md, and docs/development/build-flags.md still describe
  Vulkan as an operative backend. Tracked as T-DOC-VULKAN-STALE-POST-ADR0726.
- Gap B (MEDIUM): PR #87 (ADR-0749, VmafLegacyQualityRunner sunset) — no
  docs/development/deprecations.md entry and no migration note in
  docs/usage/python.md. Tracked as T-DOC-LEGACY-RUNNER-MISSING-DEPRECATION.
- Gap C (LOW): PR #135 (log format standardization) — removes "Error: " prefix
  from CUDA error messages with no doc note.

Three follow-up issues proposed (Issue A, B, C in Research-0848). No code
changes in this PR; audit-only per task brief.

Six deliverables (ADR-0108):
(1) Research-0848: docs/research/research-0848-per-surface-doc-compliance-audit-20260529.md
(2) Decision matrix in ADR-0848 §Alternatives considered
(3) no rebase-sensitive invariants
(4) Reproducer: git log --oneline origin/master | head -30 (auditable from commit list)
(5) changelog.d/changed/per-surface-doc-compliance-audit.md
(6) docs/rebase-notes.md entry added

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request Jun 3, 2026
@lusoris @claude
docs: hygiene bundle (#127 NOLINT audit + #216 doc-compliance audit + #…
9479bd9 
…291 state.md drift sweep + #233 changelog concat fix) (#527)

* docs(lint): NOLINT cluster audit and refactor plan (ADR-0780)

Swept all 218 NOLINT annotations in core/src/ for clusters of five or more
identical suppressions per file. Identified five clusters (71 annotations):

- 21 bare performance-no-int-to-ptr in GPU slab allocators — ADR-0278
  non-compliant (no citations); scheduled for SLAB_FIELD macro (PR B).
- 12 bugprone-implicit-widening in SYCL stride arithmetic — scheduled for
  explicit (ptrdiff_t) casts that eliminate the suppression entirely (PR A).
- 14 misc-const-correctness in SYCL atomic_ref loops — fold into existing
  NOLINTBEGIN/NOLINTEND block (PR C).
- 13 bare readability-function-size in integer_adm.c — ADR-0278 non-compliant;
  scheduled for NOLINTBEGIN block consolidation (PR C).
- 11 readability-function-size on SYCL kernel entry-points — load-bearing
  per ADR-0141; no change planned.

Deliverables: research digest, ADR-0780 (Proposed), changelog fragment.
Follow-up refactor PRs A–C are independent.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* docs(process): per-surface doc compliance audit — last 30 PRs (ADR-0848)

Audits 30 most recent merged commits (PRs #96–#174) for CLAUDE §12 r10
compliance: user-discoverable surface change must ship human-readable docs
in the same PR.

Score: 22/30 N/A (no surface change), 5/8 with surface changes compliant
(62.5 %). Three confirmed gaps:

- Gap A (HIGH): PR #47 (ADR-0726, Vulkan drop) — docs/backends/vulkan/overview.md,
  docs/metrics/features.md, and docs/development/build-flags.md still describe
  Vulkan as an operative backend. Tracked as T-DOC-VULKAN-STALE-POST-ADR0726.
- Gap B (MEDIUM): PR #87 (ADR-0749, VmafLegacyQualityRunner sunset) — no
  docs/development/deprecations.md entry and no migration note in
  docs/usage/python.md. Tracked as T-DOC-LEGACY-RUNNER-MISSING-DEPRECATION.
- Gap C (LOW): PR #135 (log format standardization) — removes "Error: " prefix
  from CUDA error messages with no doc note.

Three follow-up issues proposed (Issue A, B, C in Research-0848). No code
changes in this PR; audit-only per task brief.

Six deliverables (ADR-0108):
(1) Research-0848: docs/research/research-0848-per-surface-doc-compliance-audit-20260529.md
(2) Decision matrix in ADR-0848 §Alternatives considered
(3) no rebase-sensitive invariants
(4) Reproducer: git log --oneline origin/master | head -30 (auditable from commit list)
(5) changelog.d/changed/per-surface-doc-compliance-audit.md
(6) docs/rebase-notes.md entry added

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* docs(state): close T-LEGACY-RUNNER-ANSNR-BROKEN + T-LEGACY-RUNNER-STUB-MISSING

Both rows referenced classes / imports that no longer exist on master:

1. T-LEGACY-RUNNER-ANSNR-BROKEN — AnsnrFeatureExtractor was deleted
   by PR #283 (merged at faede7a). The class is absent from
   compat/python-vmaf/core/feature_extractor.py on master tip
   d45d503. The Netflix golden assertions still cover the
   integer-path VMAF score (Rule #1 preserved).

2. T-LEGACY-RUNNER-STUB-MISSING-2026-05-29 — VmafLegacyQualityRunner
   import-time failure: the class was removed in ADR-0749 / PR #87
   and python/test/quality_runner_test.py was updated by the
   ADR-0749 sunset PR to drop the import (only a removed-comment
   placeholder at line 49 remains).

Moved both rows from Open to Recently closed with verified-on-master
reproducer commands.

No code changes — documentation cleanup only.

Deliverables (ADR-0108):
- [x] **Research digest**: no digest needed: state.md hygiene
- [x] **Decision matrix**: no alternatives: only-one-way fix
- [x] **AGENTS.md invariant**: no rebase-sensitive invariants
- [x] **Reproducer**: `git show origin/master:compat/python-vmaf/core/feature_extractor.py | grep -c 'class AnsnrFeatureExtractor'` returns 0; `git show origin/master:python/test/quality_runner_test.py | grep -c '^from.*VmafLegacy'` returns 0.
- [x] **Changelog**: changelog.d/changed/state-md-drift-sweep-20260530.md
- [x] **Rebase notes**: no rebase impact: docs only

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

* docs(state): migrate 3 Vulkan rows to Recently closed (ADR-0726 supersession)

ADR-0726 (Vulkan backend dropped 2026-05-28, PR #47) structurally closes
three long-standing Vulkan Open rows by removing the affected code path
entirely:

- T-VK-1.4-BUMP — NVIDIA driver 595.71+ FP-contraction regression
- T-VK-CIEDE-F32-F64 — NVIDIA Vulkan f32/f64 ciede precision gap
- T-VK-VIF-1.4-RESIDUAL-ARC — Intel Arc A380 vif residual on Mesa-ANV

The entire `core/src/vulkan/`, `core/src/feature/vulkan/`, and
`core/include/libvmaf/libvmaf_vulkan.h` surface no longer exists on
master, so each row now has a verified empty-path reproducer. Native
CUDA / HIP / SYCL backends cover every vendor formerly served by Vulkan
(see ADR-0726 §Context).

Combined with the legacy-runner closures already in this PR
(T-LEGACY-RUNNER-ANSNR-BROKEN + T-LEGACY-RUNNER-STUB-MISSING-2026-05-29),
this brings the Open section from 14 → 11 rows.

Row counts (verified): Open 11 + Deferred 4 + Recently closed 137 +
Confirmed not-affected 1 = 153 total T- rows in tree.

* chore(changelog): fix concat awk boundary + consolidate perf/ fragments (ADR-0221)

Fix `concat-changelog-fragments.sh` --check/--write awk block-boundary: the
previous `/^## [^[]/` pattern terminated the Unreleased block on any `## Heading`
embedded inside a fragment (e.g. `## Added`, `## [perf] …`). Switch to
`/^## \[(Unreleased|[0-9])/` which matches only versioned-release headers and
the `[Unreleased]` sentinel, making `--check` deterministic. Tightened the
`--write` awk pass with the same fix.

Consolidate 32 fragments from non-standard `changelog.d/perf/` (27) and
`changelog.d/performance/` (5) into the recognised `changed/` section with
`perf-` filename prefix per ADR-0221 KaC convention.

Remove 3 duplicate stubs: `ort-run-stack-arrays-f3b.md`, `adm-pnorm-deferred-comment.md`,
`fixed/vmaf-tune-predictor-directory-corpus.md` (richer versions retained).
Rename `changed/fr-regressor-v3-namespace.md` → `…-adr-appendix.md` to
resolve filename collision with `added/fr-regressor-v3-namespace.md`.

Regenerate CHANGELOG.md Unreleased block via `--write`; `--check` exits 0.

no rebase impact: changelog.d/ + scripts/release/ only, no upstream C touched.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* chore(bundle): add changelog fragment for docs/hygiene bundle PR

Adds changelog.d/changed/bundle-docs-hygiene.md summarising the
four source PRs: #127 (NOLINT audit), #216 (doc compliance audit),
#291 (state.md drift sweep), #233 (changelog concat fix).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-authored-by: Lusoris <lusoris@pm.me>
@lusoris lusoris mentioned this pull request Jun 3, 2026
Merged
6 of 9 tasks
lusoris added a commit that referenced this pull request Jun 3, 2026
@lusoris @claude
perf(cuda): psnr_hvs — __ldg() + __restrict__ + __launch_bounds__(64)…
4c2c3f0 
… (ADR-0764) (#563)

Apply the F3 struct-by-value fix (PR #96 candidate #5) to the psnr_hvs
CUDA kernel in core/src/feature/cuda/integer_psnr_hvs/psnr_hvs_score.cu:

- Extract `const float *__restrict__ ref_buf` and `dist_buf` from the
  VmafCudaBuffer struct args before the cooperative 64-thread tile load.
- Apply `__ldg()` to both per-thread element reads, routing them through
  the L1 read-only texture cache (LDG.E.CONSTANT in SASS).
- Add `__launch_bounds__(64)` matching the actual 8x8 block dispatch.

No arithmetic change; bit-identical scores (ADR-0214 places=4). Predicted
-3 to -5% kernel duration at >=1080p (mirrors ADR-0754 1080p -4.2% result).

Six deep-dive deliverables: ADR-0764, Research-0764, changelog.d/perf,
AGENTS.md invariant extended, rebase-notes.md entry, state.md row.

Rebased from #107 onto current master.

Co-authored-by: Lusoris <lusoris@pm.me>
Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris lusoris added this to the 1.0.0 — First release milestone Sep 4, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

1.0.0 — First release

Development

Successfully merging this pull request may close these issues.

1 participant

@lusoris