Make the low-level math callable from device code - #257
Merged
zfergus merged 4 commits intoSep 8, 2026
Conversation
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## feature/shared-device-sources #257 +/- ##
=================================================================
+ Coverage 96.69% 96.74% +0.05%
=================================================================
Files 191 191
Lines 17293 17292 -1
Branches 928 933 +5
=================================================================
+ Hits 16722 16730 +8
+ Misses 571 562 -9
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:
|
zfergus
force-pushed
the
feature/device-callable-math
branch
from
September 8, 2026 04:58
a2580b8 to
e43e96f
Compare
This was referenced Sep 8, 2026
zfergus
force-pushed
the
feature/device-callable-math
branch
from
September 8, 2026 16:29
e43e96f to
61a719d
Compare
Extends the mechanism from the previous commit across distance/, geometry/, tangent/, barrier/, and math/, and adds GPU unit tests that check each function against its host result. Annotating the declarations is most of the diff, but a few functions had to stop using host-only constructs to compile for the device: - distance_type.cpp is split, moving the logger and the exceptions into distance_type_host.cpp. nvcc cannot parse spdlog, so a shared device source must not include it; the header picks between the host helpers and a device trap. - barrier.cpp keeps only the three free functions. The virtual Barrier hierarchy moves to barrier_classes.cpp, which stays host-only -- virtual dispatch has nothing to offer device code. - morton.hpp gets a device-safe clamp. MSVC's debug std::clamp checks bounds with _STL_VERIFY, which expands to __debugbreak(), and nvcc's NVVM backend then emits invalid IR. The comparison order matches std::clamp exactly. - angle.cpp scatters the normal-Jacobian blocks with middleCols instead of dn1_dx(Eigen::all, idx), since Eigen's index slicing is unavailable in device code. It also returns an aggregate rather than a std::pair: nvcc admits libstdc++'s constexpr std::pair into device code only under --expt-relaxed-constexpr, which covers compile-time evaluation, so a pair built at run time inside a kernel comes back zero-filled. The xsimd instantiation blocks are guarded on IPC_TOOLKIT_WITH_SIMD alone. The parent commit retracts that macro for nvcc, so the blocks no longer need an IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS term; only the autodiff instantiations in edge_edge.cpp and point_triangle.cpp still use it. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
zfergus
force-pushed
the
feature/device-callable-math
branch
from
September 8, 2026 17:42
61a719d to
511cb01
Compare
barrier_classes.cpp and distance_type_host.cpp existed only to keep code away from nvcc, which left barrier.hpp and distance_type.hpp each implemented across two .cpp files. The dual-compilation mechanism does not require a separate file, only that the host-only code be invisible to the device pass, so both are merged back behind `#ifndef __CUDACC__`. - barrier.cpp regains the virtual Barrier hierarchy. Virtual dispatch cannot cross the host/device boundary -- a host-built vtable holds host code addresses, CUDA forbids passing an object of a class with virtual functions to a __global__ function, and BarrierPotential owns its barrier through a host-only std::shared_ptr -- so the guard also keeps every class symbol, float and double included, in the host object. - distance_type.cpp regains the error reporting helpers. Here the blocker is that nvcc cannot parse spdlog, so the guard covers the includes as well as the definitions; an #include in a false branch is never processed. Every use of the logger and of fmt must stay inside it. The free functions in both files are unchanged and still emit relocatable device code: cuobjdump confirms all six barrier functions in the fatbin, and the device object holds zero class symbols and zero spdlog symbols. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
morton.hpp carried a private detail::clamp, a device-safe std::clamp usable only from that header. It now sits in simd.hpp beside select, all_of, and infinity, which pair a scalar and a batch overload the same way, and covers float, double, and an xsimd batch. - The scalar overload is IPC_TOOLKIT_HOST_DEVICE, so it works on the GPU. std::clamp does not: MSVC's debug STL checks the bounds with _STL_VERIFY, which expands to __debugbreak(), and nvcc's NVVM backend then emits invalid IR. The comparison order matches std::clamp exactly, so a NaN v still passes through unchanged. - The batch overload blends the same first-match-wins order with two xsimd::select calls, since a batch comparison answers per-lane rather than with one bool. It is the more specialized template, so it wins for a batch argument. Host only, as config.hpp retracts IPC_TOOLKIT_WITH_SIMD for a CUDA translation unit. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
The distance, geometry, and tangent GPU tests each carried a byte-identical copy of the same seven helpers -- REQUIRE_CUDA, GpuTestKernel, write_out, write_expected, run_gpu_kernel, skip_if_no_cuda_device, and check_gpu_matches_host -- and the barrier test hand-rolled a fourth version of the same launch, transfer, and comparison. They now live in tests/src/tests/gpu_utils.hpp, beside simd_utils.hpp and following its include and namespace conventions, leaving only the kernels and their fixtures in the individual files. - check_gpu_matches_host compares an infinite reference by sign rather than by margin. That was the only thing keeping the barrier test from sharing the comparison, and it is a no-op for the finite values the other three produce. - The header defines __device__ code, so it #errors when included outside a .cu rather than failing on an unrecognized __device__. - geometry's copy documented run_gpu_kernel as skipping the test when no device is present, which it never did; the comment now sits on skip_if_no_cuda_device. Net -152 lines. Verified with the full suite (345/345) and with the [gpu] filter, whose 21 cases execute rather than skip on this machine.
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
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
Applies
IPC_TOOLKIT_HOST_DEVICEand the shared-device-source helper acrossdistance/,geometry/,tangent/,barrier/, andmath/, and adds GPU unit tests that check each function's device result against its host result.Annotating declarations is most of the diff and is mechanical. Four functions had to stop using host-only constructs, and those are the parts worth reading:
distance_type.cppis split. The logger and the exceptions move todistance_type_host.cpp. nvcc cannot parse spdlog, so a shared device source must not include it; inline wrappers in the header pick between the host helpers and a device-side trap.barrier.cppkeeps only the three free functions. The virtualBarrierhierarchy moves tobarrier_classes.cpp, which stays host-only — virtual dispatch has nothing to offer device code, and keeping it out letsbarrier.cppcompile as a.cu.morton.hppgets a device-safe clamp. MSVC's debugstd::clampchecks its bounds with_STL_VERIFY, which expands to__debugbreak(), and nvcc's NVVM backend then emits invalid IR ("Terminator found in the middle of a basic block"). The comparison order matchesstd::clampexactly, so a NaN input still passes through unchanged.angle.cppscatters the normal-Jacobian blocks withmiddleColsinstead ofdn1_dx(Eigen::all, idx), because Eigen's index slicing is unavailable in device code. It also returns an aggregate rather than astd::pair: nvcc admits libstdc++'sconstexpr std::pairinto device code only under--expt-relaxed-constexpr, which covers compile-time evaluation — a pair built at run time inside a kernel silently comes back zero-filled.API changes
barrier_classes.cppanddistance_type_host.cppare new files, both host-only.Effect on the default (CUDA off) build
IPC_TOOLKIT_HOST_DEVICEexpands to nothing and theIPC_TOOLKIT_INSTANTIATE_*switches are both 1, so a non-CUDA build instantiates exactly the scalars it does today from a single translation unit per file. The GPU tests are gated onIPC_TOOLKIT_WITH_CUDAand are not compiled.How Has This Been Tested?
The GPU tests actually run
CMAKE_CUDA_ARCHITECTURES=86: 525/525 targets including the device link.CPU suites
ipc_toolkit,ipc_toolkit_tests, andipctkall compile and link, which is the check that matters for the instantiation split: a dropped scalar surfaces as an undefined symbol.CUDA (Debug),CUDA (Release)andClang-Tidy.One pre-existing bug these tests expose
Running the new GPU tests in the same process as the pre-existing
test_gpu_ccdtests fails in roughly 5 of 6 randomized orders. The first error is an illegal memory access inside Scalable CCD's narrow phase (cuda/narrow_phase/root_finder.cu); because that error is sticky, every latercudaMallocin the process then returns 700, which is why the collateral count varies with the ordering.Scoped runs separate it cleanly:
[gpu] ~[ccd]— the 19 tests added here, randomized, 3 runs[gpu]in declaration order — all 21, including CCD[gpu]in most randomized ordersNot caused by anything in this stack. I built the previous and current Scalable CCD pins and ran four fixed seeds against each; failure counts are identical. The cause is that
MAX_UNIT_SIZEis sized from currently-free GPU memory whileroot_finder.cusetsm_tailfromd_data.size()independently, so under memory pressure the buffer's tail exceeds its capacity — hence the order dependence. It is upstream and pre-existing, and it wants its own issue. CI has never caught it:cuda.ymlbuilds but never runs tests, since the runners have no GPU.Test Configuration:
CMAKE_BUILD_TYPE=ReleaseChecklist