ipc-sim · zfergus · Sep 15, 2026 · Sep 9, 2026 · Jul 25, 2026 · Jul 25, 2026
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
@@ -11,4 +11,4 @@ repos:
           - -i
           - --style=file
           - --verbose
-        files: '\.(c|cc|cpp|h|hpp|tpp|cxx|hh|inl|ipp)$'
+        files: '\.(c|cc|cpp|cu|cuh|h|hpp|tpp|cxx|hh|inl|ipp)$'
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -344,14 +344,13 @@ if(IPC_TOOLKIT_WITH_SIMD)
     # compiled twice and the nvcc pass owns the float and double
     # instantiations (see ipc_toolkit_shared_device_sources.cmake), so without
     # this they would be the only scalar code in the library built without
-    # SIMD support. Hand the flags to the host compiler nvcc drives rather
-    # than to nvcc itself.
-    set(SIMD_CUDA_HOST_FLAGS "")
-    foreach(simd_flag IN LISTS SIMD_CXX_FLAGS)
-      list(APPEND SIMD_CUDA_HOST_FLAGS "-Xcompiler=${simd_flag}")
-    endforeach()
-    target_compile_options(ipc_toolkit PUBLIC
-      "$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:${SIMD_CUDA_HOST_FLAGS}>")
+    # SIMD support. ipc_toolkit_filter_nvcc_flags() hands the flags to the
+    # host compiler nvcc drives (via -Xcompiler) rather than to nvcc itself,
+    # keeping only those that pass compiles with.
+    include(ipc_toolkit_filter_flags)
+    set(SIMD_CUDA_HOST_FLAGS ${SIMD_CXX_FLAGS})
+    ipc_toolkit_filter_nvcc_flags(SIMD_CUDA_HOST_FLAGS)
+    target_compile_options(ipc_toolkit PUBLIC ${SIMD_CUDA_HOST_FLAGS})
   endif()
 
   # Link against cross-platform xsimd library

diff --git a/cmake/ipc_toolkit/ipc_toolkit_filter_flags.cmake b/cmake/ipc_toolkit/ipc_toolkit_filter_flags.cmake
@@ -22,4 +22,27 @@ function(ipc_toolkit_filter_flags flags)
     endif()
   endforeach()
   set(${flags} ${output_flags} PARENT_SCOPE)
-endfunction()
+endfunction()
+
+# The nvcc counterpart of ipc_toolkit_filter_flags(): keep the flags of `flags`
+# that nvcc's host compiler accepts, checked by actually compiling with nvcc and
+# `-Xcompiler=<flag>`, and wrap each so it applies only to CUDA sources compiled
+# by nvcc. `-Xcompiler` is required: nvcc parses some host flags itself with a
+# different meaning (`-Werror` takes nvcc's own diagnostic names, `-march=...`
+# is read as an input file), so a bare host flag on the nvcc command line is
+# unsafe. Requires the CUDA language to be enabled.
+function(ipc_toolkit_filter_nvcc_flags flags)
+  include(CheckCompilerFlag)
+  set(output_flags)
+  foreach(FLAG IN ITEMS ${${flags}})
+    string(REPLACE "=" "-" FLAG_VAR "${FLAG}")
+    if(NOT DEFINED IS_SUPPORTED_NVCC_HOST_${FLAG_VAR})
+      check_compiler_flag(CUDA "-Xcompiler=${FLAG}" IS_SUPPORTED_NVCC_HOST_${FLAG_VAR})
+    endif()
+    if(IS_SUPPORTED_NVCC_HOST_${FLAG_VAR})
+      list(APPEND output_flags
+        "$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:-Xcompiler=${FLAG}>")
+    endif()
+  endforeach()
+  set(${flags} ${output_flags} PARENT_SCOPE)
+endfunction()
diff --git a/cmake/ipc_toolkit/ipc_toolkit_warnings.cmake b/cmake/ipc_toolkit/ipc_toolkit_warnings.cmake
@@ -195,3 +195,54 @@ add_library(ipc::toolkit::warnings ALIAS ipc_toolkit_warnings)
 include(ipc_toolkit_filter_flags)
 ipc_toolkit_filter_flags(IPC_TOOLKIT_WARNING_FLAGS)
 target_compile_options(ipc_toolkit_warnings INTERFACE ${IPC_TOOLKIT_WARNING_FLAGS})
