Skip to content

Add the host/device shared source mechanism - #255

Merged
zfergus merged 2 commits into
feature/cuda-buildfrom
feature/shared-device-sources
Sep 8, 2026
Merged

Add the host/device shared source mechanism#255
zfergus merged 2 commits into
feature/cuda-buildfrom
feature/shared-device-sources

Conversation

@zfergus

@zfergus zfergus commented Sep 8, 2026

Copy link
Copy Markdown
Member

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

@zfergus
zfergus force-pushed the feature/shared-device-sources branch 3 times, most recently from 650349e to a224c6e Compare September 8, 2026 04:12
@codecov

codecov Bot commented Sep 8, 2026

Copy link
Copy Markdown

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.

@zfergus
zfergus force-pushed the feature/shared-device-sources branch from a224c6e to c2dbe7b Compare September 8, 2026 04:58
@zfergus zfergus mentioned this pull request Sep 8, 2026
13 tasks
@zfergus
zfergus force-pushed the feature/shared-device-sources branch from c2dbe7b to 9f1c4fd Compare September 8, 2026 16:24
Groundwork for compiling the low-level math for both host C++ and CUDA
device code, applied here to point_line only as a worked example.

Three pieces:

- IPC_TOOLKIT_HOST_DEVICE annotates a declaration so the same header is valid
  in host and device code, expanding to `__host__ __device__` only under nvcc.
- 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 bookkeeping keys 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.
- The two passes split their template 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
  objects link cleanly.

We also pin Eigen's alignment to 0 when CUDA is enabled. 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 (e.g.
VectorMax<double, 18>) end up with a different sizeof/alignof in each and
smash the stack when returned by value across that boundary. Pinning to 0
makes nvcc behave the way the host already does.

The value has to be 0 rather than 16: any non-zero setting is a cap that Eigen
applies 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 it
to under-align an array of them, which clang rejects. Scoping it to CUDA keeps
the default build byte-identical to before.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
@zfergus
zfergus force-pushed the feature/shared-device-sources branch from 9f1c4fd to b2c749c Compare September 8, 2026 16:31
- Keep the generated .cu wrappers out of source_group(TREE). They live under
  the binary dir, which made configure fail for any build directory outside
  the source tree, i.e. every CPM/FetchContent consumer.
- Forward SIMD_CXX_FLAGS to nvcc's host pass via -Xcompiler. That pass owns
  the float and double instantiations of a shared device source, so those
  were the only scalar code in the library compiled without AVX/FMA.
- Move the Eigen alignment pin out of the SIMD block. The host-versus-nvcc
  layout disagreement it prevents is also reachable with SIMD off.
- Retract IPC_TOOLKIT_WITH_SIMD for nvcc in config.hpp, so SIMD-guarded
  blocks need no CUDA-specific spelling once their file becomes a shared
  device source.
- Gate --expt-relaxed-constexpr on the CUDA compiler being nvcc; clang as
  the CUDA compiler rejects it.
- Pass /utf-8 to the nvcc host pass on MSVC, making the fmt patch's
  hardcoded UTF-8 claim true rather than only silencing fmt's static_assert.
- Warn when an enclosing project supplies spdlog::spdlog with CUDA enabled,
  since the fmt patch then never applies.
- Unify the two dev-container Dockerfiles into one parameterized by
  BASE_IMAGE.
- build-cuda.sh: exclude tests/data from the rsync so --delete cannot strip
  it while its ExternalProject stamp survives, default to one CUDA
  architecture for a compile-only check, and drop two dead lines.
- .dockerignore: keep tests/data and nested .git out of the build context.
- Drop the duplicate set_source_files_properties calls and the unreachable
  rel_dir branch in the shared device sources helper.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
@zfergus
zfergus merged commit 006af56 into main Sep 8, 2026
21 checks passed
@zfergus
zfergus deleted the feature/shared-device-sources 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