Add ipc::cuda::LBVH: a GPU broad phase sharing the CPU LBVH implementation - #260
Draft
zfergus wants to merge 10 commits into
Draft
Add ipc::cuda::LBVH: a GPU broad phase sharing the CPU LBVH implementation#260zfergus wants to merge 10 commits into
zfergus wants to merge 10 commits into
Conversation
Introduce src/ipc/utils/cuda/device_utils.cuh, the header every ipc::cuda translation unit needs before it can launch anything: - IPC_TOOLKIT_CUDA_CHECK, which turns a cudaError_t into a std::runtime_error naming the file and line. - KERNEL_BLOCK_SIZE and kernel_grid_size(), the single definition of the launch geometry, so the block size is not repeated per call site. - global_dof_index(), mirroring the index math of local_gradient_to_global_gradient() for device-side gradient scatter. - A compile-time guard rejecting compute capability < 6.0, where atomicAdd(double*, double) does not exist. Include <Eigen/Core> directly: global_dof_index() compares VERTEX_DERIVATIVE_LAYOUT against Eigen::RowMajor, and config.hpp deliberately defines its own Eigen-free layout constants rather than pulling in Eigen, so the header would otherwise only compile when an includer happened to have included Eigen first. The header is CUDA-only and included from .cu files exclusively; it is wired in under IPC_TOOLKIT_WITH_CUDA so a non-CUDA build never sees it.
A first-class GPU counterpart to ipc::LBVH (not a CPU-upload adapter): builds vertex/edge/face AABBs and their BVHs entirely on the device (Morton codes + Apetrei 2014 single-pass bottom-up construction, reusing the 32-byte ipc::LBVH::Node layout for host validation/interop), then runs candidate detection with the BVH descent and mesh-connectivity (shared-vertex) exclusion both on the device. The user vertex filter is honored on the device for the common accept-all case (new CollisionFilter::accepts_all()); a non-trivial filter falls back to a host pass over the device-emitted, connectivity-filtered candidates. Either path matches the CPU ipc::LBVH's candidate set exactly. Adds DeviceCandidateView + detect_*_candidates_device() so candidates can stay device-resident for a future GPU-native pipeline (e.g. device Additive CCD) instead of always materializing to host vectors. Supporting changes: ipc::math::morton_2D/3D and expand_bits_1/2 are now IPC_TOOLKIT_HOST_DEVICE so the device Morton codes reuse the exact CPU implementation; the Morton-normalization reciprocal is now precomputed once per build and multiplied per box instead of divided (CPU and GPU changed identically so their Morton codes stay bit-matched to each other). Validation: build + detect + custom-filter-fallback GPU-run-validated on an RTX 3070 (artemis): 150517 assertions across 3 test cases, plus exact candidate-set parity against the CPU LBVH for all 6 candidate types. Benchmarked against the CPU LBVH (edge-edge detection): 1.1-1.7x faster on every real mesh tested except a trivial two-cube case. The Morton reciprocal-multiply optimization and code cleanup (Eigen::Array3d in place of a hand-rolled Vec3d, .min()/.max() in place of manual fminf/fmaxf loops) landed after artemis went offline and are Docker-compile-validated only; pending a GPU re-run. Not yet done: ipc::cuda::LBVH is not registered in BroadPhaseMethod / create_broad_phase (deferred until the device-resident candidate path is consumed by something), and the connectivity/user-filter split does not yet support device-side patch/connected-component filters (would need a label-data CollisionFilter descriptor). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
detect_*_candidates sized its output buffer from a fresh max(1024, 8 * n_source_leaves) guess on every single call, with nothing remembering what a prior call actually needed. Checking the real candidate counts on the CPU LBVH (proven exactly equal to the GPU's) showed 5 of 7 benchmarked meshes overflow that guess by up to 17x, so nearly every real mesh silently paid for two full kernel dispatches on every call: one that discovers the buffer is too small, then a full re-traversal at the corrected size. Add a predicted_capacity field to LBVH::Impl::DeviceCandidates (one per candidate type) that persists the largest count ever observed and seeds the next call's guess. It is deliberately not reset by clear(), since build() calls clear() every timestep and the hint must survive that or it never helps; it only ever grows for the object's lifetime, mirroring the predicted_*_candidates_size pattern already used by the (Slang) vulkan branch's LBVH. Also add a logger().warn() on overflow, matching that same branch, so a retry is no longer silent. Validated on artemis (RTX 3070): [lbvh][cuda] unchanged at 150517 assertions. Re-benchmarked detect_edge_edge_candidates against the CPU LBVH: the 2 meshes that never overflowed are byte-identical before/after as expected; the 5 that did are 13-29% faster (e.g. Rod-Twist 15.8ms -> 11.2ms, Puffer-Ball 1.097s -> 0.912s), widening the GPU's margin over the CPU across the board (e.g. Rod-Twist 1.23x -> 1.73x, Puffer-Ball 1.47x -> 1.76x faster than CPU). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Add BroadPhaseMethod::LBVH_CUDA, appended after SWEEP_AND_TINIEST_QUEUE to keep the existing enum values stable (the "Create broad phase" test casts consecutive integers to BroadPhaseMethod). The factory case mirrors the SWEEP_AND_TINIEST_QUEUE case exactly: returns ipc::cuda::LBVH under IPC_TOOLKIT_WITH_CUDA, otherwise throws with a message naming the CMake option to enable. Not added to tests/src/tests/utils.cpp's broad_phases() / BroadPhaseGenerator (used by most generic cross-broad-phase comparison tests): several of those exercise 2D meshes, and ipc::cuda::LBVH::build() currently throws on non-3D input (v1 scope), unlike SweepAndTiniestQueue which silently upgrades 2D to 3D via to_X3d() before building. Adding it there would break those tests immediately; left for a follow-up if 2D parity is wanted. Validated: host (non-CUDA) build passes "Create broad phase" (5 assertions, count unchanged). Artemis (CUDA, RTX 3070): same test passes with the bumped count (7 assertions); [lbvh][cuda] suite unaffected (150517 assertions, no regression). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
build() previously hard-coded dim = 3 and threw for non-3-column input.
The vertex upload also unconditionally read 3 components per vertex,
which would read out of bounds on a 2-column matrix -- dimension
support was blocked below the validation check, not just at it.
Mirror the CPU ipc::LBVH's actual semantics instead of the simpler
upgrade-to-3D-via-to_X3d approach SweepAndTiniestQueue uses. The key
subtlety: ipc::AABB's constructor zero-initializes its 3-wide array and
only assigns the first `dim` components from the already-inflated
input, so a 2D box's z bound is an exact, uninflated 0.0 -- not
nextafter(0 +/- inflation_radius, ...). build_vertex_boxes_{static,
dynamic}_kernel now take dim and, for components past it, write a hard
0.0 instead of running the inflation formula, matching that exactly.
Vertex upload now sizes to dim * n instead of a fixed 3 * n. All three
build() overloads relax to assert(dim == 2 || dim == 3) (matching the
CPU's debug-only assert, not a throw) and set dim from the real input.
The Morton-code kernel's dim == 2 branch already existed (copied from
the CPU when first written) and needed no change; the edge/face box
union kernels and the Apetrei hierarchy build are dim-agnostic and
untouched.
Add "GPU LBVH 2D build and detect" using the same mesh-2D CSV data as
the CPU's own 2D test: checks vertex/edge BVH structural and root-AABB
parity, plus exact detect_edge_vertex_candidates parity against the CPU
LBVH (the only candidate type meaningful in 2D).
Validated on artemis (RTX 3070): [lbvh][cuda] now 152785 assertions
across 4 test cases (was 150517/3) -- the existing 3D paths are
unregressed and the new 2D path matches the CPU exactly.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
nextafter(double, float) does not exist on the device, so use nextafter(double, double) with a constexpr for positive and negative infinity. INFINITY is a float macro, so nextafter(double, INFINITY) resolves to the host-only std::nextafter<double, float> promotion template instead of CUDA's __device__ nextafter(double, double).
prim_shares_vertex used two runtime-indexed index_t[3] locals. Runtime indexing forces them into local memory, and on sm_120 ptxas sized traverse_kernel's frame at 0x110 bytes while basing those arrays at frame+0x100 -- 16 bytes of room for 24 bytes of object, on top of the 0x100-byte traversal stack based at frame+0. Writes landed on stack[0..1] and destroyed the INVALID_POINTER sentinel the descent loop terminates on, so the traversal popped past the bottom of the stack and read stack[-1]. The result was cudaErrorIllegalAddress, which surfaces as an apparent hang: the driver spins in the candidate-counter readback, and the poisoned context makes every later GPU test look stuck too. Hold the vertex ids in scalars instead, filling unused slots from slot 0 so every comparison stays well defined. The frame drops to 0x100 (exactly the stack) and local traffic to the 3 stack accesses. Scope of the miscompile: sm_120 only. sm_75/86/89 allocate 0x120 as expected, identically with -rdc=true and -rdc=false, and the driver's own JIT (CUDA 13.3) reproduces the 272 vs 288 split, so it is neither an -rdc nor a 12.8 artifact. Building the unfixed source as compute_89 PTX and JIT-ing onto the sm_120 device passes clean. A provably bounded index does not help, so this is not licensed by the latent UB. Tests: [gpu] ~[!benchmark] passes (158705 assertions, 28 cases) and compute-sanitizer memcheck reports 0 errors on [lbvh][gpu]; both faulted before. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
ipc::LBVH and ipc::cuda::LBVH held line-for-line ports of the same algorithm, with agreement asserted only in comments and checked only by tests. Hoist the pieces that need no parallel abstraction into shared host/device code so the agreement is structural. New shared code: - ipc::morton_code() computes a box's Morton code from its center, normalizing by a domain whose width is passed as a reciprocal so the host and device multiply rather than divide. - ipc::count_leading_zeros() and ipc::morton_common_prefix() replace the per-platform CLZ dispatch and the duplicate-code fallback rule (Apetrei 2014's delta). - ipc::details::can_*_collide() hold the five mesh-connectivity filters. These were duplicated three times, not two: ipc::BroadPhase carries the same logic over AABB::vertex_ids. LBVH::Node's is_inner/is_leaf/is_valid/intersects are now IPC_TOOLKIT_HOST_DEVICE, so the traversal kernel calls the same predicates as the CPU instead of open-coding is_inner_marker == 0 and reimplementing the AABB overlap test. 167 duplicated lines collapse into 111 shared ones. Node::intersects() also generates better SASS than the hand-expanded aabb_intersects it replaces: traverse_kernel drops 320 -> 304 instructions, 33 -> 29 global loads, 16 -> 12 float compares and 5 -> 3 reconvergence pairs, with the register count (35-36) and the 0x100 local frame unchanged. Holding that frame is a hard requirement here -- the sm_120 miscompile fixed in c03e546 was frame-size sensitive. Morton codes are unchanged bit-for-bit. check_tree only compares root AABBs within 1e-4, so the suite cannot establish this; a standalone harness comparing the shared function against both prior forms over 200,000 random 2D and 3D cases found zero differences, and the compute_morton_codes_kernel opcode histogram is unchanged. Tested: full suite (4,348,911 assertions in 353 cases), compute-sanitizer memcheck on [lbvh][gpu] with 0 errors, clang-format and clang-tidy clean. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The Apetrei 2014 bottom-up build and the BVH descent were line-for-line ports between ipc::LBVH and ipc::cuda::LBVH. Hoist both into shared host/device code, leaving each platform only what it genuinely owns: the parallel launch, the sort, and a small policy per difference. ipc::details::build_hierarchy_from_leaf() takes the sorted-code accessor (the host stores an array of structs, the device a flat array) and the atomic arrival gate. ipc::details::traverse_lbvh() takes what to do on an overlap, which is the whole of the host/device difference there: the host filters and appends to a std::vector, the device filters against the mesh connectivity and appends through an atomic counter. Also shared: set_inflated_aabb(), init_leaf_node(), delta(), is_left_child(), swap_root_to_zero() and patch_left_pointer(). LBVH::ConstructionInfo is now a template over its counter type, so the host uses std::atomic<int> and the device a plain int, from one layout. 434 lines leave the two implementations for 197 lines of shared code. Fixes a latent race in the device build. The arrival gate had a __threadfence() on the release side but none on the acquire side, then read the sibling's child pointer, range endpoint and rightmost leaf with ordinary loads, which may be served from a stale L1 on another SM. The shared gate's contract requires both halves, and the device policy now fences after an increment that returns nonzero. The kernel's SASS gains exactly one MEMBAR.SC.GPU, giving MEMBAR.SC.GPU / ATOMG.E.ADD.STRONG.GPU / MEMBAR.SC.GPU with the paired CCTL.IVALL that invalidates L1. This would have corrupted internal-node AABBs and rightmost[] without breaking the tree structure, so check_tree's structural checks could not have caught it. The single-leaf build case is now explicit on both sides. The host previously relied on writing nodes[0].left = 0 over the lone leaf's primitive_id, which was only correct because a one-box sort always yields box_id 0. traverse_kernel's SASS is bit-identical after the change -- the templated descent and its lambda inline away completely -- and every kernel's register count is unchanged from before this series. The 0x100 frame that the sm_120 miscompile in c03e546 turned on is preserved; that was the acceptance gate for touching this kernel at all. Adds coverage for single-primitive BVHs, which no existing mesh reaches. Only face-vertex and edge-face put a one-node BVH in the traversal target position, so the test builds one face and one disjoint edge and checks both against BruteForce, and the device against the host. Verified non-vacuous by mutation: suppressing the emit in the single-node branch fails 4 of its assertions. Tested: full suite (4,348,947 assertions in 354 cases), compute-sanitizer memcheck on [lbvh][gpu] with 0 errors, clang-format and clang-tidy clean. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
- Warning falsly triggers on our own dependencies because spdlog.cmake takes precedence over downstream spdlog.cmake scripts.
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #260 +/- ##
==========================================
+ Coverage 96.74% 96.82% +0.07%
==========================================
Files 191 193 +2
Lines 17292 17306 +14
Branches 933 935 +2
==========================================
+ Hits 16730 16757 +27
+ Misses 562 549 -13
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Adds
ipc::cuda::LBVH, a GPU-native LBVH broad phase, and factors the build and traversal it shares withipc::LBVHinto a single implementation used by both.The shared code lives in
broad_phase/details/, parameterized on the few things that genuinely differ between host and device:lbvh_build.hpp— the Apetrei [2014] bottom-up build, parameterized on how sorted Morton codes are read and how a node's two children rendezvous (std::atomicpost-increment vs.atomicAddbetween two__threadfence()s).lbvh_traverse.hpp— the Karras-style explicit-stack descent, parameterized on what to do with an overlapping leaf. That functor turned out to be the entire host/device difference.connectivity_filters.hpp— thecan_*_collidepredicates, previously duplicated inBroadPhase,ipc::LBVH, andipc::cuda::LBVH.math/morton.hppgains the shared Morton code construction, count-leading-zeros dispatch, and common-prefix length.LBVH::Node's predicates are now host/device, andLBVH::ConstructionInfois templated on its counter sostd::atomic<int>andintshare one layout.Impacts:
ipc::LBVHloses 336 lines by adopting the shared build and descent.traverse_kernelat a0x100frame, 3 local-memory accesses, and 35 registers.API changes
BroadPhaseMethod::LBVH_CUDA, appended so existing enumerator values are unchanged. Without CUDA it throws, matchingSWEEP_AND_TINIEST_QUEUE.ipc::cuda::LBVH(broad_phase/cuda/lbvh.hpp), aBroadPhasesubclass supporting 2D and 3D.CollisionFilter::accepts_all()— true only for the default filter, letting a GPU broad phase skip host-side filtering when the device-emitted set is already exact.LBVH::ConstructionInfomoves from a private type to a public template.Known gaps
Left as follow-ups, hence draft:
ipc::cuda::LBVH, unlike every other broad phase (including CUDASweepAndTiniestQueue).traverse_lbvh_simdremains a third, separate traversal by design; a scalar shared function cannot express it.Type of change
How Has This Been Tested?
test_gpu_lbvh.cu: GPU candidate sets compared against the CPU LBVH in 2D and 3D, including a custom collision filter."LBVH single-primitive trees"case intest_lbvh.cpp, covering the one-node target BVH that face-vertex and edge-face detection produce. Verified non-vacuous by mutation — suppressing the single-node emit fails 4 assertions.compute-sanitizer --tool memcheckon[lbvh][gpu]: 0 errors.clang-format: 0 replacements across all 17 touched C++/CUDA files.Test Configuration:
-arch=nativeon sm_120 (RTX 5080), ReleaseChecklist