Add the host/device shared source mechanism

[zfergus]

Description

Stacked on #254. Review that one first; this diff is only the layer above it.

Adds the mechanism for compiling the low-level math for both host C++ and CUDA device code. point_line is converted as the worked example — the rest of distance/, geometry/, tangent/, and barrier/ follow in a separate PR, so the design can be argued here on one file instead of fifty.

Three pieces:

• IPC_TOOLKIT_HOST_DEVICE annotates a declaration so one header is valid in host and device code. It expands to __host__ __device__ only under nvcc and to nothing otherwise, so a non-CUDA build sees the code it sees today.
• ipc_toolkit_target_shared_device_sources() compiles such a .cpp as a genuine .cu under CUDA. Setting LANGUAGE CUDA on a .cpp is not enough: CMake's separable-compilation and device-link bookkeeping key off the .cu extension, so the .cpp would get a host symbol but no linkable __device__ symbol. The helper generates a one-line .cu wrapper in the build tree instead, so nothing extra is committed. With CUDA off, each .cpp is compiled normally.
• Split instantiation. Under CUDA the file is compiled twice, and the two passes divide their explicit instantiations via IPC_TOOLKIT_INSTANTIATE_{DEVICE,HOST}_SCALARS: nvcc takes float and double, the host compiler takes the autodiff and xsimd batch scalars, which have no device-callable operations. Each symbol is emitted exactly once, so the two objects link cleanly.

Definitions stay out of the headers, so editing one rebuilds a single translation unit plus a link step rather than every TU that includes the header.

API changes

• IPC_TOOLKIT_HOST_DEVICE and the two IPC_TOOLKIT_INSTANTIATE_* switches are new public macros in ipc/config.hpp. All three are no-ops without nvcc.
• No function signature, type, or behavior changes.

Effect on the default (CUDA off) build

None. With IPC_TOOLKIT_WITH_CUDA=OFF the only Eigen define reaching the compiler is EIGEN_DONT_VECTORIZE=1, exactly as on main. IPC_TOOLKIT_HOST_DEVICE expands to nothing and both IPC_TOOLKIT_INSTANTIATE_* switches are 1, so a non-CUDA build compiles the same code from the same single translation unit per file.

Eigen alignment under CUDA

When CUDA is enabled, EIGEN_MAX_STATIC_ALIGN_BYTES and EIGEN_MAX_ALIGN_BYTES are pinned to 0. EIGEN_DONT_VECTORIZE alone gives align=0 under the host compiler but align=16 under nvcc, which forces 16-byte alignment when compiling CUDA — so fixed-max types such as VectorMax<double, 18> get a different sizeof/alignof in each, and returning one by value across that boundary smashes the stack. Pinning to 0 makes nvcc behave the way the host already does.

The value has to be 0, not 16. Any non-zero setting is a cap, and Eigen applies it without consulting the scalar's own alignment — so a cap below the natural alignment of an xsimd batch (32 bytes on AVX2, 64 on AVX-512) asks Eigen to under-align an array of them, which clang rejects outright:

error: requested alignment is less than minimum alignment of 32
       for type 'xsimd::batch<float>[9]'

With no cap, an inline batch array simply gets its element type's natural alignment, which is correct on every architecture. Every SimdBatch instantiation in the library lands on fixed-size or max-size types (VectorMax<T, N> is Eigen::Matrix<T, Dynamic, 1, ColMajor, N, 1>, so its storage is an inline array, not heap), so no heap path is involved.

How Has This Been Tested?

The alignment pin, measured on hardware

The choice of 0 is not a guess. On an AVX2 machine with clang, compiling an Eigen::Matrix<xsimd::batch<float>, 36, 12> under each candidate pin:

architecture	alignof(batch<float>)	pin=16	pin=32	pin=0
AVX2 (-march=native)	32	✗ error	✓ ok	✓ ok
AVX-512 (-march=x86-64-v4)	64	✗ error	✗ error	✓ ok

A cap of 16 is what broke Clang-Tidy CI. A cap of 32 fixes AVX2 but relocates the same failure to AVX-512 — and since SIMD_CXX_FLAGS is -march=native, that is a real host, not a hypothetical. Only 0 holds everywhere, which is the point: 0 means "emit no alignas", so an inline batch array simply keeps its element type's natural alignment.

Builds and suites

• Full CPU suite, Linux x86 AVX2 / clang 22: 4,189,135 assertions in 324 test cases. This is the platform where the pin matters; earlier arm64 runs could not have caught the bug, since NEON batches are 16 bytes and satisfy a 16-byte cap.
• Full CPU suite, macOS arm64 / AppleClang 21: 4,173,897 assertions in 324 test cases.
• With IPC_TOOLKIT_WITH_CUDA=OFF, the only Eigen define reaching the compiler is EIGEN_DONT_VECTORIZE=1 — identical to main. Confirmed by inspecting the generated compile line.
• Full CUDA build with nvcc 13.3, CMAKE_CUDA_ARCHITECTURES=86 — 525/525 targets including the device link. This is the check that matters for this PR's mechanism: it exercises the generated .cu wrappers, the relocatable-device-code link, and the split instantiation. A dropped or duplicated scalar would surface as an undefined or duplicate symbol at that link; neither appeared. Built at the stack tip (Make the low-level math callable from device code #257), which contains this PR.
• Clang-Tidy, CUDA (Debug), CUDA (Release) and all remaining CI checks pass.

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 Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 96.69%. Comparing base (415a8bb) to head (d67670a).

Additional details and impacted files

@@                  Coverage Diff                   @@
##           feature/cuda-build     #255      +/-   ##
======================================================
- Coverage               96.71%   96.69%   -0.02%     
======================================================
  Files                     191      191              
  Lines                   17295    17293       -2     
  Branches                  928      928              
======================================================
- Hits                    16726    16722       -4     
- Misses                    569      571       +2     

Flag	Coverage Δ
unittests	96.69% <100.00%> (-0.02%)	⬇️

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.