From b2c749cc1f82b6fcda92f0dd0571cef439bfd865 Mon Sep 17 00:00:00 2001 From: Zachary Ferguson Date: Tue, 8 Sep 2026 00:01:03 -0400 Subject: [PATCH 1/2] Add the host/device shared source mechanism 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) 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 --- CMakeLists.txt | 23 +++++ .../ipc_toolkit_shared_device_sources.cmake | 94 +++++++++++++++++++ src/ipc/config.hpp.in | 35 +++++++ src/ipc/distance/CMakeLists.txt | 10 +- src/ipc/distance/point_line.cpp | 12 ++- src/ipc/distance/point_line.hpp | 23 +++-- 6 files changed, 181 insertions(+), 16 deletions(-) create mode 100644 cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 54a6df7d4..89787c6b9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -129,6 +129,7 @@ list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake/find/") # General CMake utils include(ipc_toolkit_cpm_cache) include(ipc_toolkit_use_colors) +include(ipc_toolkit_shared_device_sources) # Generate position-independent code by default set(CMAKE_POSITION_INDEPENDENT_CODE ON) @@ -328,6 +329,28 @@ if(IPC_TOOLKIT_WITH_SIMD) # Disable vectorization in Eigen since I've found it to have alignment issues. # NOTE: I don't know why this needs to be public, but it crashes if I make it private. target_compile_definitions(ipc_toolkit PUBLIC EIGEN_DONT_VECTORIZE=1) + + if(IPC_TOOLKIT_WITH_CUDA) + # Pin Eigen's alignment so C++ and CUDA translation units agree on object + # layout. EIGEN_DONT_VECTORIZE alone yields align=0 under the host compiler + # but align=16 under nvcc, which forces 16-byte alignment when compiling + # CUDA. Fixed-max types (e.g. VectorMax) then get a different + # sizeof/alignof in each, which smashes the stack when one is returned by + # value across that boundary. PUBLIC because every TU exchanging Eigen + # objects with the library must agree, or it is an ODR violation. + # + # NOTE: We pin to 0, not 16. Any non-zero pin is a cap, and Eigen applies + # it without checking the scalar's own alignment -- so a cap below the + # natural alignment of an xsimd batch (32 bytes on AVX2, 64 on AVX-512) + # makes Eigen ask to under-align an array of them, which clang rejects + # outright. Since EIGEN_DONT_VECTORIZE is already set, alignment buys us + # nothing here anyway; 0 means "emit no alignas", which is both what the + # host compiler already does and the one value that cannot conflict with + # any scalar on any architecture. + target_compile_definitions(ipc_toolkit PUBLIC + EIGEN_MAX_STATIC_ALIGN_BYTES=0 + EIGEN_MAX_ALIGN_BYTES=0) + endif() endif() # For MSVC, do not use the min and max macros. diff --git a/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake b/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake new file mode 100644 index 000000000..5fd9a5468 --- /dev/null +++ b/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake @@ -0,0 +1,94 @@ +# Add C++ sources whose definitions are shared between host C++ and CUDA device +# code. +# +# The matching header declares the functions with IPC_TOOLKIT_HOST_DEVICE and +# does NOT textually include the .cpp, so editing a definition rebuilds a single +# translation unit plus a link step instead of every TU that includes the +# header. +# +# Without CUDA each .cpp is compiled normally. With CUDA each .cpp must be +# compiled as a genuine .cu so nvcc emits *relocatable* device code that other +# TUs can device-link against. 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 gets a host symbol but no linkable __device__ symbol. +# For each .cpp this function generates a one-line .cu wrapper in the build tree +# (so nothing extra is committed) and adds it to the target. +# +# Under CUDA the .cpp is added *as well*, compiled by the host compiler. These +# files instantiate their templates for more than float and double -- autodiff +# and xsimd batch scalars -- and neither has device-callable operations, so +# nvcc cannot compile those instantiations. The two passes therefore split the +# instantiations rather than duplicating them: the .cu wrapper gets +# IPC_TOOLKIT_DEVICE_SCALARS_ONLY and the .cpp gets IPC_TOOLKIT_HOST_SCALARS_ONLY, +# which drive the IPC_TOOLKIT_INSTANTIATE_* switches in config.hpp. Each symbol +# is emitted exactly once, so the two objects link together cleanly. +# +# Usage: +# ipc_toolkit_target_shared_device_sources(ipc_toolkit +# point_point.cpp +# distance_type.cpp +# ...) +function(ipc_toolkit_target_shared_device_sources target) + set(sources_to_add "") + foreach(src IN LISTS ARGN) + if(IPC_TOOLKIT_WITH_CUDA) + get_filename_component(src_abs "${src}" ABSOLUTE + BASE_DIR "${CMAKE_CURRENT_SOURCE_DIR}") + + # Mirror the source's path (relative to the current source dir) under the + # wrapper directory so that duplicate base names in different subfolders + # (e.g. a/point.cpp and b/point.cpp) get distinct wrapper paths instead of + # colliding on a single flattened name. + file(RELATIVE_PATH src_rel "${CMAKE_CURRENT_SOURCE_DIR}" "${src_abs}") + get_filename_component(rel_dir "${src_rel}" DIRECTORY) + get_filename_component(src_name_we "${src_rel}" NAME_WE) + + set(wrapper_dir "${CMAKE_CURRENT_BINARY_DIR}/cuda_device_wrappers") + if(rel_dir) + set(wrapper_dir "${wrapper_dir}/${rel_dir}") + endif() + set(wrapper "${wrapper_dir}/${src_name_we}.cu") + + # file(GENERATE) only rewrites when the content changes, so this does not + # trigger spurious rebuilds on reconfigure. The wrapper #includes the + # original .cpp by absolute path; the compiler's dependency scan then ties + # the wrapper TU to the .cpp, so editing the .cpp rebuilds just this TU. + file(GENERATE OUTPUT "${wrapper}" CONTENT +"// Generated by ipc_toolkit_target_shared_device_sources() -- do not edit. +// Compiles ${src_rel} as a CUDA translation unit so nvcc emits relocatable +// device code linkable from other TUs. Edit the .cpp instead. +#include \"${src_abs}\" +") + + # Source file properties are directory-scoped and are read from the + # directory that created the target, which is not this subdirectory. Set + # them in both scopes (TARGET_DIRECTORY for the generator, the plain call + # for anything reading them here) and use absolute paths so the lookup + # resolves from either scope. Without TARGET_DIRECTORY the defines are + # silently dropped and both passes instantiate every scalar, so nvcc sees + # the xsimd batch instantiations and fails. + set_source_files_properties("${wrapper}" + TARGET_DIRECTORY ${target} + PROPERTIES + GENERATED TRUE + COMPILE_DEFINITIONS IPC_TOOLKIT_DEVICE_SCALARS_ONLY) + set_source_files_properties("${wrapper}" PROPERTIES + GENERATED TRUE + COMPILE_DEFINITIONS IPC_TOOLKIT_DEVICE_SCALARS_ONLY) + list(APPEND sources_to_add "${wrapper}") + + # The same .cpp, compiled by the host compiler for the scalars nvcc + # cannot handle. + set_source_files_properties("${src_abs}" + TARGET_DIRECTORY ${target} + PROPERTIES COMPILE_DEFINITIONS IPC_TOOLKIT_HOST_SCALARS_ONLY) + set_source_files_properties("${src_abs}" PROPERTIES + COMPILE_DEFINITIONS IPC_TOOLKIT_HOST_SCALARS_ONLY) + list(APPEND sources_to_add "${src_abs}") + else() + list(APPEND sources_to_add "${src}") + endif() + endforeach() + + target_sources(${target} PRIVATE ${sources_to_add}) +endfunction() diff --git a/src/ipc/config.hpp.in b/src/ipc/config.hpp.in index 581b1abb6..56f6e414d 100644 --- a/src/ipc/config.hpp.in +++ b/src/ipc/config.hpp.in @@ -22,6 +22,41 @@ #cmakedefine IPC_TOOLKIT_WITH_MESHFEM_SPARSE // #define IPC_TOOLKIT_DEBUG_AUTODIFF +// Execution-space annotation for functions shared between C++ and CUDA. +// IPC_TOOLKIT_HOST_DEVICE expands to `__host__ __device__` under nvcc and to +// nothing otherwise, so the same declaration is valid in host and device code. +// The definitions live in .cpp files compiled once per configuration (as a +// generated .cu under CUDA; see ipc_toolkit_shared_device_sources.cmake). +#if defined(IPC_TOOLKIT_WITH_CUDA) && defined(__CUDACC__) +#define IPC_TOOLKIT_HOST_DEVICE __host__ __device__ +#else +#define IPC_TOOLKIT_HOST_DEVICE +#endif + +// Which scalar types a shared device source should instantiate. +// +// Templates like edge_edge_distance are instantiated for several scalars, but +// only float and double mean anything on the GPU: autodiff and xsimd batches +// have no device-callable operations, so instantiating a __host__ __device__ +// template with them makes nvcc compile host-only calls into device code. +// +// With CUDA off there is one translation unit and it instantiates everything. +// With CUDA on the file is compiled twice -- once by nvcc for the device +// scalars, once by the host compiler for the rest -- and +// ipc_toolkit_target_shared_device_sources() defines one of the two macros +// below per pass so the instantiations are split between them rather than +// duplicated. See ipc_toolkit_shared_device_sources.cmake. +#if defined(IPC_TOOLKIT_HOST_SCALARS_ONLY) +#define IPC_TOOLKIT_INSTANTIATE_DEVICE_SCALARS 0 +#define IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS 1 +#elif defined(IPC_TOOLKIT_DEVICE_SCALARS_ONLY) +#define IPC_TOOLKIT_INSTANTIATE_DEVICE_SCALARS 1 +#define IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS 0 +#else +#define IPC_TOOLKIT_INSTANTIATE_DEVICE_SCALARS 1 +#define IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS 1 +#endif + namespace ipc { // Type definitions diff --git a/src/ipc/distance/CMakeLists.txt b/src/ipc/distance/CMakeLists.txt index 1897af881..f327d575f 100644 --- a/src/ipc/distance/CMakeLists.txt +++ b/src/ipc/distance/CMakeLists.txt @@ -9,7 +9,6 @@ set(SOURCES line_line.hpp point_edge.cpp point_edge.hpp - point_line.cpp point_line.hpp point_plane.cpp point_plane.hpp @@ -20,6 +19,15 @@ set(SOURCES target_sources(ipc_toolkit PRIVATE ${SOURCES}) +# These definitions are shared between host C++ and CUDA device code. The +# header only declares them (IPC_TOOLKIT_HOST_DEVICE), so editing a definition +# rebuilds one TU plus a link step rather than every TU that includes the +# header. Under CUDA the helper generates a .cu wrapper per file so nvcc emits +# relocatable device code linkable from other TUs. +ipc_toolkit_target_shared_device_sources(ipc_toolkit + point_line.cpp +) + ################################################################################ # Subfolders ################################################################################ diff --git a/src/ipc/distance/point_line.cpp b/src/ipc/distance/point_line.cpp index be2c4b04f..c5f2d7b5d 100644 --- a/src/ipc/distance/point_line.cpp +++ b/src/ipc/distance/point_line.cpp @@ -8,7 +8,7 @@ namespace ipc::autogen { // This function was generated by the Symbolic Math Toolbox version 8.3. // 16-Mar-2020 15:56:29 template -void point_line_distance_gradient_2D( +IPC_TOOLKIT_HOST_DEVICE void point_line_distance_gradient_2D( T v01, T v02, T v11, T v12, T v21, T v22, T g[6]) { T t13, t14, t23, t25, t24, t26, t27; @@ -33,7 +33,7 @@ void point_line_distance_gradient_2D( // This function was generated by the Symbolic Math Toolbox version 8.3. // 10-Jun-2019 18:02:37 template -void point_line_distance_gradient_3D( +IPC_TOOLKIT_HOST_DEVICE void point_line_distance_gradient_3D( T v01, T v02, T v03, T v11, T v12, T v13, T v21, T v22, T v23, T g[9]) { T t17, t18, t19, t20, t21, t22, t23, t24, t25, t42, t44, t45, t46, t43, t50, @@ -71,7 +71,7 @@ void point_line_distance_gradient_3D( // This function was generated by the Symbolic Math Toolbox version 8.3. // 16-Mar-2020 15:56:29 template -void point_line_distance_hessian_2D( +IPC_TOOLKIT_HOST_DEVICE void point_line_distance_hessian_2D( T v01, T v02, T v11, T v12, T v21, T v22, T H[36]) { T t15, t16, t17, t18, t19, t20, t21, t22, t23, t24, t31, t34, t32, t33, t35, @@ -170,7 +170,7 @@ void point_line_distance_hessian_2D( // This function was generated by the Symbolic Math Toolbox version 8.3. // 10-Jun-2019 18:02:39 template -void point_line_distance_hessian_3D( +IPC_TOOLKIT_HOST_DEVICE void point_line_distance_hessian_3D( T v01, T v02, T v03, T v11, T v12, T v13, T v21, T v22, T v23, T H[81]) { T t17, t18, t19, t20, t21, t22, t23, t24, t25, t26, t27, t28, t35, t36, t37, @@ -407,9 +407,11 @@ void point_line_distance_hessian_3D( template void point_line_distance_hessian_3D( \ T, T, T, T, T, T, T, T, T, T[81]) +#if IPC_TOOLKIT_INSTANTIATE_DEVICE_SCALARS IPC_INSTANTIATE_POINT_LINE_AUTOGEN(float); IPC_INSTANTIATE_POINT_LINE_AUTOGEN(double); -#ifdef IPC_TOOLKIT_WITH_SIMD +#endif +#if defined(IPC_TOOLKIT_WITH_SIMD) && IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS IPC_INSTANTIATE_POINT_LINE_AUTOGEN(SimdBatch); IPC_INSTANTIATE_POINT_LINE_AUTOGEN(SimdBatch); #endif diff --git a/src/ipc/distance/point_line.hpp b/src/ipc/distance/point_line.hpp index 988f53bfb..1d768b375 100644 --- a/src/ipc/distance/point_line.hpp +++ b/src/ipc/distance/point_line.hpp @@ -1,5 +1,6 @@ #pragma once +#include #include namespace ipc { @@ -7,16 +8,16 @@ namespace ipc { // Symbolically generated derivatives namespace autogen { template - void point_line_distance_gradient_2D( + IPC_TOOLKIT_HOST_DEVICE void point_line_distance_gradient_2D( T v01, T v02, T v11, T v12, T v21, T v22, T g[6]); template - void point_line_distance_gradient_3D( + IPC_TOOLKIT_HOST_DEVICE void point_line_distance_gradient_3D( T v01, T v02, T v03, T v11, T v12, T v13, T v21, T v22, T v23, T g[9]); template - void point_line_distance_hessian_2D( + IPC_TOOLKIT_HOST_DEVICE void point_line_distance_hessian_2D( T v01, T v02, T v11, T v12, T v21, T v22, T H[36]); template - void point_line_distance_hessian_3D( + IPC_TOOLKIT_HOST_DEVICE void point_line_distance_hessian_3D( T v01, T v02, T v03, T v11, T v12, T v13, T v21, T v22, T v23, T H[81]); } // namespace autogen @@ -30,7 +31,7 @@ namespace detail { /// @param e1 The second vertex of the edge defining the line. /// @return The distance between the point and line. template - inline T point_line_distance( + IPC_TOOLKIT_HOST_DEVICE inline T point_line_distance( Eigen::ConstRef> p, Eigen::ConstRef> e0, Eigen::ConstRef> e1) @@ -58,7 +59,8 @@ namespace detail { /// @param e1 The second vertex of the edge defining the line. /// @return The gradient of the distance wrt p, e0, and e1. template - inline Eigen::Vector point_line_distance_gradient( + IPC_TOOLKIT_HOST_DEVICE inline Eigen::Vector + point_line_distance_gradient( Eigen::ConstRef> p, Eigen::ConstRef> e0, Eigen::ConstRef> e1) @@ -85,7 +87,8 @@ namespace detail { /// @param e1 The second vertex of the edge defining the line. /// @return The hessian of the distance wrt p, e0, and e1. template - inline Eigen::Matrix point_line_distance_hessian( + IPC_TOOLKIT_HOST_DEVICE inline Eigen::Matrix + point_line_distance_hessian( Eigen::ConstRef> p, Eigen::ConstRef> e0, Eigen::ConstRef> e1) @@ -111,7 +114,7 @@ namespace detail { /// @param e1 The second vertex of the edge defining the line. /// @return The distance between the point and line. template -inline auto point_line_distance( +IPC_TOOLKIT_HOST_DEVICE inline auto point_line_distance( const Eigen::MatrixBase& p, const Eigen::MatrixBase& e0, const Eigen::MatrixBase& e1) @@ -137,7 +140,7 @@ inline auto point_line_distance( /// @param e1 The second vertex of the edge defining the line. /// @return The gradient of the distance wrt p, e0, and e1. template -inline auto point_line_distance_gradient( +IPC_TOOLKIT_HOST_DEVICE inline auto point_line_distance_gradient( const Eigen::MatrixBase& p, const Eigen::MatrixBase& e0, const Eigen::MatrixBase& e1) @@ -170,7 +173,7 @@ inline auto point_line_distance_gradient( /// @param e1 The second vertex of the edge defining the line. /// @return The hessian of the distance wrt p, e0, and e1. template -inline auto point_line_distance_hessian( +IPC_TOOLKIT_HOST_DEVICE inline auto point_line_distance_hessian( const Eigen::MatrixBase& p, const Eigen::MatrixBase& e0, const Eigen::MatrixBase& e1) From d67670a9c5cafccc512e5f167bb7a86c34c84f62 Mon Sep 17 00:00:00 2001 From: Zachary Ferguson Date: Tue, 8 Sep 2026 13:33:50 -0400 Subject: [PATCH 2/2] Fix out-of-source CUDA builds and host/device compiler flag gaps - 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) --- .devcontainer/Dockerfile | 82 ++++++++------- .devcontainer/cuda/Dockerfile | 76 -------------- .devcontainer/cuda/build-cuda.sh | 21 ++-- .devcontainer/cuda/devcontainer.json | 36 +++++-- .dockerignore | 12 ++- CLAUDE.md | 7 +- CMakeLists.txt | 99 ++++++++++++++----- .../ipc_toolkit_shared_device_sources.cmake | 29 ++---- cmake/recipes/spdlog.cmake | 12 +++ docs/source/build/c++.rst | 2 + src/ipc/config.hpp.in | 11 +++ src/ipc/distance/point_line.cpp | 2 +- src/ipc/utils/simd.hpp | 4 +- 13 files changed, 212 insertions(+), 181 deletions(-) delete mode 100644 .devcontainer/cuda/Dockerfile diff --git a/.devcontainer/Dockerfile b/.devcontainer/Dockerfile index 4bae1e862..f87a2a814 100644 --- a/.devcontainer/Dockerfile +++ b/.devcontainer/Dockerfile @@ -1,4 +1,23 @@ -FROM ubuntu:22.04 +# Development container for the IPC Toolkit. +# +# One file serves both containers. BASE_IMAGE selects which: +# +# ubuntu:22.04 -> .devcontainer/devcontainer.json +# nvidia/cuda:*-devel-ubuntu22.04 -> .devcontainer/cuda/devcontainer.json +# +# The CUDA "devel" images are Ubuntu 22.04 with nvcc and a host gcc already +# installed, so everything below applies unchanged to both. Compiling CUDA does +# not require a GPU; only running it does. On a Linux host with the NVIDIA +# Container Toolkit you can expose the GPU by adding "--gpus=all" to runArgs in +# the CUDA devcontainer.json. +# +# This does NOT copy the source: the workspace is bind-mounted at runtime by +# the Dev Containers tooling, and .devcontainer/cuda/build-cuda.sh rsyncs it +# into a named volume. All of the project's own dependencies are fetched by +# CMake/CPM, so no development libraries are installed here. + +ARG BASE_IMAGE=ubuntu:22.04 +FROM ${BASE_IMAGE} # Set environment variables ENV DEBIAN_FRONTEND=noninteractive @@ -6,11 +25,8 @@ ENV CCACHE_DIR=/home/devuser/.ccache ENV CCACHE_MAXSIZE=1G ENV CXX_STANDARD=17 -# Update package lists -RUN apt-get update - # Install essential packages -RUN apt-get install -y --no-install-recommends \ +RUN apt-get update && apt-get install -y --no-install-recommends \ build-essential \ git \ wget \ @@ -19,23 +35,17 @@ RUN apt-get install -y --no-install-recommends \ zsh \ ninja-build \ ccache \ - libeigen3-dev \ - libtbb-dev \ - libspdlog-dev \ + rsync \ python3 \ python3-pip \ python3-dev \ libgmp-dev \ libssl-dev \ - libncurses5-dev \ - libncursesw5-dev \ - libxml2-dev \ - libjsoncpp-dev \ - libz3-dev \ sudo \ software-properties-common \ lsb-release \ gnupg \ + ca-certificates \ && rm -rf /var/lib/apt/lists/* # Create a new user with sudo privileges @@ -47,36 +57,34 @@ RUN useradd -m devuser \ && chown devuser:devuser $CCACHE_DIR # Set up Python tools -RUN pip3 install --upgrade pip setuptools wheel pre-commit +RUN pip3 install --no-cache-dir --upgrade pip setuptools wheel pre-commit -# Add Kitware APT repository for CMake +# Add Kitware APT repository for a recent CMake (project requires >= 3.24) RUN wget -qO- https://apt.kitware.com/keys/kitware-archive-latest.asc | \ - gpg --dearmor -o /usr/share/keyrings/kitware-archive-keyring.gpg - -RUN echo "deb [signed-by=/usr/share/keyrings/kitware-archive-keyring.gpg] https://apt.kitware.com/ubuntu/ $(lsb_release -cs) main" | \ - tee /etc/apt/sources.list.d/kitware.list > /dev/null - -# Update package lists and install CMake -RUN apt-get update -RUN apt-get install -y cmake + gpg --dearmor -o /usr/share/keyrings/kitware-archive-keyring.gpg \ + && echo "deb [signed-by=/usr/share/keyrings/kitware-archive-keyring.gpg] https://apt.kitware.com/ubuntu/ $(lsb_release -cs) main" | \ + tee /etc/apt/sources.list.d/kitware.list > /dev/null \ + && apt-get update && apt-get install -y --no-install-recommends cmake \ + && rm -rf /var/lib/apt/lists/* \ + && cmake --version \ + && if command -v nvcc > /dev/null; then nvcc --version; fi -# Verify CMake installation -RUN cmake --version - -# Install LLVM/Clang and Clang-Format version 18 -RUN wget -q https://apt.llvm.org/llvm.sh -O /tmp/llvm.sh -RUN chmod +x /tmp/llvm.sh +# Install LLVM/Clang and clang-format 18 so pre-commit's clang-format hook +# matches CI. +RUN wget -q https://apt.llvm.org/llvm.sh -O /tmp/llvm.sh \ + && chmod +x /tmp/llvm.sh +# NOTE: llvm.sh exits non-zero in cases where it has still registered the apt +# source, so its status alone is not a failure. Keep it on its own line: in a +# `&&` chain the trailing `|| true` would also swallow a failed download and +# defer the error to the confusing "Unable to locate package clang-18" below. RUN /tmp/llvm.sh 18 || true +RUN apt-get update \ + && apt-get install -y --no-install-recommends clang-18 clang-tools-18 clang-format-18 \ + && rm -rf /var/lib/apt/lists/* \ + && clang-format-18 --version -RUN apt-get update && apt-get install -y clang-18 clang-tools-18 clang-format-18 - -RUN clang-18 --version -RUN clang++-18 --version -RUN clang-format-18 --version - -# 12. Set the default user and working directory +# Set the default user and working directory USER devuser - WORKDIR /home/devuser/workspace CMD ["bash"] diff --git a/.devcontainer/cuda/Dockerfile b/.devcontainer/cuda/Dockerfile deleted file mode 100644 index 1882bdacd..000000000 --- a/.devcontainer/cuda/Dockerfile +++ /dev/null @@ -1,76 +0,0 @@ -# CUDA-enabled development container for the IPC Toolkit. -# -# This mirrors ../Dockerfile (the default Ubuntu dev container) but is based on -# NVIDIA's CUDA "devel" image so nvcc is available. Like the default container -# it does NOT copy the source: the workspace is bind-mounted at runtime by the -# Dev Containers tooling. Compiling CUDA does not require a GPU; running it -# does. On a Linux host with the NVIDIA Container Toolkit you can expose the GPU -# by adding "--gpus=all" to runArgs in devcontainer.json. - -ARG CUDA_IMAGE=nvidia/cuda:12.6.2-devel-ubuntu22.04 -FROM ${CUDA_IMAGE} - -# Set environment variables -ENV DEBIAN_FRONTEND=noninteractive -ENV CCACHE_DIR=/home/devuser/.ccache -ENV CCACHE_MAXSIZE=1G -ENV CXX_STANDARD=17 - -# Install essential packages (nvcc + host gcc ship in the CUDA devel image). -RUN apt-get update && apt-get install -y --no-install-recommends \ - build-essential \ - git \ - wget \ - curl \ - fish \ - zsh \ - ninja-build \ - ccache \ - rsync \ - python3 \ - python3-pip \ - python3-dev \ - libgmp-dev \ - libssl-dev \ - sudo \ - software-properties-common \ - lsb-release \ - gnupg \ - ca-certificates \ - && rm -rf /var/lib/apt/lists/* - -# Create a new user with sudo privileges -RUN useradd -m devuser \ - && echo "devuser:password" | chpasswd \ - && usermod -aG sudo devuser \ - && echo "devuser ALL=(ALL) NOPASSWD:ALL" >> /etc/sudoers \ - && mkdir -p $CCACHE_DIR \ - && chown devuser:devuser $CCACHE_DIR - -# Set up Python tools -RUN pip3 install --no-cache-dir --upgrade pip setuptools wheel pre-commit - -# Add Kitware APT repository for a recent CMake (project requires >= 3.24) -RUN wget -qO- https://apt.kitware.com/keys/kitware-archive-latest.asc | \ - gpg --dearmor -o /usr/share/keyrings/kitware-archive-keyring.gpg \ - && echo "deb [signed-by=/usr/share/keyrings/kitware-archive-keyring.gpg] https://apt.kitware.com/ubuntu/ $(lsb_release -cs) main" | \ - tee /etc/apt/sources.list.d/kitware.list > /dev/null \ - && apt-get update && apt-get install -y --no-install-recommends cmake \ - && rm -rf /var/lib/apt/lists/* \ - && cmake --version && nvcc --version - -# Install LLVM/Clang and clang-format 18 (matches the default dev container so -# pre-commit's clang-format hook behaves identically). -RUN wget -q https://apt.llvm.org/llvm.sh -O /tmp/llvm.sh \ - && chmod +x /tmp/llvm.sh \ - && /tmp/llvm.sh 18 || true -RUN apt-get update \ - && apt-get install -y --no-install-recommends clang-18 clang-tools-18 clang-format-18 \ - && rm -rf /var/lib/apt/lists/* \ - && clang-format-18 --version - -# Set the default user and working directory -USER devuser -WORKDIR /home/devuser/workspace - -CMD ["bash"] diff --git a/.devcontainer/cuda/build-cuda.sh b/.devcontainer/cuda/build-cuda.sh index 82bf40866..d7b72a7e9 100755 --- a/.devcontainer/cuda/build-cuda.sh +++ b/.devcontainer/cuda/build-cuda.sh @@ -12,9 +12,9 @@ # hermetic and never writes into your host working tree. # # Usage: -# .devcontainer/cuda/build-cuda.sh # cuda-release, arch 75;80;86;89 +# .devcontainer/cuda/build-cuda.sh # cuda-release, arch 75 # PRESET=test .devcontainer/cuda/build-cuda.sh # test preset (CUDA + tests) -# CUDA_ARCH="86" .devcontainer/cuda/build-cuda.sh # single architecture +# CUDA_ARCH="75;80;86;89" .devcontainer/cuda/build-cuda.sh # several archs # JOBS=4 .devcontainer/cuda/build-cuda.sh # limit parallelism (memory) # set -euo pipefail @@ -23,7 +23,11 @@ REPO_ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" IMAGE_NAME="${IMAGE_NAME:-ipc-toolkit-cuda-dev}" PRESET="${PRESET:-cuda-release}" -CUDA_ARCH="${CUDA_ARCH:-75;80;86;89}" +# One architecture is enough to answer "does it compile": nvcc runs the whole +# device front-end per architecture in the list, so the extra ones only repeat +# codegen. 75 is the oldest we support, hence the strictest. Override to build +# a list when you want to check architecture-specific codegen. +CUDA_ARCH="${CUDA_ARCH:-75}" CUDA_IMAGE="${CUDA_IMAGE:-nvidia/cuda:12.6.2-devel-ubuntu22.04}" # Heavy TUs (headers textually include implementations under CUDA) can OOM the # VM at full parallelism; default below nproc. @@ -31,9 +35,9 @@ JOBS="${JOBS:-4}" echo ">> Building CUDA dev image '${IMAGE_NAME}'" docker build \ - -f "${REPO_ROOT}/.devcontainer/cuda/Dockerfile" \ + -f "${REPO_ROOT}/.devcontainer/Dockerfile" \ -t "${IMAGE_NAME}" \ - --build-arg "CUDA_IMAGE=${CUDA_IMAGE}" \ + --build-arg "BASE_IMAGE=${CUDA_IMAGE}" \ "${REPO_ROOT}" echo ">> Compiling (preset=${PRESET}, arch=${CUDA_ARCH})" @@ -50,21 +54,24 @@ docker run --rm --user root \ "${IMAGE_NAME}" \ bash -euo pipefail -c ' export CPM_SOURCE_CACHE=/cpm-cache CCACHE_DIR=/root/.ccache - mkdir -p /workspace # /workspace is a persistent named volume: rsync copies only files # that changed since the last run (the macOS<->VM file-share is slow, # so minimizing reads matters) and ninja can then build incrementally. # The excludes also shield the persistent build/ dir from --delete. + # tests/data is excluded for the same reason: it is cloned by an + # ExternalProject whose stamp lives under build/, so deleting the data + # while keeping the stamp makes the next build fail in gitupdate.cmake + # instead of re-cloning. echo ">> [1/3] Syncing source into the container (delta copy)..." time rsync -a --delete \ --exclude=/build \ --exclude=/.git \ --exclude=/.ccache \ + --exclude=/tests/data \ --exclude=/docs \ --exclude=/notebooks \ --exclude=/IPCToolkitOptions.cmake \ /src/ /workspace/ - rm -f /workspace/IPCToolkitOptions.cmake cd /workspace echo ">> [2/3] Configuring (preset=${PRESET})..." cmake --preset="${PRESET}" -G Ninja \ diff --git a/.devcontainer/cuda/devcontainer.json b/.devcontainer/cuda/devcontainer.json index b22b5fae9..9a6ff54b8 100644 --- a/.devcontainer/cuda/devcontainer.json +++ b/.devcontainer/cuda/devcontainer.json @@ -1,11 +1,18 @@ +// CUDA variant of ../devcontainer.json. It shares ../Dockerfile and differs +// only in the base image it selects, the GPU pass-through note, and the +// CMake-on-open setting; keep everything else in step with the default. { "$schema": "https://raw.githubusercontent.com/devcontainers/spec/main/schemas/devContainer.schema.json", "name": "IPCToolkit CUDA Development Container", "build": { - "dockerfile": "Dockerfile", - // Context is the repository root (two levels up from this file) so the - // build sees the whole project, matching the default dev container. - "context": "../.." + // Shared with the default dev container; BASE_IMAGE is what makes this + // one CUDA-enabled. Context is the repository root (two levels up from + // this file), matching the default. + "dockerfile": "../Dockerfile", + "context": "../..", + "args": { + "BASE_IMAGE": "nvidia/cuda:12.6.2-devel-ubuntu22.04" + } }, // On a Linux host with the NVIDIA Container Toolkit, uncomment the next line // to pass the GPU into the container so the CUDA code can actually run. @@ -14,15 +21,22 @@ "customizations": { "vscode": { "settings": { - "terminal.integrated.shell.linux": "/bin/zsh", + "terminal.integrated.shell.linux": "/bin/fish", "C_Cpp.default.configurationProvider": "ms-vscode.cmake-tools", "C_Cpp.default.intelliSenseMode": "gcc-x64", "C_Cpp.default.compilerPath": "/usr/bin/clang++-18", "C_Cpp.clang_format_path": "/usr/bin/clang-format-18", "C_Cpp.clang_format_style": "file", + // Unlike the default container: opening this one would kick off + // a CUDA configure, which fetches and builds the GPU + // dependencies. Configure when you mean to. "cmake.configureOnOpen": false, "cmake.buildDirectory": "${workspaceFolder}/build", - "python.pythonPath": "/usr/bin/python3" + "python.pythonPath": "/usr/bin/python3", + "python.linting.enabled": true, + "python.linting.pylintEnabled": true, + "python.formatting.provider": "black", + "prettier.requireConfig": true }, "extensions": [ "ms-vscode.cpptools", @@ -31,10 +45,18 @@ "ms-python.python", "ms-azuretools.vscode-docker", "eamodio.gitlens", + "esbenp.prettier-vscode", + "mhutchie.git-graph", "twxs.cmake", "jeff-hykin.better-cpp-syntax", "vadimcn.vscode-lldb", - "ms-python.vscode-pylance" + "cschlosser.doxdocgen", + "ms-python.vscode-pylance", + "mutantdino.resourcemonitor", + "randomfractalsinc.vscode-data-preview", + "oderwat.indent-rainbow", + "formulahendry.code-runner", + "donjayamanne.git-extension-pack" ] } }, diff --git a/.dockerignore b/.dockerignore index 5a09fcf55..b2697f06a 100644 --- a/.dockerignore +++ b/.dockerignore @@ -1,16 +1,20 @@ # Shrinks the build context uploaded to the Docker daemon. Both dev containers -# (.devcontainer/Dockerfile and .devcontainer/cuda/Dockerfile) use the repo -# root as their build context but never COPY it in (the workspace is mounted at -# runtime), so this only speeds up context transfer — it does not affect image -# contents. +# build .devcontainer/Dockerfile with the repo root as their context but never +# COPY it in (the workspace is mounted at runtime), so this only speeds up +# context transfer — it does not affect image contents. build/ docs/_build/ +# The test data is ~120 MB and carries its own nested .git. It is fetched +# inside the container when it is needed, never from the build context. +tests/data/ +.ccache/ *.o *.so *.a __pycache__/ *.pyc .git/ +**/.git/ .venv/ venv/ .DS_Store diff --git a/CLAUDE.md b/CLAUDE.md index 094cfe013..f9582b2f0 100644 --- a/CLAUDE.md +++ b/CLAUDE.md @@ -39,12 +39,17 @@ exits with the build's status: ./.devcontainer/cuda/build-cuda.sh ``` -`PRESET` (default `cuda-release`), `CUDA_ARCH`, and `JOBS` override the +`PRESET` (default `cuda-release`), `CUDA_ARCH` (default `75`, the oldest +architecture we support and so the strictest; a list such as `"75;80;86;89"` +costs one full device codegen pass per entry), and `JOBS` override the defaults, e.g. `PRESET=test ./.devcontainer/cuda/build-cuda.sh` to build the CUDA tests too. The source is mounted read-only and rsynced into a named volume, so a run never writes into your working tree, and later runs are incremental. +The script builds `.devcontainer/Dockerfile`, the same file the non-CUDA dev +container uses, passing `BASE_IMAGE` to select an nvcc-equipped base. + **This needs a running Docker daemon.** On macOS that means Docker Desktop or a colima VM (`colima start`); the failure mode otherwise is a confusing `/var/run/docker.sock` connection error rather than a clear diagnostic. A VM diff --git a/CMakeLists.txt b/CMakeLists.txt index 89787c6b9..68fd32051 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -194,7 +194,7 @@ if(IPC_TOOLKIT_WITH_CUDA) # std::array). PUBLIC so downstream CUDA consumers (e.g. the tests) inherit # the flag. target_compile_options(ipc_toolkit PUBLIC - "$<$:--expt-relaxed-constexpr>") + "$<$,$>:--expt-relaxed-constexpr>") endif() # Fill in configuration options @@ -214,7 +214,25 @@ target_include_directories(ipc_toolkit PUBLIC # Folder name for IDE set_target_properties(ipc_toolkit PROPERTIES FOLDER "SRC") get_target_property(IPC_TOOLKIT_SOURCES ipc_toolkit SOURCES) -source_group(TREE "${PROJECT_SOURCE_DIR}" FILES ${IPC_TOOLKIT_SOURCES}) +# source_group(TREE) hard-errors on any file that is not under ROOT, and with +# CUDA enabled the target owns generated wrappers under the binary dir (see +# ipc_toolkit_shared_device_sources.cmake). The binary dir is outside the +# source tree for any out-of-source build, which is every consumer that pulls +# us in via CPM/FetchContent, so split the two groups apart. +set(IPC_TOOLKIT_IN_TREE_SOURCES "") +set(IPC_TOOLKIT_GENERATED_SOURCES "") +foreach(source IN LISTS IPC_TOOLKIT_SOURCES) + cmake_path(IS_PREFIX PROJECT_SOURCE_DIR "${source}" NORMALIZE is_in_tree) + if(is_in_tree) + list(APPEND IPC_TOOLKIT_IN_TREE_SOURCES "${source}") + else() + list(APPEND IPC_TOOLKIT_GENERATED_SOURCES "${source}") + endif() +endforeach() +source_group(TREE "${PROJECT_SOURCE_DIR}" FILES ${IPC_TOOLKIT_IN_TREE_SOURCES}) +if(IPC_TOOLKIT_GENERATED_SOURCES) + source_group("Generated" FILES ${IPC_TOOLKIT_GENERATED_SOURCES}) +endif() ################################################################################ # Dependencies @@ -317,11 +335,25 @@ if(IPC_TOOLKIT_WITH_SIMD) # unit's own flags. ipc/utils/simd.hpp exposes batch types in the public API, # and a consumer compiled without these flags would name a different type # than the one instantiated in the library, failing to link. - # NOTE: Restricted to C++ sources because these are MSVC-style /arch flags - # that nvcc would otherwise treat as input files, and device code has no use - # for CPU SIMD anyway. + # NOTE: Restricted to C++ sources because nvcc would treat a bare /arch or + # -march flag as an input file, and device code has no use for CPU SIMD. target_compile_options(ipc_toolkit PUBLIC "$<$:${SIMD_CXX_FLAGS}>") + if(IPC_TOOLKIT_WITH_CUDA) + # The *host* pass of a .cu still needs them. A shared device source is + # 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 + "$<$,$>:${SIMD_CUDA_HOST_FLAGS}>") + endif() + # Link against cross-platform xsimd library include(xsimd) target_link_libraries(ipc_toolkit PUBLIC xsimd::xsimd) @@ -329,33 +361,48 @@ if(IPC_TOOLKIT_WITH_SIMD) # Disable vectorization in Eigen since I've found it to have alignment issues. # NOTE: I don't know why this needs to be public, but it crashes if I make it private. target_compile_definitions(ipc_toolkit PUBLIC EIGEN_DONT_VECTORIZE=1) +endif() - if(IPC_TOOLKIT_WITH_CUDA) - # Pin Eigen's alignment so C++ and CUDA translation units agree on object - # layout. EIGEN_DONT_VECTORIZE alone yields align=0 under the host compiler - # but align=16 under nvcc, which forces 16-byte alignment when compiling - # CUDA. Fixed-max types (e.g. VectorMax) then get a different - # sizeof/alignof in each, which smashes the stack when one is returned by - # value across that boundary. PUBLIC because every TU exchanging Eigen - # objects with the library must agree, or it is an ODR violation. - # - # NOTE: We pin to 0, not 16. Any non-zero pin is a cap, and Eigen applies - # it without checking the scalar's own alignment -- so a cap below the - # natural alignment of an xsimd batch (32 bytes on AVX2, 64 on AVX-512) - # makes Eigen ask to under-align an array of them, which clang rejects - # outright. Since EIGEN_DONT_VECTORIZE is already set, alignment buys us - # nothing here anyway; 0 means "emit no alignas", which is both what the - # host compiler already does and the one value that cannot conflict with - # any scalar on any architecture. - target_compile_definitions(ipc_toolkit PUBLIC - EIGEN_MAX_STATIC_ALIGN_BYTES=0 - EIGEN_MAX_ALIGN_BYTES=0) - endif() +if(IPC_TOOLKIT_WITH_CUDA) + # Pin Eigen's alignment so C++ and CUDA translation units agree on object + # layout. Eigen picks its alignment per translation unit from the flags that + # unit was compiled with, and nvcc never sees the host compiler's -march + # flags -- so the two sides disagree whenever anything makes them differ. + # Two ways that happens here: EIGEN_DONT_VECTORIZE yields align=0 under the + # host compiler but align=16 under nvcc, and -march=native yields align=32 + # (AVX) or 64 (AVX-512) in C++ against nvcc's 16. Fixed-max types (e.g. + # VectorMax) then get a different sizeof/alignof in each, which + # smashes the stack when one is returned by value across that boundary. + target_compile_definitions(ipc_toolkit PUBLIC + EIGEN_MAX_STATIC_ALIGN_BYTES=0 + EIGEN_MAX_ALIGN_BYTES=0) endif() # For MSVC, do not use the min and max macros. if(MSVC) target_compile_definitions(ipc_toolkit PRIVATE NOMINMAX) + + # Compile CUDA translation units as UTF-8. + # + # The bundled fmt hard-fails on MSVC unless narrow literals are UTF-8: + # static_assert(!FMT_UNICODE || use_utf8, "Unicode support requires + # compiling with /utf-8"). It detects this by encoding U+00A7 and checking + # for the two-byte form, which the active code page does not produce. + # + # spdlog already hands us /utf-8 PUBLIC, but only for COMPILE_LANGUAGE:CXX, + # so a .cu never receives it. Our fmt patch (cmake/patches/) keeps the + # assert quiet there by hardcoding the probe true under __CUDACC__, which + # silences the error without making it correct: fmt then skips its + # vprint_mojibake fallback and writes code-page bytes as though they were + # UTF-8. This flag makes the hardcoded answer the true one. + # + # NOTE: -Xcompiler because options reach nvcc verbatim and it forwards no + # MSVC flags of its own. It reaches the host pass only -- the device + # front-end cannot be given /utf-8 at all, which is why the patch stays. + # PUBLIC so a consumer's own .cu files agree, and gated on nvcc because + # clang as the CUDA compiler has no -Xcompiler. + target_compile_options(ipc_toolkit PUBLIC + "$<$,$>:-Xcompiler=/utf-8>") endif() # Use C++17 diff --git a/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake b/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake index 5fd9a5468..ba6e9be15 100644 --- a/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake +++ b/cmake/ipc_toolkit/ipc_toolkit_shared_device_sources.cmake @@ -38,16 +38,12 @@ function(ipc_toolkit_target_shared_device_sources target) # Mirror the source's path (relative to the current source dir) under the # wrapper directory so that duplicate base names in different subfolders # (e.g. a/point.cpp and b/point.cpp) get distinct wrapper paths instead of - # colliding on a single flattened name. + # colliding on a single flattened name. Swapping the extension on the + # relative path covers a bare filename and a subdirectory path alike. file(RELATIVE_PATH src_rel "${CMAKE_CURRENT_SOURCE_DIR}" "${src_abs}") - get_filename_component(rel_dir "${src_rel}" DIRECTORY) - get_filename_component(src_name_we "${src_rel}" NAME_WE) - - set(wrapper_dir "${CMAKE_CURRENT_BINARY_DIR}/cuda_device_wrappers") - if(rel_dir) - set(wrapper_dir "${wrapper_dir}/${rel_dir}") - endif() - set(wrapper "${wrapper_dir}/${src_name_we}.cu") + string(REGEX REPLACE "\\.[^./\\\\]*$" "" src_rel_we "${src_rel}") + set(wrapper + "${CMAKE_CURRENT_BINARY_DIR}/cuda_device_wrappers/${src_rel_we}.cu") # file(GENERATE) only rewrites when the content changes, so this does not # trigger spurious rebuilds on reconfigure. The wrapper #includes the @@ -61,20 +57,15 @@ function(ipc_toolkit_target_shared_device_sources target) ") # Source file properties are directory-scoped and are read from the - # directory that created the target, which is not this subdirectory. Set - # them in both scopes (TARGET_DIRECTORY for the generator, the plain call - # for anything reading them here) and use absolute paths so the lookup - # resolves from either scope. Without TARGET_DIRECTORY the defines are - # silently dropped and both passes instantiate every scalar, so nvcc sees - # the xsimd batch instantiations and fails. + # directory that created the target, which is not this subdirectory, so + # they must be set with TARGET_DIRECTORY and with absolute paths. + # Without it the defines are silently dropped and both passes instantiate + # every scalar, so nvcc sees the xsimd batch instantiations and fails. set_source_files_properties("${wrapper}" TARGET_DIRECTORY ${target} PROPERTIES GENERATED TRUE COMPILE_DEFINITIONS IPC_TOOLKIT_DEVICE_SCALARS_ONLY) - set_source_files_properties("${wrapper}" PROPERTIES - GENERATED TRUE - COMPILE_DEFINITIONS IPC_TOOLKIT_DEVICE_SCALARS_ONLY) list(APPEND sources_to_add "${wrapper}") # The same .cpp, compiled by the host compiler for the scalars nvcc @@ -82,8 +73,6 @@ function(ipc_toolkit_target_shared_device_sources target) set_source_files_properties("${src_abs}" TARGET_DIRECTORY ${target} PROPERTIES COMPILE_DEFINITIONS IPC_TOOLKIT_HOST_SCALARS_ONLY) - set_source_files_properties("${src_abs}" PROPERTIES - COMPILE_DEFINITIONS IPC_TOOLKIT_HOST_SCALARS_ONLY) list(APPEND sources_to_add "${src_abs}") else() list(APPEND sources_to_add "${src}") diff --git a/cmake/recipes/spdlog.cmake b/cmake/recipes/spdlog.cmake index d3be8c70c..5e083f888 100644 --- a/cmake/recipes/spdlog.cmake +++ b/cmake/recipes/spdlog.cmake @@ -1,6 +1,18 @@ # 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/build/c++.rst b/docs/source/build/c++.rst index 596ecbc07..23fd828cd 100644 --- a/docs/source/build/c++.rst +++ b/docs/source/build/c++.rst @@ -63,4 +63,6 @@ Dependencies **All required dependencies are downloaded through CMake** depending on the build options, and are built automatically when you build the IPC Toolkit. You do not need to install them separately. +Fetching them does need ``git`` and the ``patch`` command-line utility on your ``PATH`` (one dependency is fetched with a patch applied). Both ship with the usual developer tool sets -- ``build-essential`` on Debian/Ubuntu, the Xcode command line tools on macOS, Git for Windows on Windows -- but a minimal container image may need them installed. + A full list of dependencies can be found on the `dependencies page `_. diff --git a/src/ipc/config.hpp.in b/src/ipc/config.hpp.in index 56f6e414d..e901e9eb4 100644 --- a/src/ipc/config.hpp.in +++ b/src/ipc/config.hpp.in @@ -22,6 +22,17 @@ #cmakedefine IPC_TOOLKIT_WITH_MESHFEM_SPARSE // #define IPC_TOOLKIT_DEBUG_AUTODIFF +// CPU SIMD is a host-only facility, so retract the switch for nvcc. xsimd is +// not on nvcc's include path and ipc/utils/simd.hpp defines no batch types +// there, so any `#ifdef IPC_TOOLKIT_WITH_SIMD` block naming SimdBatch would +// otherwise fail to compile the moment its file becomes a shared device +// source. Retracting once here means those blocks need no CUDA-specific +// spelling. Note that this is a property of the translation unit, not of the +// build: host TUs of a CUDA-enabled build still see SIMD as on. +#if defined(IPC_TOOLKIT_WITH_SIMD) && defined(__CUDACC__) +#undef IPC_TOOLKIT_WITH_SIMD +#endif + // Execution-space annotation for functions shared between C++ and CUDA. // IPC_TOOLKIT_HOST_DEVICE expands to `__host__ __device__` under nvcc and to // nothing otherwise, so the same declaration is valid in host and device code. diff --git a/src/ipc/distance/point_line.cpp b/src/ipc/distance/point_line.cpp index c5f2d7b5d..73293237d 100644 --- a/src/ipc/distance/point_line.cpp +++ b/src/ipc/distance/point_line.cpp @@ -411,7 +411,7 @@ IPC_TOOLKIT_HOST_DEVICE void point_line_distance_hessian_3D( IPC_INSTANTIATE_POINT_LINE_AUTOGEN(float); IPC_INSTANTIATE_POINT_LINE_AUTOGEN(double); #endif -#if defined(IPC_TOOLKIT_WITH_SIMD) && IPC_TOOLKIT_INSTANTIATE_HOST_SCALARS +#ifdef IPC_TOOLKIT_WITH_SIMD IPC_INSTANTIATE_POINT_LINE_AUTOGEN(SimdBatch); IPC_INSTANTIATE_POINT_LINE_AUTOGEN(SimdBatch); #endif diff --git a/src/ipc/utils/simd.hpp b/src/ipc/utils/simd.hpp index 3955c9e66..e4ad700f5 100644 --- a/src/ipc/utils/simd.hpp +++ b/src/ipc/utils/simd.hpp @@ -4,7 +4,7 @@ #include #include -#if defined(IPC_TOOLKIT_WITH_SIMD) && !defined(__CUDACC__) +#ifdef IPC_TOOLKIT_WITH_SIMD #include #endif @@ -109,7 +109,7 @@ inline auto select_lazy(const Mask& mask, F&& value, Rest&&... rest) } // namespace ipc -#if defined(IPC_TOOLKIT_WITH_SIMD) && !defined(__CUDACC__) +#ifdef IPC_TOOLKIT_WITH_SIMD namespace Eigen {