+
+# nvcc forwards none of the flags above to the host compiler it drives, so a
+# .cu is otherwise compiled with no warnings at all. Hand the host pass the
+# flags that catch bugs in OUR code (the device pass has no use for them).
+#
+# The full set is deliberately not forwarded: nvcc's generated host code and the
+# CUDA headers are not clean under -Wpedantic, -Wold-style-cast or -Wsign-promo
+# (thousands of hits in its stubs and in crt/device_functions.hpp), and its
+# rewriting of aggregate initializers trips -Werror=missing-braces.
+#
+# ipc_toolkit_filter_nvcc_flags() checks each flag against nvcc and scopes it
+# to CUDA sources compiled by nvcc (clang as the CUDA compiler has no -Xcompiler
+# and takes the C++ flags above directly).
+if(IPC_TOOLKIT_WITH_CUDA AND NOT MSVC)
+  set(IPC_TOOLKIT_NVCC_HOST_WARNING_FLAGS
+    -Wall
+    -Wextra
+    -Wshadow
+    -Woverloaded-virtual
+    -Wuninitialized
+    -Wcast-qual
+    -Wpointer-arith
+    -Werror=return-type
+    -Werror=non-virtual-dtor
+    -Werror=delete-non-virtual-dtor
+    -Wno-unused-parameter
+    -Wno-sign-compare
+    -Wno-unknown-pragmas # nvcc's own pragmas (unroll) reach the host pass
+  )
+  ipc_toolkit_filter_nvcc_flags(IPC_TOOLKIT_NVCC_HOST_WARNING_FLAGS)
+  target_compile_options(ipc_toolkit_warnings INTERFACE ${IPC_TOOLKIT_NVCC_HOST_WARNING_FLAGS})
+endif()
+
+# The device pass has no -Wall; nvcc's own diagnostics are few, and these are
+# the ones with a bug-finding record here. They go to nvcc directly (not via
+# -Xcompiler): nvcc's -Werror takes nvcc's own diagnostic names.
+#  * all-warnings: the front end's own warnings (e.g. "variable used before its
+#    value is set") become errors.
+#  * cross-execution-space-call: calling a host-only function from device code
+#    is an error, the typical mistake in code shared via IPC_TOOLKIT_HOST_DEVICE.
+# Not used: -Xptxas -warn-spills. Under separable compilation the shared device
+# sources are relocatable __device__ functions that follow the ABI, and ptxas
+# reports their callee-saved register traffic as spills (64 hits in the barrier
+# functions alone), so it is noise here; check kernels with --resource-usage on
+# a non-rdc compile instead. -warn-lmem-usage likewise: the traversal stack is
+# intentional local memory.
+if(IPC_TOOLKIT_WITH_CUDA)
+  target_compile_options(ipc_toolkit_warnings INTERFACE
+    "$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:SHELL:-Werror all-warnings>"
+    "$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CUDA_COMPILER_ID:NVIDIA>>:SHELL:-Werror cross-execution-space-call>")
+endif()
diff --git a/cmake/recipes/spdlog.cmake b/cmake/recipes/spdlog.cmake
@@ -1,18 +1,6 @@
 # spdlog (https://github.com/gabime/spdlog)
 # License: MIT
 if(TARGET spdlog::spdlog)
-    # Someone else created the target, so the fmt patch below never runs. That
-    # is fine without CUDA, but the patch is what makes the bundled fmt
-    # compile under nvcc at all, so warn rather than fail at the first .cu.
-    if(IPC_TOOLKIT_WITH_CUDA)
-        message(WARNING
-            "spdlog::spdlog was provided by an enclosing project, so IPC "
-            "Toolkit's cmake/patches/fmt-nvcc-compat.patch was not applied. "
-            "The bundled fmt does not compile under nvcc unpatched: its "
-            "literal-encoding probe misfires and a char32_t table uses hex "
-            "escapes with the high bit set. Apply the same patch to your "
-            "spdlog, or let IPC Toolkit fetch its own.")
-    endif()
     return()
 endif()
 

