Skip to content

perf(cuda): skip the far LOD pass when the far band is empty - #5

Open
bam93 wants to merge 3 commits into
cross-platformfrom
perf/lod-skip-empty-far-pass
Open

perf(cuda): skip the far LOD pass when the far band is empty#5
bam93 wants to merge 3 commits into
cross-platformfrom
perf/lod-skip-empty-far-pass

Conversation

@bam93

@bam93 bam93 commented Jun 29, 2026

Copy link
Copy Markdown
Collaborator

What

API_computeSES_lod runs computeSlicedSES for both the near and far bands unconditionally. When the molecule lies entirely in the near band, the far pass still pays a full pass of setup (grid build + atom thrust::sort) + the slab-enumeration loop + aabbInFrustum tests, only to produce 0 geometry. The per-slab visible-only cull (mode==0) already skips the heavy refine/MC kernels, but not this setup/enumeration cost.

This adds a pass-level pre-check: test the model's padded world AABB (getMinMax, expanded by maxAtomRad + 2*probe to match the grid extent, so the test is conservative and never skips a band that could contribute) against farPlanes using the existing aabbInFrustum. The far computeSlicedSES is skipped when the far band is provably empty.

Correctness — output-identical (verified byte-exact)

The skip only removes a pass that produced nothing (when the far AABB is outside the far frustum, the near pass already covered every visible slab). Verified on 1CRN and 3EAM, both the empty-far and the real-split cases: NVert/NTri are identical with vs without the skip.

case structure NVert (no-skip) NVert (skip)
empty far band 3EAM 391977 391977
real split 3EAM 485153 485153

Performance (dgx, RTX A1000, reso 0.5, coarseMul 2)

case structure no-skip with-skip delta
empty far band 3EAM (12625 atoms) 129 ms 115 ms ~11%
empty far band 1CRN (327 atoms) 4.0 ms 3.6 ms ~noise
real split 3EAM 146 ms 150 ms unchanged (skip correctly does not fire)

A modest win, larger on big multi-slab systems; the skip never fires when the far band has real work. Tested via a small harness calling API_computeSES_lod directly (the CLI doesn't exercise the LOD API); harness not included in the PR.

Provenance

This mirrors the same empty-far-pass skip applied in the UnityMol Unity-ComputeShader port (where it was ~2×, since that path re-ran the whole GPU dispatch; CUDA's in-loop cull already makes the empty pass cheaper, hence the smaller win here).

🤖 Generated with Claude Code

API_computeSES_lod ran computeSlicedSES for BOTH near and far bands unconditionally. When
the molecule lies entirely in the near band the far pass still pays a full pass of setup
(grid build + atom thrust::sort) + the slab-enumeration loop + aabbInFrustum tests, only to
produce 0 geometry — the per-slab visible-only cull skips the heavy refine/MC kernels but not
this setup/enumeration cost.

Add a pass-level pre-check: test the model's padded world AABB (getMinMax, expanded by
maxAtomRad + 2*probe to match the grid extent so the test is conservative) against farPlanes
with the existing aabbInFrustum; skip the far computeSlicedSES when the far band is provably
empty. Output-identical (the near pass already covers every visible slab when the far AABB is
outside) — verified byte-identical NVert/NTri with vs without the skip on 1CRN and 3EAM, both
the empty-far and real-split cases.

Measured (dgx RTX A1000, reso 0.5, coarseMul 2): empty-far band 3EAM (12625 atoms) 129->115 ms
(~11%); 1CRN small ~noise; real-split case unchanged (skip correctly does not fire). Modest
win, larger on big multi-slab systems; the skip never fires when the far band has work.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Copilot AI review requested due to automatic review settings June 29, 2026 22:03

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR optimizes the CUDA LOD SES pipeline by avoiding the far-band computeSlicedSES pass when it can be proven to generate no geometry, reducing unnecessary setup and slab-enumeration overhead for near-only cases.

Changes:

  • Added a pass-level frustum pre-check using a conservatively padded model AABB against farPlanes.
  • Skips the far computeSlicedSES call entirely when the padded AABB is outside the far frustum (i.e., far band is empty).

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread CudaSurf.cu Outdated
Comment on lines +1093 to +1094
// this setup+enumeration cost. Test the model AABB (padded by the same maxAtomRad + 2*probe the grid
// uses, so the test is conservative — never skips a band that could contribute) against farPlanes
…xact)

Per PR review: the comment claimed the AABB is padded by the SAME amount the grid uses, but
the grid extent comes from computeMaxDist/originGridNeighbor (per-side maxAtomRad+probeRadius)
and doesn't equal maxAtomRad+2*probeRadius exactly. Reword to state the pad is >= the grid's
per-side reach — a deliberate over-cover that never under-covers, so the skip never drops a
band that could contribute. Comment-only; the pad value and logic are unchanged.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@bam93

bam93 commented Jun 29, 2026

Copy link
Copy Markdown
Collaborator Author

Good catch — fixed in a7ad2e4. The comment now states the relationship accurately: the per-side pad maxAtomRad + 2*probeRadius is >= the grid's actual reach (the SES grid origin offsets each side by maxAtomRad + probeRadius via originGridNeighbor, and computeMaxDist adds 2*maxAtomRad + 4*probeRadius along the longest axis), i.e. a deliberate over-cover that never under-covers — so the skip can never drop a band that could contribute. Comment-only change; the pad value and logic are unchanged, so the verified parity/perf results still hold.

@bam93
bam93 requested a review from nezix June 29, 2026 22:12
…with Metal)

The per-vertex nearest-atom search used fast_distance (sqrt) per candidate while the Metal
twin (nearestAtomSorted) already uses dot/squared distance. sqrt is monotonic so the argmin
is identical without it. Switch to sqr_distance (the squared helper already present, matching
the commented-out intent). Output byte-identical: verified md5-identical OBJ on 1CRN/1UBQ/3EAM
(baseline build vs this change, same -v 0.5). Brings the two backends into agreement;
correctness-neutral (fast-math already softened the sqrt, so this is a tidy-up, not a perf win).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
@bam93

bam93 commented Jun 30, 2026

Copy link
Copy Markdown
Collaborator Author

Added a small related parity tidy-up (8eadb93): computeClosestAtom now uses squared distance for the per-vertex nearest-atom argmin (sqrt is monotonic, so the argmin is identical without the per-candidate sqrt) — matching the Metal port which already does this. Byte-identical verified: md5-identical OBJ on 1CRN/1UBQ/3EAM (baseline build vs change, same -v 0.5). Correctness-neutral; brings the two backends into agreement. Happy to split it into its own PR if you'd prefer.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants