Skip to content

docs(cuda): ncu hotpath profiles for adm, motion, ssim, ms_ssim (Research-0734-0738) - #77

Merged
lusoris merged 1 commit into
masterfrom
research/cuda-other-kernels-ncu-profile-20260528
May 28, 2026
Merged

docs(cuda): ncu hotpath profiles for adm, motion, ssim, ms_ssim (Research-0734-0738)#77
lusoris merged 1 commit into
masterfrom
research/cuda-other-kernels-ncu-profile-20260528

Conversation

@lusoris

@lusoris lusoris commented May 28, 2026

Copy link
Copy Markdown
Contributor

Summary

  • ncu --set basic profiles collected on RTX 4090 (sm_89, CUDA 13.3) for all remaining CUDA metric families at 576x324 Netflix golden pair
  • Research-0734 (ADM), Research-0735 (motion), Research-0736 (SSIM), Research-0737 (MS-SSIM), Research-0738 (cross-metric summary)
  • No source code changes: research-only PR with docs, changelog fragment, rebase note, and state.md entry

Per-metric diagnosis (one line each)

Metric Diagnosis
ADM Launch-starved across all 5 kernels (< 1 wave / 128 SMs); adm_cm_line_kernel_8 also register-limited (114 regs, 33% theoretical occ)
Motion ~62% achieved occupancy (~5.9 waves), best performer; DRAM at 10-12%, no secondary bottleneck
SSIM vert_combine DRAM-bound at 55.8%; P0 bug: integer_ssim_score.cu missing extern "C" crashes int64 SSIM CUDA path at runtime
MS-SSIM Severely launch-starved at pyramid levels 2-4 (0.06-0.25 waves); no shared-memory staging in 9-tap LPF decimate

Top-3 cross-metric optimization candidates

  1. Fix extern "C" in integer_ssim_score.cu (P0 correctness, < 5 LOC) — unblocks int64 SSIM CUDA path which currently crashes with CUDA_ERROR_NOT_FOUND. File: core/src/feature/cuda/integer_ssim/integer_ssim_score.cu.
  2. Shared-memory tiling for ms_ssim_decimate — 81 global reads per output pixel with no tile staging. Estimated +30-40% DRAM throughput reduction at 1080p+. ncu estimates 68-93% local speedup from grid coverage alone.
  3. Register reduction in adm_cm_line_kernel_8 — 114 regs/thread caps theoretical occupancy at 33%. Reducing to 56-64 regs lifts theoretical occupancy to 75%, estimated +20-30% throughput at >= 1080p.

Reproducer commands

ADM:

docker run --rm --gpus all --privileged --entrypoint bash \
  -v <worktree>:/workspace -v <repo>/python:/workspace/python:ro \
  -w /workspace/core vmaf-dev-mcp:cuda13.3 -c \
  'ncu -k "regex:adm_cm_line|i4_adm_cm_line|adm_csf_den|adm_dwt2" --set basic --launch-count 4 \
   build-ncu/tools/vmaf \
     --reference ../python/test/resource/yuv/src01_hrc00_576x324.yuv \
     --distorted ../python/test/resource/yuv/src01_hrc01_576x324.yuv \
     --width 576 --height 324 --pixel_format 420 --bitdepth 8 \
     --feature adm --backend cuda -o /dev/null'

Motion / SSIM / MS-SSIM: see per-metric reproducer in Research-0735/0736/0737.

Deep-dive deliverables checklist

  • (a) Research digests: docs/research/0734-0738-*.md
  • (b) no ADR needed: research-only, no code change proposed
  • (c) no rebase-sensitive invariants: docs/rebase-notes.md entry added
  • (d) Reproducer: ncu commands in each digest
  • (e) changelog.d/changed/cuda-other-kernels-ncu-profile.md
  • (f) docs/rebase-notes.md entry appended

Per-surface docs sentinel

no user-discoverable surface change: research

State.md

State entry prepended for T-CUDA-HOTPATH-PROFILES-ADM-MOTION-SSIM-2026-05-28.

🤖 Generated with Claude Code

@lusoris
lusoris enabled auto-merge (squash) May 28, 2026 20:41
@lusoris lusoris mentioned this pull request May 28, 2026
Merged
8 of 11 tasks
@lusoris @claude
docs(cuda): ncu hotpath profiles for adm, motion, ssim, ms_ssim kernels
811f6d9 
Research-0734 to 0738: ncu --set basic profiles on RTX 4090 (sm_89, CUDA 13.3)
for all remaining CUDA metric families at 576x324 Netflix golden pair.

ADM (Research-0734): all 5 kernels launch-starved (< 1 wave / 128 SMs).
adm_cm_line_kernel_8 register-limited at 114 regs/thread (33% theoretical occ).

Motion (Research-0735): 62-64% achieved occupancy (~5.9 waves). DRAM 10-12%.