diff --git a/docs/source/Doxyfile b/docs/source/Doxyfile
@@ -2315,7 +2315,7 @@ INCLUDE_FILE_PATTERNS  =
 # recursively expanded use the := operator instead of the = operator.
 # This tag requires that the tag ENABLE_PREPROCESSING is set to YES.
 
-PREDEFINED             = IPC_TOOLKIT_WITH_INEXACT_CCD IPC_TOOLKIT_WITH_ROBIN_MAP IPC_TOOLKIT_WITH_ABSEIL IPC_TOOLKIT_WITH_FILIB IPC_TOOLKIT_WITH_MESHFEM_SPARSE
+PREDEFINED             = IPC_TOOLKIT_WITH_CUDA IPC_TOOLKIT_WITH_INEXACT_CCD IPC_TOOLKIT_WITH_ROBIN_MAP IPC_TOOLKIT_WITH_ABSEIL IPC_TOOLKIT_WITH_FILIB IPC_TOOLKIT_WITH_MESHFEM_SPARSE
 
 # If the MACRO_EXPANSION and EXPAND_ONLY_PREDEF tags are set to YES then this
 # tag can be used to specify a list of macro names that should be expanded. The

diff --git a/docs/source/about/release_notes.rst b/docs/source/about/release_notes.rst
@@ -25,10 +25,17 @@ New Features |:rocket:|
 
 - Expose the intersection coordinates of an edge–triangle intersection through a new :cpp:func:`ipc::edge_triangle_intersection` overload, which reports the barycentric coordinates :math:`(u, v)` on the triangle and the parameter :math:`t` along the edge (`#245 <https://github.com/ipc-sim/ipc-toolkit/pull/245>`_).
 - Add :cpp:func:`ipc::CollisionMesh::face_normals`, computing the unit normal of each face for a given set of vertex positions (3D only) (`#245 <https://github.com/ipc-sim/ipc-toolkit/pull/245>`_).
+- Add :cpp:class:`ipc::cuda::LBVH`, a GPU broad phase that builds the vertex/edge/face AABBs and BVHs and runs the traversal and shared-vertex filtering on the device, producing the same candidates as :cpp:class:`ipc::LBVH` for any vertex filter (`#260 <https://github.com/ipc-sim/ipc-toolkit/pull/260>`_). Requires ``IPC_TOOLKIT_WITH_CUDA``; selectable through ``BroadPhaseMethod::LBVH_CUDA``.
+
+  - The Apetrei :cite:p:`Apetrei2014FastAS` bottom-up build, the BVH descent, and the shared-vertex exclusion are one ``ipc::details`` implementation shared by the CPU and CUDA broad phases; each platform supplies only its parallel launch, its sort, and its atomics.
+  - ``detect_*_candidates_device()`` return a view of the candidate pairs left on the device for a GPU-native pipeline.
 
 API Changes |:wrench:|
 ~~~~~~~~~~~~~~~~~~~~~~
 
+- ``BroadPhase::detect_*_candidates()`` now uniformly **clear** their output vector first on every broad phase, so it holds exactly that detection's result (`#260 <https://github.com/ipc-sim/ipc-toolkit/pull/260>`_). Previously half the implementations overwrote and half appended; all in-library callers pass an empty vector, so their results are unchanged.
+- Add :cpp:func:`ipc::CollisionFilter::accepts_all`, true for a filter that holds no predicate (the default), so a broad phase can skip per-pair filtering entirely (`#260 <https://github.com/ipc-sim/ipc-toolkit/pull/260>`_). Composing with an accept-all filter now short-circuits: ``f | accept_all`` is accept-all and ``f & accept_all`` is ``f``.
+- Add ``BroadPhaseMethod::NUM_BROAD_PHASE_METHODS`` as a sentinel for the number of methods (`#260 <https://github.com/ipc-sim/ipc-toolkit/pull/260>`_).
 - Update Tight Inclusion from ``1.0.6`` to ``1.1.0`` (`#248 <https://github.com/ipc-sim/ipc-toolkit/pull/248>`_).
 
   - Adds a ``BUCKET_DEPTH_FIRST_SEARCH`` root-finding method, which upstream makes the default for ``edgeEdgeCCD`` and ``vertexFaceCCD``.
