Skip to content

Make the low-level math callable from device code - #257

Merged
zfergus merged 4 commits into
feature/shared-device-sourcesfrom
feature/device-callable-math
Sep 8, 2026
Merged

Make the low-level math callable from device code#257
zfergus merged 4 commits into
feature/shared-device-sourcesfrom
feature/device-callable-math

Conversation

@zfergus

@zfergus zfergus commented Sep 8, 2026

Copy link
Copy Markdown
Member

Description

Stacked on #255, which introduces the mechanism this PR applies. Review that one first.

Applies IPC_TOOLKIT_HOST_DEVICE and the shared-device-source helper across distance/, geometry/, tangent/, barrier/, and math/, 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.cpp is split. The logger and the exceptions move to distance_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.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, and keeping it out lets barrier.cpp compile as a .cu.
  • morton.hpp gets a device-safe clamp. MSVC's debug std::clamp checks 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 matches std::clamp exactly, so a NaN input still passes through unchanged.
  • angle.cpp scatters the normal-Jacobian blocks with middleCols instead of dn1_dx(Eigen::all, idx), because 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 — a pair built at run time inside a kernel silently comes back zero-filled.

API changes

  • No public function signature, type, or behavior changes. The two source splits are internal reorganization; every declaration keeps its current spelling.
  • barrier_classes.cpp and distance_type_host.cpp are new files, both host-only.

Effect on the default (CUDA off) build

IPC_TOOLKIT_HOST_DEVICE expands to nothing and the IPC_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 on IPC_TOOLKIT_WITH_CUDA and are not compiled.

How Has This Been Tested?

The GPU tests actually run

  • 19 GPU test cases executed on an RTX 3070: 5,430 assertions, all passing, over 3 consecutive randomized-order runs. Each case checks a device result against the host result for the same inputs, so this is what confirms the annotated functions compute the same thing in both execution spaces — not merely that they compile.
  • Built with nvcc 13.3, CMAKE_CUDA_ARCHITECTURES=86: 525/525 targets including the device link.

CPU suites

  • Linux x86 AVX2 / clang 22: 4,189,135 assertions in 324 test cases.
  • macOS arm64 / AppleClang 21: 4,173,877 assertions in 324 test cases. (Add the host/device shared source mechanism #255 reports 4,173,897 on the same suite; the delta is Catch2's random seed, not a behavior change — two consecutive runs of one unchanged build differ by ~38 assertions.)
  • ipc_toolkit, ipc_toolkit_tests, and ipctk all compile and link, which is the check that matters for the instantiation split: a dropped scalar surfaces as an undefined symbol.
  • All CI checks pass, including CUDA (Debug), CUDA (Release) and Clang-Tidy.

One pre-existing bug these tests expose

Running the new GPU tests in the same process as the pre-existing test_gpu_ccd tests 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 later cudaMalloc in the process then returns 700, which is why the collateral count varies with the ordering.

Scoped runs separate it cleanly:

selection result
[gpu] ~[ccd] — the 19 tests added here, randomized, 3 runs all pass (5,430 assertions)
[gpu] in declaration order — all 21, including CCD all pass (5,445 assertions)
either CCD case on its own, 3 runs each all pass
[gpu] in most randomized orders CCD fails, rest cascade

Not 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_SIZE is sized from currently-free GPU memory while root_finder.cu sets m_tail from d_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.yml builds but never runs tests, since the runners have no GPU.

Test Configuration:

  • OS and Version: Arch Linux x86_64 (AMD Ryzen 9 5900XT, AVX2) and macOS 26.6 (arm64)
  • Compiler and Version: clang 22.1.8, gcc 16.2.1, nvcc 13.3 (V13.3.73), AppleClang 21.0; CMake 4.4.2, CMAKE_BUILD_TYPE=Release
  • GPU: NVIDIA GeForce RTX 3070 (compute capability 8.6)

Checklist

  • I have followed the project style guide
  • My code follows the clang-format style guidelines of this project
  • I have performed a self-review of my code
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes
  • Any dependent changes have been merged and published in downstream modules

@codecov

codecov Bot commented Sep 8, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 96.72727% with 9 lines in your changes missing coverage. Please review.
✅ Project coverage is 96.74%. Comparing base (d67670a) to head (785174b).

Files with missing lines Patch % Lines
src/ipc/distance/distance_type.cpp 58.33% 5 Missing ⚠️
src/ipc/distance/distance_type.hpp 60.00% 4 Missing ⚠️
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     
Flag Coverage Δ
unittests 96.74% <96.72%> (+0.05%) ⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

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
zfergus force-pushed the feature/device-callable-math branch from 61a719d to 511cb01 Compare September 8, 2026 17:42
zfergus and others added 3 commits September 8, 2026 16:02
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.
@zfergus
zfergus merged commit 929086d into main Sep 8, 2026
19 checks passed
@zfergus
zfergus deleted the feature/device-callable-math branch September 8, 2026 20:37
@zfergus zfergus added this to the v2.0.0 milestone Sep 8, 2026
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.

1 participant