SSIM (Research-0736): calculate_ssim_vert_combine DRAM-bound at 55.8%.
P0 bug found: integer_ssim_score.cu missing extern "C" causes
CUDA_ERROR_NOT_FOUND at runtime for --feature ssim --backend cuda.

MS-SSIM (Research-0737): ms_ssim_decimate severely launch-starved at pyramid
levels 2-4 (0.06-0.25 waves); no shared-memory staging in 9-tap LPF.

Cross-metric summary (Research-0738) ranks top-3 candidates:
1. Fix extern "C" in integer_ssim_score.cu (P0 correctness, < 5 LOC).
2. Shared-memory tiling for ms_ssim_decimate (+30-40% DRAM at 1080p+).
3. Register reduction in adm_cm_line_kernel_8 (33% -> 75% theoretical occ).

no user-discoverable surface change: research

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris
lusoris force-pushed the research/cuda-other-kernels-ncu-profile-20260528 branch from 086a9b8 to 811f6d9 Compare May 28, 2026 21:34
@lusoris
lusoris merged commit f10e79b into master May 28, 2026
15 of 20 checks passed
@lusoris
lusoris deleted the research/cuda-other-kernels-ncu-profile-20260528 branch May 28, 2026 21:34
lusoris added a commit that referenced this pull request May 28, 2026
@lusoris @claude
fix(cuda)!: wrap all host-looked-up __global__ kernels in extern "C" …
be587bb 
…(P0 silent-corruption sweep)

Full audit of all 24 .cu kernel files in core/src/feature/cuda/ and
core/src/cuda/ against every cuModuleGetFunction host-side lookup.

Finding: integer_ssim/integer_ssim_score.cu is the only file with
__global__ kernels referenced by cuModuleGetFunction but not wrapped
in extern "C". nvcc compiles .cu as C++ by default; the three entry
points (integer_ssim_horiz_8bpc, integer_ssim_horiz_16bpc,
integer_ssim_vert_combine) received C++ name-mangling, causing the
driver to return CUDA_ERROR_NOT_FOUND for all three lookups.
init_fex_cuda returned -EINVAL, silently disabling --feature ssim
--backend cuda since the file was introduced (PR #77 fixed the
analogous break in ssim_score.cu).

Fix: add extern "C" { } around all three __global__ entry points.
Device-only helpers (__device__ static, __constant__) are not affected.
All 23 other kernel files are confirmed safe.

Deliverables (ADR-0108):
  (a) Research digest: docs/research/research-0747-cuda-extern-c-sweep.md
  (b) ADR-0747: docs/adr/0747-cuda-extern-c-sweep.md
  (c) AGENTS.md invariant + CI script: scripts/dev/check-cuda-extern-c.sh
  (d) Reproducer: bash scripts/dev/check-cuda-extern-c.sh
  (e) Changelog: changelog.d/fixed/cuda-extern-c-sweep.md
  (f) Rebase notes: docs/rebase-notes.md

Affected features: ssim (--backend cuda).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 28, 2026
@lusoris @claude
fix(cuda)!: wrap all host-looked-up __global__ kernels in extern "C" …
396e8ac 
…(P0 silent-corruption sweep) (#80)

Full audit of all 24 .cu kernel files in core/src/feature/cuda/ and
core/src/cuda/ against every cuModuleGetFunction host-side lookup.

Finding: integer_ssim/integer_ssim_score.cu is the only file with
__global__ kernels referenced by cuModuleGetFunction but not wrapped
in extern "C". nvcc compiles .cu as C++ by default; the three entry
points (integer_ssim_horiz_8bpc, integer_ssim_horiz_16bpc,
integer_ssim_vert_combine) received C++ name-mangling, causing the
driver to return CUDA_ERROR_NOT_FOUND for all three lookups.
init_fex_cuda returned -EINVAL, silently disabling --feature ssim
--backend cuda since the file was introduced (PR #77 fixed the
analogous break in ssim_score.cu).

Fix: add extern "C" { } around all three __global__ entry points.
Device-only helpers (__device__ static, __constant__) are not affected.
All 23 other kernel files are confirmed safe.

Deliverables (ADR-0108):
  (a) Research digest: docs/research/research-0747-cuda-extern-c-sweep.md
  (b) ADR-0747: docs/adr/0747-cuda-extern-c-sweep.md
  (c) AGENTS.md invariant + CI script: scripts/dev/check-cuda-extern-c.sh
  (d) Reproducer: bash scripts/dev/check-cuda-extern-c.sh
  (e) Changelog: changelog.d/fixed/cuda-extern-c-sweep.md
  (f) Rebase notes: docs/rebase-notes.md

Affected features: ssim (--backend cuda).

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris lusoris mentioned this pull request May 29, 2026
Merged
6 tasks done
lusoris added a commit that referenced this pull request May 29, 2026
@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 added a commit that referenced this pull request May 29, 2026
@lusoris @claude
docs(cuda): ADR-0756 F3 struct-by-value kernel audit — 20 kernels, to…
950d2af 
…p-5 dispatch (#96)

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 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