@@ -186,6 +193,7 @@ Python |:snake:|
 
 - Validate preconditions in the bindings instead of relying on the C++ ``assert``\ s, which are compiled out under ``NDEBUG`` and would let a release build silently accept a bad value (`#247 <https://github.com/ipc-sim/ipc-toolkit/pull/247>`_). ``BarrierPotential`` now raises ``ValueError`` for a non-positive or NaN ``dhat``/``stiffness`` and for a null barrier.
 - Bind ``edge_triangle_intersection()``, returning an ``(intersects, u, v, t)`` tuple since Python has no out-parameters, and ``CollisionMesh.face_normals()``, returning an (#F × 3) array to match the other per-element accessors (`#245 <https://github.com/ipc-sim/ipc-toolkit/pull/245>`_). ``face_normals()`` raises ``ValueError`` on a 2D mesh rather than invoking undefined behavior.
+- Add the ``ipctk.cuda`` submodule, mirroring the C++ ``ipc::cuda`` namespace, with ``ipctk.cuda.LBVH`` (``ipc::cuda::LBVH``) in CUDA builds (`#260 <https://github.com/ipc-sim/ipc-toolkit/pull/260>`_).
 
 Documentation
 ~~~~~~~~~~~~~

diff --git a/docs/source/cpp-api/broad_phase.rst b/docs/source/cpp-api/broad_phase.rst
@@ -40,7 +40,18 @@ Sweep and Prune
 Sweep and Tiniest Queue
 -----------------------
 
-.. .. doxygenclass:: ipc::SweepAndTiniestQueueGPU
+Requires CUDA (``IPC_TOOLKIT_WITH_CUDA``).
+
+.. doxygenclass:: ipc::SweepAndTiniestQueue
+    :allow-dot-graphs:
+
+LBVH (CUDA)
+-----------
+
+Requires CUDA (``IPC_TOOLKIT_WITH_CUDA``).
+
+.. doxygenclass:: ipc::cuda::LBVH
+    :allow-dot-graphs:
 
 AABB
 ----

diff --git a/docs/source/python-api/broad_phase.rst b/docs/source/python-api/broad_phase.rst
@@ -46,8 +46,24 @@ Sweep and Prune
 Sweep and Tiniest Queue
 -----------------------
 
-.. .. autoclass:: ipctk.SweepAndTiniestQueueGPU
-..     :members:
+``ipctk.SweepAndTiniestQueue`` is available only when ``ipctk`` is built with
+CUDA (``IPC_TOOLKIT_WITH_CUDA``), which the documentation build is not.
+
+.. .. autoclass:: ipctk.SweepAndTiniestQueue
+..
+..     .. autoclasstoc::
+
+LBVH (CUDA)
+-----------
+
+``ipctk.cuda.LBVH`` is the GPU counterpart of ``ipctk.LBVH`` (C++:
+``ipc::cuda::LBVH``). The ``ipctk.cuda`` submodule mirrors the C++ ``ipc::cuda``
+namespace; its classes exist only when ``ipctk`` is built with CUDA
+(``IPC_TOOLKIT_WITH_CUDA``), which the documentation build is not.
+
+.. .. autoclass:: ipctk.cuda.LBVH
+..
+..     .. autoclasstoc::
 
 AABB
 ----

diff --git a/docs/source/tutorials/getting_started.rst b/docs/source/tutorials/getting_started.rst
@@ -539,7 +539,7 @@ The ``Candidates`` class represents the culled set of candidate pairs and is bui
                 collision_mesh, vertices_t0, vertices_t1,
                 broad_phase=ipctk.LBVH())
 
