Skip to content

docs(research): hardware measurement of CUDA ms_ssim_decimate + adm_cm (Research-0749 / ADR-0750) - #89

Merged
lusoris merged 1 commit into
masterfrom
docs/cuda-ms-ssim-decimate-adm-cm-measure-0749
May 28, 2026
Merged

docs(research): hardware measurement of CUDA ms_ssim_decimate + adm_cm (Research-0749 / ADR-0750)#89
lusoris merged 1 commit into
masterfrom
docs/cuda-ms-ssim-decimate-adm-cm-measure-0749

Conversation

@lusoris

@lusoris lusoris commented May 28, 2026

Copy link
Copy Markdown
Contributor

Summary

  • Hardware A/B measurement on RTX 4090 (CUDA 13.3, ncu 2026.2.0) for branch perf/cuda-ms-ssim-decimate-adm-cm-ncu-driven-20260528
  • adm_cm_line_kernel_8 __launch_bounds__: confirmed -9.3% kernel duration at 1080p, registers 114→64 (-44%). Keep.
  • ms_ssim_decimate smem tiling: measured +8-24% kernel regression (baseline L1 hit rate already 95%; cooperative-load breaks hardware prefetcher). Recommend revert.
  • End-to-end throughput: +4.8% WL1 (576p), +3.9% WL2 (1080p) — both driven by adm_cm, not ms_ssim.
  • Correctness: bit-exact at ADR-0214 places=4 (max_diff=0.0).

Test plan

  • ncu verified under --privileged container (RmProfilingAdminOnly=1 on host)
  • 3-run median end-to-end timing per variant per workload
  • Correctness A/B: float_ms_ssim_cuda + vmaf model, max_diff=0.0 for all 48f/3f

Reproducer

# see docs/research/0749-cuda-ms-ssim-decimate-adm-cm-1080p-measure.md for full commands
docker run --rm --privileged --runtime=nvidia --gpus all \
  vmaf-dev-mcp:cuda13.3 \
  -c "ncu --target-processes all -k regex:ms_ssim_decimate \
      --metrics gpu__time_duration.sum,l1tex__t_sector_hit_rate.pct \
      --csv vmaf --backend cuda -r ref.yuv -d dis.yuv ..."

Deep-dive checklist (ADR-0108)

  • Research digest: docs/research/0749-cuda-ms-ssim-decimate-adm-cm-1080p-measure.md
  • Decision matrix in ADR-0750 ## Alternatives considered
  • no AGENTS.md invariant needed: doc-only PR
  • Reproducer command in PR description above
  • Changelog fragment: changelog.d/perf/cuda-ms-ssim-adm-cm-measure-0749.md
  • docs/rebase-notes.md updated with adm_cm register budget invariant

no rebase impact: doc-only, no source changes

🤖 Generated with Claude Code

@lusoris lusoris mentioned this pull request May 28, 2026
Merged
9 tasks done
@lusoris
lusoris enabled auto-merge (squash) May 28, 2026 22:14
lusoris added a commit that referenced this pull request May 28, 2026
@lusoris @claude
perf(cuda): partial revert PR #79 — keep adm_cm __launch_bounds__, dr…
5c83151 
…op ms_ssim smem tiling

Per PR #89 hardware measurement on RTX 4090 (vmaf-dev-mcp:cuda13.3):

- REVERT: ms_ssim_decimate smem tiling.  The baseline kernel was already
  L1-resident (95% hit rate).  Cooperative load + __syncthreads() broke
  the hardware prefetcher and raised kernel duration +8–24% at all
  tested resolutions (576p and 1080p).  The non-tiled kernel is restored.

- KEEP: adm_cm_line_kernel_8 __launch_bounds__(128, 8).  Confirmed −9.3%
  kernel duration at 1080p; registers 114→64 per thread.  End-to-end
  net: +3.9–4.8% fps driven solely by this occupancy gain.

ADR-0744 updated from stub to Accepted with full revert rationale and
alternatives table.  docs/state.md row added.

Correctness: ADR-0214 places=4 parity gate confirmed bit-exact before
and after (max_diff=0.0, per ADR-0750 Research-0749).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 28, 2026
@lusoris @claude
perf(cuda): partial revert PR #79 — keep adm_cm __launch_bounds__, dr…
3009610 
…op ms_ssim smem tiling

Per PR #89 hardware measurement on RTX 4090 (vmaf-dev-mcp:cuda13.3):

- REVERT: ms_ssim_decimate smem tiling.  The baseline kernel was already
  L1-resident (95% hit rate).  Cooperative load + __syncthreads() broke
  the hardware prefetcher and raised kernel duration +8–24% at all
  tested resolutions (576p and 1080p).  The non-tiled kernel is restored.

- KEEP: adm_cm_line_kernel_8 __launch_bounds__(128, 8).  Confirmed −9.3%
  kernel duration at 1080p; registers 114→64 per thread.  End-to-end
  net: +3.9–4.8% fps driven solely by this occupancy gain.

