From ba8359da8025d743cc95a4662c0a1fdb392e04f9 Mon Sep 17 00:00:00 2001 From: Marc Baaden Date: Tue, 30 Jun 2026 00:02:22 +0200 Subject: [PATCH 1/3] perf(cuda): skip the far LOD pass when the far band is empty MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- CudaSurf.cu | 40 ++++++++++++++++++++++++++++++++-------- 1 file changed, 32 insertions(+), 8 deletions(-) diff --git a/CudaSurf.cu b/CudaSurf.cu index 7592679..5e79617 100644 --- a/CudaSurf.cu +++ b/CudaSurf.cu @@ -1086,14 +1086,38 @@ API void API_computeSES_lod(float resoSES, float coarseMul, std::vector nearMeshes = computeSlicedSES(atomPos, atomRad, N, resoSES, doSmoothing, &vNear); // Pass 2 — far band, coarse voxel. - ViewParams vFar; - vFar.enabled = (farPlanes != NULL); - vFar.mode = 0; - vFar.camPos = camPos; - vFar.cb = slabCb; - vFar.userData = userData; - if (farPlanes != NULL) for (int p = 0; p < 24; p++) vFar.planes[p] = farPlanes[p]; - std::vector farMeshes = computeSlicedSES(atomPos, atomRad, N, resoSES * coarseMul, doSmoothing, &vFar); + // FAST-SKIP an entirely-empty far band: if the model's (padded) world AABB is fully OUTSIDE the + // far frustum, no slab can fall in the far band, so the far pass would run a whole computeSlicedSES + // (grid build + atom thrust::sort + the full slab-enumeration loop) only to produce 0 geometry. + // The per-slab visible-only cull inside computeSlicedSES skips the heavy refine/MC kernels but NOT + // 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 + // with the existing aabbInFrustum. Output-identical: when the far AABB is outside, the near pass + // (which already covered every visible slab) is the whole surface. (Mirrors the UnityMol port's + // LOD empty-far-pass skip; the per-slab empty-cell early-out CUDA already has, this is pass-level.) + bool runFar = true; + if (farPlanes != NULL) + { + float3 aabbMin, aabbMax; float maxAtomRad; + getMinMax(atomPos, atomRad, N, &aabbMin, &aabbMax, &maxAtomRad); + float pad = maxAtomRad + 2.0f * probeRadius; + aabbMin.x -= pad; aabbMin.y -= pad; aabbMin.z -= pad; + aabbMax.x += pad; aabbMax.y += pad; aabbMax.z += pad; + runFar = aabbInFrustum(farPlanes, aabbMin, aabbMax); + } + + std::vector farMeshes; + if (runFar) + { + ViewParams vFar; + vFar.enabled = (farPlanes != NULL); + vFar.mode = 0; + vFar.camPos = camPos; + vFar.cb = slabCb; + vFar.userData = userData; + if (farPlanes != NULL) for (int p = 0; p < 24; p++) vFar.planes[p] = farPlanes[p]; + farMeshes = computeSlicedSES(atomPos, atomRad, N, resoSES * coarseMul, doSmoothing, &vFar); + } // Consolidate near+far into the global arrays for callers that fetch one mesh at the end. for (size_t i = 0; i < farMeshes.size(); i++) nearMeshes.push_back(farMeshes[i]); From a7ad2e422f0267f8e74fc5413d262cac9edae106 Mon Sep 17 00:00:00 2001 From: Marc Baaden Date: Tue, 30 Jun 2026 00:08:02 +0200 Subject: [PATCH 2/3] docs(cuda): correct the LOD-skip padding comment (conservative, not exact) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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) --- CudaSurf.cu | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/CudaSurf.cu b/CudaSurf.cu index 5e79617..fbe1e56 100644 --- a/CudaSurf.cu +++ b/CudaSurf.cu @@ -1090,17 +1090,20 @@ API void API_computeSES_lod(float resoSES, float coarseMul, // far frustum, no slab can fall in the far band, so the far pass would run a whole computeSlicedSES // (grid build + atom thrust::sort + the full slab-enumeration loop) only to produce 0 geometry. // The per-slab visible-only cull inside computeSlicedSES skips the heavy refine/MC kernels but NOT - // 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 - // with the existing aabbInFrustum. Output-identical: when the far AABB is outside, the near pass - // (which already covered every visible slab) is the whole surface. (Mirrors the UnityMol port's - // LOD empty-far-pass skip; the per-slab empty-cell early-out CUDA already has, this is pass-level.) + // this setup+enumeration cost. Test the model AABB against farPlanes with the existing aabbInFrustum, + // padded CONSERVATIVELY so the test never skips a band that could contribute: the SES grid origin + // offsets each side by maxAtomRad + probeRadius (see originGridNeighbor) and the grid extent adds + // 2*maxAtomRad + 4*probeRadius along the longest axis (computeMaxDist), so a per-side pad of + // maxAtomRad + 2*probeRadius is >= the grid's actual reach — it OVER-covers, never under-covers. + // Output-identical: when the far AABB is outside, the near pass (which already covered every visible + // slab) is the whole surface. (Mirrors the UnityMol port's LOD empty-far-pass skip; the per-slab + // empty-cell early-out CUDA already has, this is the pass-level one.) bool runFar = true; if (farPlanes != NULL) { float3 aabbMin, aabbMax; float maxAtomRad; getMinMax(atomPos, atomRad, N, &aabbMin, &aabbMax, &maxAtomRad); - float pad = maxAtomRad + 2.0f * probeRadius; + float pad = maxAtomRad + 2.0f * probeRadius; // >= the grid's per-side reach (conservative) aabbMin.x -= pad; aabbMin.y -= pad; aabbMin.z -= pad; aabbMax.x += pad; aabbMax.y += pad; aabbMax.z += pad; runFar = aabbInFrustum(farPlanes, aabbMin, aabbMax); From 8eadb93039db151c3566d3e5dce11ad4116c20a6 Mon Sep 17 00:00:00 2001 From: Marc Baaden Date: Tue, 30 Jun 2026 11:00:07 +0200 Subject: [PATCH 3/3] cleanup(cuda): squared-distance argmin in computeClosestAtom (parity 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) --- Kernels.cu | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/Kernels.cu b/Kernels.cu index d791ee3..249c423 100644 --- a/Kernels.cu +++ b/Kernels.cu @@ -550,7 +550,8 @@ inline __host__ __device__ int computeClosestAtom(float3 vert, int3 id3DNeigh, i int closestId = -1; - float minD = 999999.0f; + float minDsq = 999999.0f * 999999.0f; // squared-distance argmin (sqrt is monotonic; no per-candidate + // sqrt needed). Matches the Metal nearestAtomSorted (dot-based). int3 curgridId; // #pragma unroll for (int x = -1; x <= 1; x++) { @@ -570,10 +571,9 @@ inline __host__ __device__ int computeClosestAtom(float3 vert, int3 id3DNeigh, i for (int id = idStart; id < idStop; id++) { float4 xyzr = sorted_xyzr[id]; float3 pos = make_float3(xyzr.x, xyzr.y, xyzr.z); - // float d = sqr_distance(pos, vert); - float d = fast_distance(pos, vert); - if(d < minD){ - minD = d; + float dsq = sqr_distance(pos, vert); // squared distance: drop the per-candidate sqrt + if(dsq < minDsq){ + minDsq = dsq; closestId = id; } }