-Possible values for ``broad_phase`` are: ``BruteForce`` (parallel brute force culling), ``HashGrid``, ``SpatialHash`` (implementation from the original IPC codebase), ``LBVH`` (CPU implementation of :cite:t:`Karras2012HPG` using TBB), ``SweepAndPrune`` (a.k.a. Sort-and-Sweep from :cite:t:`Baraff1992PhD`), or ``SweepAndTiniestQueue`` (method of :cite:t:`Belgrod2023Time`; requires CUDA). The default is ``LBVH``.
+Possible values for ``broad_phase`` are: ``BruteForce`` (parallel brute force culling), ``HashGrid``, ``SpatialHash`` (implementation from the original IPC codebase), ``LBVH`` (CPU implementation of :cite:t:`Karras2012HPG` using TBB), ``SweepAndPrune`` (a.k.a. Sort-and-Sweep from :cite:t:`Baraff1992PhD`), ``SweepAndTiniestQueue`` (method of :cite:t:`Belgrod2023Time`; requires CUDA), or ``cuda.LBVH`` (``ipc::cuda::LBVH`` in C++; the GPU counterpart of ``LBVH``, building and traversing the same hierarchy on the device; requires CUDA). The default is ``LBVH``.
 
 Narrow-Phase
 ^^^^^^^^^^^^

diff --git a/python/src/bindings.cpp b/python/src/bindings.cpp
@@ -32,6 +32,12 @@ PYBIND11_MODULE(ipctk, m)
     define_sweep_and_tiniest_queue(m);
     define_voxel_size_heuristic(m);
 
+    // GPU implementations, mirroring the C++ ipc::cuda namespace. The submodule
+    // always exists; its classes are defined only in CUDA builds.
+    py::module_ cuda = m.def_submodule(
+        "cuda", "GPU (CUDA) implementations; populated only in CUDA builds.");
+    define_cuda_lbvh(cuda);
+
     // candidates
     define_candidates(m);
     define_collision_stencil(m);