ADR-0744 updated from stub to Accepted with full revert rationale and
alternatives table.  docs/state.md row added.

Correctness: ADR-0214 places=4 parity gate confirmed bit-exact before
and after (max_diff=0.0, per ADR-0750 Research-0749).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 28, 2026
@lusoris @claude
perf(cuda): ms_ssim_decimate smem tiling + adm_cm register reduction …
c1a8508 
…(ADR-0744) (#79)

* perf(cuda): ms_ssim_decimate smem tiling + adm_cm register reduction (ADR-0744)

Opt A: ms_ssim_decimate shared-memory tile
  Converts 81 global/L2 reads per output pixel to L1 hits via a
  (2*BLOCK_X+2*LPF_HALF+1) x (2*BLOCK_Y+2*LPF_HALF) = 41x24 float smem
  tile per CTA (3936 B). Cooperative load applies mirror_idx once in the
  load phase only; the hot 9x9 LPF convolution loop reads smem at
  tile[2*ty+kv][2*tx+ku] unconditionally. TILE_W_PAD=41 (+1 pad) prevents
  bank aliasing per the ADR-0454 / filter1d.cu convention.
  Estimated: -30 to -40% DRAM throughput; +68-93% local speedup at 1080p+.

Opt B: adm_cm_line_kernel_8 __launch_bounds__(128, 8)
  Hints ptxas to target <=64 regs/thread (65536/(8x128)) matching the
  fused scale 1-3 kernel, raising theoretical occupancy from 33% to ~67%
  on Ampere. Block size 128 (BLOCKX=32 x BLOCKY=4) is fixed by the host
  launcher in integer_adm_cuda.c.
  Estimated: +66.7% local kernel throughput at 1080p+.

Both changes are correctness-neutral; verified by ADR-0214 parity gate
at places=4. ncu measurement commands documented in Research-0744.

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

* perf(cuda): partial revert PR #79 — keep adm_cm __launch_bounds__, drop ms_ssim smem tiling

Per PR #89 hardware measurement on RTX 4090 (vmaf-dev-mcp:cuda13.3):

- REVERT: ms_ssim_decimate smem tiling.  The baseline kernel was already
  L1-resident (95% hit rate).  Cooperative load + __syncthreads() broke
  the hardware prefetcher and raised kernel duration +8–24% at all
  tested resolutions (576p and 1080p).  The non-tiled kernel is restored.

- KEEP: adm_cm_line_kernel_8 __launch_bounds__(128, 8).  Confirmed −9.3%
  kernel duration at 1080p; registers 114→64 per thread.  End-to-end
  net: +3.9–4.8% fps driven solely by this occupancy gain.

ADR-0744 updated from stub to Accepted with full revert rationale and
alternatives table.  docs/state.md row added.

Correctness: ADR-0214 places=4 parity gate confirmed bit-exact before
and after (max_diff=0.0, per ADR-0750 Research-0749).

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

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris @claude
docs(research): hardware measurement of CUDA ms_ssim_decimate + adm_c…
02ef22f 
…m (Research-0749 / ADR-0750)

Add hardware benchmark results for PR #79 (ms_ssim_decimate smem tiling +
adm_cm register reduction) on 1080p content, research digest 0749, and
ADR-0750 documenting the measurement methodology.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris
lusoris force-pushed the docs/cuda-ms-ssim-decimate-adm-cm-measure-0749 branch from 6801d45 to 02ef22f Compare May 28, 2026 22:51
@lusoris
lusoris merged commit 34a95b0 into master May 28, 2026
10 of 15 checks passed
@lusoris
lusoris deleted the docs/cuda-ms-ssim-decimate-adm-cm-measure-0749 branch May 28, 2026 22:51
@lusoris lusoris added this to the 1.0.0 — First release milestone Sep 4, 2026
@lusoris lusoris mentioned this pull request Oct 7, 2026
Merged
8 of 11 tasks
lusoris added a commit that referenced this pull request Oct 7, 2026
@lusoris
refactor(interop): re-vendor Pelorus at the commit that opens the qp-…
61c9016 
…report CSV with _wfsopen (ADR-1113) (#2424)

* refactor(interop): re-vendor Pelorus at the commit that opens the qp-report CSV with _wfsopen (ADR-1113)

VMAFx/pelorus #89 (fixing #88) moves open_utf8() from the deprecated
_wfopen() to _wfsopen(..., _SH_DENYNO), the local edit the MSVC
zero-warnings series carried in core/src/interop/pelorus_qp_report_csv.c.
PELORUS_VENDOR_SHA moves to 4aae30711c65 and --update re-renders the ten
vendored files; the mirror is byte-identical to pelorus again apart from
the banner and the include rewrite, and the drift check passes.
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