Skip to content

Cache CMake configure results and precompile the DaCe runtime header - #2453

Merged
tbennun merged 34 commits into
mainfrom
cmake-build-cache
Aug 5, 2026
Merged

Cache CMake configure results and precompile the DaCe runtime header#2453
tbennun merged 34 commits into
mainfrom
cmake-build-cache

Conversation

@ThrudPrimrose

@ThrudPrimrose ThrudPrimrose commented Jul 21, 2026

Copy link
Copy Markdown
Collaborator

Precompiling the DaCe runtime header and caching CMake configure + build results lets repeated builds skip work that can't change between two SDFGs compiled the same way. 10s → 0.5s per repeat on my machine. It also fixes several compiler-flag bugs surfaced along the way.

Changes

  1. Every TU re-parses <dace/dace.h> (most of a small kernel's compile) -> precompiled header, force-included into host sources only.
  2. Every SDFG re-runs project() detection + find_package -> configure cache: transplant CMakeCache.txt.
  3. CMake re-derives identical compile/link commands per SDFG -> ninja -t compdb record -> replay (no CMake/Ninja); Ninja default, jobs = build_jobs.
  4. Default -ffast-math unsafe. It breaks isnan/isinf, NPBench LU ⇒ granular safe subset -fno-math-errno -fno-trapping-math -fno-signed-zeros -freciprocal-math.
  5. Configured build type silently overridden (-O3 args beat CMake's -O2) -> drop -O, default build type comes from the cmake config.
  6. Flags handed to wrong compiler -> empty cpu.executable → CMake's $CXX is used, which can be different than c++ if $CXX set to something else -> always pass DaCe's compiler as -DCMAKE_CXX_COMPILER (never fall back to $CXX), pin to CMake's id.
  7. PAPI/LIKWID append flags to global config -> leak per-SDFG (-lpapi everywhere) -> scoped library environments (just like other proper libraries).
  8. PAPI vec-report flag hardcoded to GCC spelling (others would fail compiling) -> chosen in CMake per COMPILER_ID as a papi_vectorization.cmake CMake module.
  9. Enforce no cmake extensions other than the flag provided (avoids defaulting to gnu++20 on GCC compilers and accidentally generating gnu-only code AFAIK)

Safety

Since I could not test on Windows, the PCH and ninja is not enabled there.

Building an SDFG re-runs work that cannot differ between two programs compiled
the same way. A fresh build folder repeats the project() compiler and ABI
detection plus every find_package, and every generated translation unit parses
<dace/dace.h> from scratch -- which is most of the compile time of a small
kernel.

Both results are now cached outside the build folder, keyed on everything that
can change them, so only the first SDFG pays:

- The configure cache transplants CMakeCache.txt (the find_package results) and
  CMakeFiles/<version>/ (the compiler detection) into a new build folder.
  Seeding one without the other is close to worthless, since each still forces
  the other half of the work. CMake refuses a cache file found in a directory
  other than the one it was written in, so CMAKE_CACHEFILE_DIR is retargeted at
  the new folder; without that the configure aborts instead of being reused.
  Only the compiler-detection subdirectory of CMakeFiles/ is published -- the
  rest holds the program's object files.

- The precompiled header is built once per compiler and flag set and
  force-included into generated host sources only, since forcing <dace/dace.h>
  into an environment's auxiliary translation unit, or a .cu, can break a file
  that never asked for it.

Both are advisory: a mismatched header is ignored by the compiler and a stale
configure seed is simply re-derived, so a bad or missing entry costs speed and
never correctness. compiler.configure_cache and compiler.precompiled_header
turn them off.

Also exports compile_commands.json, which tooling can use and which is how the
test checks the header actually reaches the compile line.

Measured on one small CPU kernel, five builds each, interleaved: 9.26s -> 2.55s
median per SDFG after the first.
CXX_STANDARD leaves CXX_EXTENSIONS on, so CMake was emitting -std=gnu++20 for
the generated library. Turn extensions off and build the precompiled header
with the matching -std=c++NN; a header precompiled under a different dialect is
silently declined, which would cost the speedup with no diagnostic.
@tbennun
tbennun self-requested a review July 22, 2026 02:52
Comment thread dace/codegen/CMakeLists.txt
Comment thread dace/codegen/compiler.py Outdated
CMake re-derives the same compile and link commands for every SDFG built with
the same configuration. `ninja -t compdb` reports all of them -- unlike
compile_commands.json, which holds compiles only -- so the first build of a
given shape can record them and every later one just runs them, with no CMake
and no Ninja.

Median per SDFG over five distinct programs, interleaved: 2.09s before, 0.53s
with all three caches on.

Recordings are templated on the program name and its two folders, and are only
replayed when the generated sources they compile are exactly the ones the
program has; anything else falls back to a full build. Ordering comes from the
graph rather than from file names, since under CUDA separable compilation
cmake_device_link.o is an object by name but consumes the objects.

Ninja is now the generator where one is available, and `cmake --build` is
bounded by compiler.build_jobs rather than defaulting to serial (Make) or to
every core (Ninja).

Addresses review feedback: says why CXX_EXTENSIONS is off, renames the
precompiled-header helper to prepare_precompiled_header, prefers RAM-backed
storage with a reserved fallback inside the build folder, and disables all
three caches on Windows, which none of this is tested on.

Three tests named their SDFGs in ways that put separate programs in one build
folder -- second-resolution timestamps, and parametrizations differing only in
generated code. Faster builds make them collide; they now get distinct names.
…y own its flags

compiler.cpu.args shipped -O3 and -ffast-math while CMake appended the build
type's own -O2 -g -DNDEBUG after it. Two optimization levels on one command
line is unspecified; every compiler we target happens to take the last, so the
configured build type quietly lost to a flag string that was not supposed to
carry a level at all. Drop -O from the args, default build_type to Release, and
pin the split with a test that reads the compiler's own compile_commands.json.

-ffast-math goes with it. Its finite-math and reassociation parts are too
aggressive -- NPBench's LU breaks under them, its initialization not being
diagonally dominant -- so the args now carry the granular subset that still lets
the vectorizer through libm calls: -fno-math-errno -fno-trapping-math
-fno-signed-zeros -freciprocal-math. Same for nvcc's --use_fast_math.

Those switches are GCC/Clang spellings, and nvc++ rejects four of them outright,
so an NVHPC user could not compile at all. config.py resolves default_<platform>
at import because the OS is always knowable; a compiler is not, since
compiler.cpu.executable defaults to empty and CMake picks. So the family is
resolved where the flags are assembled instead: compiler_family detects it from
the macros the compiler predefines -- checking __NVCOMPILER and __clang__ before
__GNUC__, which all three define -- and reads a default_<family> sibling key. An
explicitly set args string is passed through untouched.

Also write compile_commands.json when a build is replayed from the command
cache. A replay never runs CMake, so the database CMake would have emitted was
missing on exactly the builds that hit the cache -- clangd stopped working on
generated code, and the record of how a translation unit was compiled vanished.
The recipe already holds every compile.

Portability test parametrizes over the compilers present on the box, so an
absent toolchain yields no test case rather than a skip. CI carries no NVHPC:
the nvc++ leg runs only where the SDK is installed.
Six ufunc tests read compiler.cpu.args, swapped -ffast-math for
-fno-finite-math-only and restored it afterwards, because -ffinite-math-only
told the compiler no operand is ever inf or nan -- which is exactly what
isfinite/isinf/isnan are asked to detect. The shipped args no longer carry
-ffast-math, so the branch never fires and those tests were reading a config
they no longer patch. Removing it lets them exercise the flags DaCe actually
ships, which is the thing worth testing. The debug prints go with it.

The HIP example in doc/optimization/gpu.rst still advertised -O3 -ffast-math;
the level comes from the build type now, and the granular flags replace
fast-math.
ThrudPrimrose added a commit that referenced this pull request Jul 22, 2026
Resolving #2453's CMakeLists.txt conflict took the PR's file wholesale, which
silently undid three pieces of work this branch already had. The PR predates
them, so its version could not have known about any of it.

The CUDA code generator defaults to the experimental implementation, so dropping
"experimental_cuda" from the target test meant a generated CUDA file stopped
being recognised as a CUDA target. It never got LANGUAGE CUDA, was compiled as
plain C++, and every CUDA symbol it emits -- dace::cuda::Context, cudaMalloc,
DACE_GPU_CHECK, gpuStream_t, gpu_streams -- became undeclared. That is the whole
failure: test_gpu_callback compiled cleanly before the merge and stopped after.

CMAKE_CUDA_STANDARD went back to a hardcoded 17, so device code and the nvcc host
pass silently lost the dialect DACE_CPP_STANDARD selects for the host side.

The auxiliary-sources block was dropped entirely. It compiles the .cu wrappers
contributed by library environments for libnodes whose host-side tasklet cannot
emit a kernel launch directly, and marks them LANGUAGE CUDA.

CUDA::nvtx3, which the PR adds, is kept.

None of this surfaced as a conflict or a build error -- only as a C++ compile
failure inside a test marked @pytest.mark.gpu, which most sweeps deselect.
callback_autodetect_test is back to 44 passed / 2 skipped under the legacy CPU
generator, matching the pre-merge baseline.
compiler.cpu.executable is DaCe's knob and wins, but when it is empty DaCe
passes no -DCMAKE_CXX_COMPILER and CMake falls back to $CXX. Detection probed
'c++' in that case, so CXX=nvc++ produced a build where CMake picked nvc++ and
the flags had been chosen for GCC; the configure died on
'nvc++-Error-Unknown switch: -fno-math-errno'.

Stopgap: the compiler is still inferred rather than read back from CMake, which
also consults toolchain files and its own search order.
The precompiled header is compiled outside CMake from DaCe's own copy of
CMAKE_CXX_FLAGS_<CONFIG>, and a header built with flags the translation unit
does not have is silently declined -- the build succeeds, just without the
speedup the cache exists for. That copy described GNU only, and NVHPC differs in
every configuration: Release adds -fast, RelWithDebInfo omits -DNDEBUG and
spells debug info -gopt, MinSizeRel is -O2 -s rather than -Os. So on nvc++ the
PCH never applied.

Nothing here is passed to a compiler and no default changes; this is a
prediction of what CMake will append, now made per family and verified by asking
CMake rather than reading its documentation.

The test is parametrized by family for the same reason: checking only the host's
default compiler passed while all four NVHPC entries were wrong, which is
precisely the case a PCH flag test exists to catch.
Two detectors decide two halves of one build: DaCe's picks the flags, CMake's
picks the compiler they are handed to. They agree today by hand. If they ever
diverge, flags chosen for one compiler reach another -- the failure mode that
made CXX=nvc++ die on 'Unknown switch: -fno-math-errno'. Family names are
CMake's ids lowercased, so the comparison is exact and needs no mapping table
that could rot.
cpu_args() treated any difference from the shipped default as a deliberate user
override and handed the value back untouched. But appending is how DaCe itself
adds flags: LIKWID appends -DLIKWID_PERFMON -fopenmp so its headers stop
compiling to no-ops, PAPI appends a vectorization-report flag, and Config.append
is literally current += value. So enabling either instrumentation on an NVHPC
host silently restored the GCC/Clang defaults and the configure died on
nvc++-Error-Unknown switch: -fno-math-errno.

Substitute the family default for the shipped one as a prefix instead, which
distinguishes the two cases exactly: an append is default+suffix and keeps its
suffix, while a hand-written value does not start with the default and is still
returned untouched.
LIKWID and PAPI added their compile flags and libraries with
Config.append('compiler', 'cpu', ...), which mutates process-wide state and
never restores it. Both did it from on_sdfg_begin, so the value grew once per
SDFG compiled -- five instrumented programs left five copies of the same flag on
the compile line -- and every later SDFG in the session inherited those flags
and -lpapi whether or not it was instrumented.

DaCe already has the mechanism: a library environment's cmake_compile_flags and
cmake_libraries are a deduplicated set, scoped to the SDFGs that declare it.
LIKWID even had the class already, with an empty cmake_compile_flags list, while
the provider appended to the config beside it.

The CPU and GPU providers get one environment each, because they need different
defines: likwid.h compiles its marker macros to no-ops without LIKWID_PERFMON,
the GPU path needs LIKWID_NVMON instead, and defining both would activate a
marker API whose initialization the generated code never emits. Both depend on
LIKWID for the headers and library.

-fopenmp is dropped rather than moved: CMakeLists.txt already does
find_package(OpenMP REQUIRED) and links OpenMP::OpenMP_CXX, so appending it only
duplicated the flag.

PAPI's -fopt-info-vec-optimized-missed=../perf/vecreport.txt is deleted outright
rather than moved to its environment. It is a GCC-only spelling writing to a
build-relative path, and nothing in DaCe ever read the report -- the whole path
was write-only. The instrumentation.papi.vectorization_analysis key that gated
it goes too, since it fed nothing else.

The tests assert against the provider sources and the environment classes rather
than by running the providers, since neither library is installed on CI -- and
on a machine without them the provider returns early, so a reintroduced append
would go unnoticed exactly where CI runs.
The report flag is spelled differently by every compiler -- GCC
-fopt-info-vec-optimized-missed, Clang -Rpass=loop-vectorize, NVHPC -Minfo=vect,
IntelLLVM -qopt-report -- and the hardcoded GCC spelling was simply rejected by
the others, so the analysis was GCC-only in practice.

CMake is what picks the compiler, so the choice is made there, from
CMAKE_CXX_COMPILER_ID, via a cmake fragment the PAPI environment contributes
through cmake_files. That avoids a second detector in Python having to predict
what CMake is about to select, which is exactly how flags chosen for one
compiler end up handed to another. add_compile_options places the flags after
CMAKE_CXX_FLAGS_<CONFIG>, so the build type cannot override them.

cmake_files is evaluated per build, so toggling
instrumentation.papi.vectorization_analysis between builds is honoured instead
of being baked in at the first instrumented SDFG the way the old one-shot global
append was.

Only GCC writes the report to perf/vecreport.txt; the others emit remarks on
stderr and so land in the build log.
Shorten the docstrings and comments added across the build-cache work, dedupe
the two LIKWID marker environments behind a shared base, drop the dead
PLACEHOLDERS constant, and rewrite papi_vectorization.cmake's preamble (the
generator expressions stay: the per-language guard keeps host-only flags off
nvcc for programs with .cu files).

Behavior change: DaCe now always passes -DCMAKE_CXX_COMPILER, resolved from
compiler.cpu.executable, else $CXX, else c++. Previously it was passed only
when compiler.cpu.executable was set, which forced compiler_family to guess
which compiler CMake would pick. Pinning it makes family detection exact.
ThrudPrimrose added a commit that referenced this pull request Jul 23, 2026
The PR #2453 merge into extended left two copies each of precompiled_header,
configure_cache and build_jobs under `compiler` (pyyaml silently keeps the last,
so both defaults were already true, but duplicate keys are a latent defect).
Kept the better-documented copies and the unique command_cache entry; all three
cache knobs remain enabled by default.
ThrudPrimrose and others added 2 commits July 23, 2026 02:11
Compress the multi-line docstrings across the four build-cache/flag/
instrumentation test files to one line each; no assertions changed.

Remove DACE_cache=single from every gpu-ci step: single-folder mode shares
one build directory across SDFGs, which defeats the recorded-build replay
(each replay wants a fresh folder, as the tests pin cache=name to get).
@ThrudPrimrose
ThrudPrimrose marked this pull request as ready for review July 23, 2026 13:48
@ThrudPrimrose
ThrudPrimrose marked this pull request as draft July 23, 2026 14:45
ThrudPrimrose and others added 3 commits July 24, 2026 14:17
… installed mpi4py

The bare directory tests/python_frontend/mpi4py/ (no __init__.py) is a namespace
package. During full-tree collection pytest prepends tests/python_frontend onto
sys.path, so 'import mpi4py' resolves to this empty dir and shadows the installed
mpi4py -> 'cannot import name MPI from mpi4py (unknown location)'. build_cache_test's
test_distributed_and_local_builds_interleave imports mpi4py inside the test body (after
collection, shadow active) and is the only 'gpu and mpi' test, so it alone surfaced it.
Rename to mpi4py_frontend to remove the top-level name collision.
@ThrudPrimrose
ThrudPrimrose marked this pull request as ready for review July 26, 2026 00:41
ThrudPrimrose and others added 2 commits July 27, 2026 11:44
The default cpu args carry -march=native, and cache_key can only see that as a
literal string -- it cannot know which instruction set the compiler selected.
Identical inputs on two different CPUs therefore produce the same key.

That is harmless while build_cache_root() resolves to /dev/shm or /tmp, both
node-local. It stops being harmless on the two paths that leave a node:
DACE_BUILD_CACHE_DIR pointed at shared scratch (which is exactly what someone
does to get hits across jobs) and the default_build_folder fallback, which on a
cluster is typically a parallel file system. A configure cache transplanted
that way replays one host's try_compile answers on another.

Fold a host identity into cache_key: the CPU flag list is what the compiler's
native detection reads, hashed with the model name for the tuning half, falling
back to the platform triple where there is no /proc. A coarse key
over-separates, which costs a miss; it cannot merge two hosts that differ.
@ThrudPrimrose
ThrudPrimrose marked this pull request as draft July 27, 2026 10:43
An mpi4py wheel carries no MPI of its own: it dlopens libmpi when the MPI submodule is
first imported. On a machine with the package but no MPI runtime that import raises
RuntimeError("cannot load MPI library"), not ImportError.

Two frontend sites guarded only for ImportError -- resolve_names around MPIResolver, and
ProgramVisitor.defined, which every operand type-check reaches. The RuntimeError escaped
both, so a stray `pip install mpi4py` made EVERY @dace.program fail to parse, including
programs that never mention MPI.

Both now ask one cached mpi4py_is_usable() helper, which guards the import alone so a
genuine error inside the visitor still surfaces. Importing the submodule does not
initialize MPI, so the parse path stays side-effect-free.

The two CI jobs that install mpi4py verify it right after installing, because this failure
mode otherwise appears as unrelated failures across the whole suite rather than at its
cause.
ThrudPrimrose and others added 5 commits July 27, 2026 15:56
The cache is machine-global, so two DaCe checkouts sharing a compiler shared one
dace_prewarm.h.gch and the second compiled against the first one's headers. The
mtime guard cannot catch it: it walks this tree's runtime and compares against a
header built from another's, so it passes while being wrong.
Relocating the runtime is what a second checkout looks like; without the include
path in the key both trees get one .gch and the second compiles against foreign
declarations.
Ninja's compdb declares the program link before the static-library
archive it consumes (the nanobind extension links libnanobind-static.a),
so replaying in declaration order failed at the link and every build
fell back to CMake. An entry whose command mentions another pending
entry's output now waits for it.

Co-authored-by: Philip Mueller <claude@philipmueller.ch>
@ThrudPrimrose
ThrudPrimrose requested a review from tbennun August 4, 2026 09:27
@ThrudPrimrose
ThrudPrimrose marked this pull request as ready for review August 4, 2026 09:27
@ThrudPrimrose
ThrudPrimrose marked this pull request as draft August 4, 2026 09:48
The recorded-build cache only publishes under the Ninja generator, so an
environment without ninja never replays. `ninja_requires` was computed but
never reached `install_requires`.
@ThrudPrimrose
ThrudPrimrose marked this pull request as ready for review August 4, 2026 14:25
…nstexpr functions

Only the kernel bodies in a generated ``.cu`` are device code; the rest of the file is host code --
the state struct, the launchers, the stream setup -- and both it and the ``.cpp`` include the same
header-only runtime. The two were built against different standards, C++17 for CUDA against C++20
for the host, with the state struct defined in both translation units. Take the standard from
DACE_CPP_STANDARD so the two cannot drift again. hipcc was never given one at all and fell back to
its own default, currently gnu++17.

``--expt-relaxed-constexpr`` goes next to the standard rather than into the configurable nvcc
arguments, since device code calling a host constexpr function must not stop compiling because
someone overrode those.
CMake hands a device compiler nothing from CMAKE_CXX_FLAGS -- CUDA and HIP are languages of their
own, and the automatic propagation people remember belonged to the removed FindCUDA module. So the
host half of every ``.cu`` was built without the flags the ``.cpp`` got, and the inline functions
both take from the shared runtime header came out with different arithmetic in each object: which
one survives is the linker's choice.

Derived from ``compiler.cpu.args`` rather than listed again, so the two cannot drift, and through
``compiler_family`` so an nvhpc host compiler still gets its own set. Warnings that only fire inside
the device headers stay behind, as does ``-fPIC``, which CMake already adds for a shared library
(Compiler/NVIDIA.cmake: ``-Xcompiler=-fPIC``). One ``-Xcompiler`` per flag, since nvcc splits the
comma-separated form on commas; hipcc is a single driver and takes them unwrapped.
return len(LIKWID.cmake_libraries()) > 0


class LIKWIDMarkers:

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nice find!

Comment thread tests/codegen/map_launch_test.py Outdated
@tbennun
tbennun enabled auto-merge August 5, 2026 05:38
@tbennun
tbennun added this pull request to the merge queue Aug 5, 2026
Merged via the queue into main with commit 2e54d8d Aug 5, 2026
16 checks passed
@tbennun
tbennun deleted the cmake-build-cache branch August 5, 2026 07:05
ThrudPrimrose added a commit that referenced this pull request Aug 5, 2026
Resolved per hunk, preferring main's design and keeping our additions on top.

Taken from main: the CMake configure/command cache and precompiled header
(#2453) throughout compiler.py, CMakeLists and the schema; the GPU codegen fixes
(#2480), whose state.py hunk is the reviewed evolution of work that originated
here; the optimized LiftTrivialIf (#2460); dropping -fPIC from the cpu/hip
default args, since CMake supplies it via CMAKE_POSITION_INDEPENDENT_CODE; and
the CI install blocks, which are supersets of ours.

Kept ours, where we add on top of that design:
  - cudacommon.cuh -- our DACE_GPU_CHECK_RETURN/_RETURN_VAL need a bool, and
    report_error() is already main's macro body factored into a function.
  - reduce.py -- the per-stream CUB scratch pool. Main's single per-node
    __cub_storage is shared by every stream that runs the node, so concurrent
    streams would race on one workspace.
  - CMakeLists static_archive target, config_schema build_mode, setup.py's
    ordered-set/pygments/numpy pin, gpu-ci's cutensor include/lib wiring,
    cpu.py's CodeObject import (still used four times here).

Two things the merge itself surfaced:
  - config_schema had build_jobs, precompiled_header and configure_cache defined
    on BOTH sides; keeping both would have been a duplicate YAML key. Main's
    definitions are kept, ours dropped.
  - get_scratch() returned a null pointer for a zero-byte request, because
    0 > 0 skips the allocation and cudaMalloc(0) yields null anyway. CUB reads a
    null workspace as "only report the size", so the reduction became a silent
    no-op leaving the output untouched -- the same failure #2480 guards on its
    own path. It now allocates a one-byte floor and reallocates whenever the
    entry is null. The default/null stream needs no special case: pool_map keys
    on the stream, so 0 is simply another key.

Gates: import clean, config_schema parses, ruff/yapf/pre-commit clean;
124 passed across the conflicted test files, 47 across config/frontend.
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.

3 participants