diff --git a/python/src/broad_phase/CMakeLists.txt b/python/src/broad_phase/CMakeLists.txt
@@ -2,6 +2,7 @@ set(SOURCES
   aabb.cpp
   broad_phase.cpp
   brute_force.cpp
+  cuda_lbvh.cpp
   hash_grid.cpp
   lbvh.cpp
   spatial_hash.cpp

diff --git a/python/src/broad_phase/bindings.hpp b/python/src/broad_phase/bindings.hpp
@@ -5,6 +5,7 @@
 void define_aabb(py::module_& m);
 void define_broad_phase(py::module_& m);
 void define_brute_force(py::module_& m);
+void define_cuda_lbvh(py::module_& cuda); // into the ipctk.cuda submodule
 void define_hash_grid(py::module_& m);
 void define_lbvh(py::module_& m);
 void define_spatial_hash(py::module_& m);

diff --git a/python/src/broad_phase/cuda_lbvh.cpp b/python/src/broad_phase/cuda_lbvh.cpp
@@ -0,0 +1,32 @@
+#include <common.hpp>
+
+#include <ipc/broad_phase/cuda/lbvh.hpp>
+
+#ifdef IPC_TOOLKIT_WITH_CUDA
+using namespace ipc; // not defined if IPC_TOOLKIT_WITH_CUDA is not defined
+#endif
+
+void define_cuda_lbvh(py::module_& m) // m is the ipctk.cuda submodule
+{
+#ifdef IPC_TOOLKIT_WITH_CUDA
+    py::class_<cuda::LBVH, BroadPhase, std::shared_ptr<cuda::LBVH>>(
+        m, "LBVH",
+        R"ipc_Qu8mg5v7(
+        GPU Linear Bounding Volume Hierarchy (LBVH) broad phase (ipc::cuda::LBVH).
+
+        Builds the vertex/edge/face AABBs and BVHs, and runs the traversal and
+        mesh-connectivity filtering, on the device. Produces the same candidates
+        as ipctk.LBVH for any vertex filter. Available only in CUDA builds.
+        )ipc_Qu8mg5v7")
+        .def(py::init())
+        .def_property_readonly(
+            "num_vertex_nodes", &cuda::LBVH::num_vertex_nodes,
+            "Number of nodes in the vertex BVH (2 * n_leaves - 1, or 0).")
+        .def_property_readonly(
+            "num_edge_nodes", &cuda::LBVH::num_edge_nodes,
+            "Number of nodes in the edge BVH (2 * n_leaves - 1, or 0).")
+        .def_property_readonly(
+            "num_face_nodes", &cuda::LBVH::num_face_nodes,
+            "Number of nodes in the face BVH (2 * n_leaves - 1, or 0).");
+#endif
+}
diff --git a/python/tests/utils.py b/python/tests/utils.py
@@ -45,6 +45,8 @@ def broad_phases():
     yield ipctk.SpatialHash()
     yield ipctk.LBVH()
     yield ipctk.SweepAndPrune()
+    if hasattr(ipctk.cuda, "LBVH"):  # only in CUDA builds
+        yield ipctk.cuda.LBVH()
 
 
 def finite_jacobian(x, f, h=1e-8):

diff --git a/src/ipc/broad_phase/CMakeLists.txt b/src/ipc/broad_phase/CMakeLists.txt
@@ -23,3 +23,8 @@ set(SOURCES
 )
 
 target_sources(ipc_toolkit PRIVATE ${SOURCES})
+
+add_subdirectory(details)
+if(IPC_TOOLKIT_WITH_CUDA)
+  add_subdirectory(cuda)
+endif()
diff --git a/src/ipc/broad_phase/aabb.cpp b/src/ipc/broad_phase/aabb.cpp
@@ -44,13 +44,11 @@ void AABB::conservative_inflation(
     // Nudge the bounds outward to ensure conservativity.
 
     min = min.unaryExpr([inflation_radius](double v) {
-        return std::nextafter(
-            v - inflation_radius, -std::numeric_limits<double>::infinity());
+        return conservative_lower_bound(v, inflation_radius);
     });
 
     max = max.unaryExpr([inflation_radius](double v) {
-        return std::nextafter(
-            v + inflation_radius, std::numeric_limits<double>::infinity());
+        return conservative_upper_bound(v, inflation_radius);
     });
 }
 

diff --git a/src/ipc/broad_phase/aabb.hpp b/src/ipc/broad_phase/aabb.hpp
@@ -3,8 +3,10 @@
 #include <ipc/config.hpp>
 #include <ipc/utils/default_init_allocator.hpp>
 #include <ipc/utils/eigen_ext.hpp>
+#include <ipc/utils/simd.hpp> // for infinity<T>()
 
 #include <array>
+#include <cmath> // for nextafter
 
 namespace ipc {
 
@@ -60,6 +62,33 @@ class alignas(64) AABB {
         Eigen::Ref<ArrayMax3d> max,
         const double inflation_radius);
 
+    /// @brief Conservatively inflate one lower bound.
+    ///
+    /// The single-coordinate policy behind conservative_inflation(): the bound
+    /// is moved out by the radius, then nudged to the next representable double
+    /// away from the box so rounding can never shrink it. Host/device so the
+    /// GPU broad phases build bit-identical boxes.
+    ///
+    /// @param v The coordinate to bound from below.
+    /// @param inflation_radius The radius to inflate by.
+    /// @return The conservative lower bound.
+    IPC_TOOLKIT_HOST_DEVICE static double
+    conservative_lower_bound(const double v, const double inflation_radius)
+    {
+        return nextafter(v - inflation_radius, -infinity<double>());
+    }
+
+    /// @brief Conservatively inflate one upper bound.
+    /// @see conservative_lower_bound
+    /// @param v The coordinate to bound from above.
+    /// @param inflation_radius The radius to inflate by.
+    /// @return The conservative upper bound.
+    IPC_TOOLKIT_HOST_DEVICE static double
+    conservative_upper_bound(const double v, const double inflation_radius)
+    {
+        return nextafter(v + inflation_radius, infinity<double>());
+    }
+
 public:
     /// @brief Minimum corner of the AABB.
     Eigen::Array3d min;