Adding LAMMPS computes for MACE descriptors and descriptor gradients. - #70
Adding LAMMPS computes for MACE descriptors and descriptor gradients.#70Fraser-Birks wants to merge 22 commits into
Conversation
|
Definitely open to this, thanks.
Are you sure this is what you'd want? I would have guessed concatenated invariants from all layers. |
I'm not sure! Is there a precedent for what makes a 'MACE descriptor'? |
|
Maybe take a look at what the torch ASE calculator returns? https://github.com/ACEsuit/mace/blob/main/mace/calculators/mace.py#L580 |
Yep - it's the concatenated invariants from all layers. I'll update the computes! |
|
Just dropped by this issue. I think this is exciting because one potentially could use these descriptor as the CVs for enhanced MD (metadynamics, umbrella sampling, biased MD, etc.) Ideas similar to what has been done somewhere lately https://pubs.acs.org/doi/full/10.1021/acs.jctc.5c01767 |
|
Descriptors now include the first layer invariants, exactly matching what's returned by However, to get descriptors that matched the python I had to undo some of @wcwitt 's clever fusion with the first layer invariants. That means I'm now exporting and reading an extra property from the model json file. As that property doesn't exist in the test json files, I've broken the C.I. @wcwitt - any chance I could request that you re-generate the .json files used in the tests with the new version of |
Profiling the silicon_surface SKMD system on Avon showed fix_skmd's
END_OF_STEP hook invokes compute_peratom() on the (previously CPU-only)
macedesc descriptor compute every single MD step while the fix is
active -- not intermittently -- which explains most of the observed
483% MD-block overhead once switching the pair style to /kk made
everything else fast.
Port the descriptor compute to a device-resident version, following
pair_symmetrix_mace_kokkos's own template closely:
- Two-mode dispatch (no_domain_decomposition / mpi_message_passing),
auto-selected by comm->nprocs exactly like the pair style's
settings(). Single-rank runs never need H1 forward-comm at all --
neighbor indices resolve straight to their owning local atom via the
atom map (AtomKokkos::map_kokkos), so H1/H2 stay purely node-indexed
and device-resident throughout. Multi-rank runs get the real
forward-comm path (pack/unpack_forward_comm_kokkos), copied from the
pair style's existing implementation.
- Graph-build kernels (node/edge counting, first_neigh scan, edge
population) copied near-verbatim from
pair_symmetrix_mace_kokkos::compute_{no_domain_decomposition,
mpi_message_passing}.
- linear_up_l0_inv (the l=0 invariant "restoration" matrix needed for
descriptors but not exposed by MACEKokkos, which only computes
energies/forces) loaded via a transient CPU MACE instance at
construction -- the same dependency the plain CPU compute already
uses, rather than a new direct JSON-parsing path.
- Per-atom output via a 2D Kokkos::DualView-backed array_atom
(memoryKK::create_kokkos), with a single modify+sync_host per
compute_peratom() call -- replacing what was an entirely CPU-resident
independent MACE forward pass running on every MD step.
Registers as symmetrix/mace/atom/kk (+ /kk/device, /kk/host); the
downstream `compute macedesc all symmetrix/mace/atom ...` command
string needs no changes since -sf kk auto-resolves it. install.sh
updated to symlink the new files into a target LAMMPS's src/KOKKOS/,
alongside the existing pair_symmetrix_mace_kokkos entries.
Unverified: no Kokkos-enabled LAMMPS build available to compile
against in the environment this was written in. Needs a real build +
bit-identical-vs-CPU dump comparison on Avon before being trusted.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
…ix-public into compute_descriptors
1. ComputeStyle(key, Class) is a 2-arg preprocessor macro; the raw
ComputeSymmetrixMACEatomKokkos<LMPDeviceType,double> template
argument's comma got parsed as an extra macro argument ("passed 3
arguments, but takes just 2"). Alias to a single token first via
#define/#undef, exactly the trick pair_symmetrix_mace_kokkos.h
already uses for its own PairStyle(...) invocations -- missed
copying that part over.
2. nvcc: "The enclosing parent function ... cannot have private or
protected access within its class" for compute_no_domain_decomposition
and compute_mpi_message_passing, both of which contain Kokkos device
lambdas. CUDA extended __host__ __device__ lambdas require their
enclosing function to not be private/protected. Move both from
protected to public, matching where the pair style declares its own
equivalent methods.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
1. mace_kokkos.hpp had no include guard at all. It compiled fine while
only pair_symmetrix_mace_kokkos.h included it, but now
compute_symmetrix_mace_atom_kokkos.h also does -- LAMMPS's
aggregated style header pulls both into the same translation unit,
so mace_kokkos.hpp's `class MACEKokkos { ... }` got pasted twice:
"class template MACEKokkos has already been defined". Add
`#pragma once`, matching most of its own included sibling Kokkos
headers (cubic_spline_kokkos.hpp etc.).
2. `Kokkos::DualView<double**, DeviceType> k_array_atom` left its
layout to the DeviceType default, which is LayoutLeft on CUDA --
memoryKK::create_kokkos/destroy_kokkos require LayoutRight to alias
with the legacy `double **array_atom` pointer (see memory_kokkos.h's
static_assert). Name it explicitly:
`DualView<double**, Kokkos::LayoutRight, DeviceType>`, matching the
`k_params` pattern in the KOKKOS porting guide's pair-style skeleton.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
Full-Jacobian mode only (VJP mode isn't ported -- fix_skmd always needs the complete Jacobian, never uses VJP). Now that the descriptor compute is ported (compute_symmetrix_mace_atom/kk), this was confirmed as the larger of the two remaining per-outer-iteration costs on Avon (~9.05s vs the MD block's ~9.11s per iteration). Shares its graph-build and two-mode dispatch (no_domain_decomposition / mpi_message_passing) with compute_symmetrix_mace_atom_kokkos. The forward pass runs once per compute_peratom() call (through H2), then run_channel_loop() reuses those same forward activations across all 2*num_channels reverse passes -- exactly the efficiency point of adjoint/reverse-mode differentiation, and exactly what the CPU compute's "SLOW FULL JACOBIAN MODE" already does with plain C++ loops instead of device kernels: - H1-block (output columns [0, num_channels)): seed H1_adj directly from linear_up_l0_inv's column kc, then reverse_H1/M0/A0_scaled/A0. Never needs comm regardless of mode -- this leg never touches Phi1 (no cross-atom message passing), matching why the CPU compute doesn't call comm->reverse_comm here either. - H2-block (output columns [num_channels, 2*num_channels)): seed H2_adj to a one-hot vector, then reverse_H2/M1/A1_scaled/A1/Phi1 -- this is what actually populates H1_adj, potentially with ghost-atom contributions, so mpi_message_passing mode reverse-comms it (pack/unpack_reverse_comm_kokkos, same pattern pair_symmetrix_mace_kokkos uses for the same purpose) before continuing with reverse_H1/M0/A0_scaled/A0. - Each channel's edge-force output is scattered into array_atom via the same team-parallel/atomic-add pattern pair_symmetrix_mace_kokkos uses for its own force reduction, just writing into a per-channel column instead of atom->f. Registers as symmetrix/maced/atom/kk (+ /kk/device, /kk/host); no lammps_setup.py changes needed (-sf kk auto-resolves it). install.sh updated to symlink the new files. Applied lessons from the descriptor compute's build-fix round: the ComputeStyle macro alias trick, public (not protected) mode-dispatch methods for nvcc's extended-lambda restriction, and explicit Kokkos::LayoutRight on the 2D output DualView. Unverified: no Kokkos-enabled LAMMPS build available to compile against in the environment this was written in. Needs a real build + bit-identical-vs-CPU dump comparison on Avon -- this is a higher-risk port than the descriptor compute (adjoint seeding, edge-force sign conventions, comm timing are all easy to get subtly wrong) and should not be trusted without that comparison. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
test_descriptor: compares compute_symmetrix_mace_atom(/kk)'s per-atom [h1_restored | H2] output against a reference computed by driving symmetrix.MACE directly through the same forward-pass sequence (compute_R0/Y/A0/A0_scaled/M0/H1/R1/Phi1/A1/A1_scaled/M1/H2), applying the same linear_up_l0_inv correction the C++ compute uses -- same pattern symmetrix/test/test_mace.py already uses for the raw layer values, carried through to the actual descriptor. test_jacobian: spot-checks compute_symmetrix_maced_atom(/kk)'s Jacobian against central finite differences of that same reference descriptor's column sum -- full 2*num_channels x 3*num_atoms coverage would be correct but slow, so this samples a handful of columns across both the H1 and H2 blocks (enough to catch a wrong adjoint seed, a transposed sign, or a swapped block ordering). Both parametrized over cmdargs (plain CPU vs -k on -sf kk), the same pattern test_pair_symmetrix_mace.py already uses -- under -sf kk, `compute ... symmetrix/mace(d)/atom ...` auto-resolves to the /kk variant, so this is what actually exercises compute_symmetrix_mace_atom_kokkos / compute_symmetrix_maced_atom_kokkos. Known gap: pytest here always runs single-process, so comm->nprocs==1 and only the no_domain_decomposition code path gets exercised (both CPU and Kokkos) -- mpi_message_passing (multi-rank forward/reverse comm) is untested. Not wired into any CI workflow, per request. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
CPU cases pass (confirms the reference computation and comparison logic are correct). Kokkos cases failed with "Kokkos has been compiled with GPU-enabled backend but no GPUs are requested" -- bare -k on -sf kk (copied from test_pair_symmetrix_mace.py, which only needs to work on a host/serial-only Kokkos build) isn't enough on a GPU-enabled build; it needs an explicit GPU count. Add g 1 -pk kokkos newton on neigh half, matching skmd.lammps_setup.make_lammps's cmdargs. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
Same issue found while running the new compute tests: bare -k on -sf kk errors "Kokkos has been compiled with GPU-enabled backend but no GPUs are requested" on a GPU-enabled build -- these existing kokkos- parametrized cases apparently were never actually run against a real GPU build before. Add g 1 -pk kokkos newton on neigh half, matching skmd.lammps_setup.make_lammps's cmdargs, in all 4 occurrences. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
…Kokkos/CUDA bug Multiple create_atoms calls in one session (with atom_modify map active) hits an unresolved upstream LAMMPS bug: AtomKokkos::map_set_device() segfaults with cudaErrorIllegalAddress -- confirmed via bisection and matches a reported, maintainer-acknowledged issue: https://matsci.org/t/using-lammps-create-atoms-and-run-0-in-a-kokkos-cuda-interface/59054 The suggested workaround (drop atom_modify map) isn't usable here -- compute_symmetrix_mace_atom_kokkos/compute_symmetrix_maced_atom_kokkos require it (AtomKokkos::map_kokkos needs it for the no_domain_decomposition mode's neighbor->local-atom resolution). Switch build_lammps() to write a small LAMMPS data file via ASE's write_lammps_data and read_data it in a single call instead of three create_atoms calls -- sidesteps the bug, and incidentally matches skmd.lammps_setup.load_config_into_lammps's actual production code path more closely than create_atoms did anyway. (test_pair_symmetrix_mace.py has the same latent issue in its own kokkos-parametrized cases, but fixing that is left to the upstream maintainer rather than done here.) Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
ASE's write_lammps_data doesn't write a Masses section by default, so read_data left both types without mass -- "Not all per-type masses are set. Type 1 is missing." Add explicit mass commands after read_data, same values (H=1.008, O=15.999) the create_atoms version used to set directly. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
ATOMS's positions ((0,-2,0) etc.) were chosen for a box centered on the
origin (the old create_atoms version used region ... -10 10 -10 10 -10
10). Setting cell=[20,20,20] implies a [0,20) box instead, so that
atom landed outside it and got wrapped onto a periodic boundary by
ASE's writer -- creating a spurious near-boundary ghost neighbor LAMMPS
then picked up, producing one extra id and an out-of-bounds index when
reindexing the compute output by id ("index 3 is out of bounds for
axis 0 with size 3").
Translate all positions by +10 in each dimension before writing the
data file -- a rigid shift doesn't change the descriptor (translation
invariant), just keeps every atom comfortably inside the box.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Kp2WNbCTxZAGeEXSEpHqf6
…ouple from skmd
compute_symmetrix_mace_atom.cpp (CPU) zeroes descriptor output for atoms
outside the compute's group, per LAMMPS convention. The Kokkos port never
applied this check and wrote real values for every atom in the neighbor
list regardless of group -- silently harmless only because every current
caller invokes it on group "all". Fixed by syncing MASK_MASK and gating
both dispatch paths on mask & groupbit, matching the CPU compute exactly.
Added test_descriptor_group_subset (CPU + Kokkos) as a regression test.
Also removed stray comments referencing fix_skmd/skmd's internal naming
from compute_symmetrix_mace_atom_kokkos.cpp, compute_symmetrix_maced_atom_kokkos.{cpp,h},
and the test file -- symmetrix should not reference its downstream consumers.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…ests
compute_symmetrix_maced_atom_kokkos.{cpp,h} previously hard-errored on the
optional `vjp <compute-id>` args, only supporting the full-Jacobian mode.
Added run_vjp(): seeds H2_adj with v2 (broadcast to every node) and
propagates it back through the second layer into H1_adj, then adds the
v1-weighted linear_up_l0_inv seed onto that same H1_adj before a single
combined first-layer backward pass -- mirrors compute_symmetrix_maced_atom.cpp's
VJP branch (accumulate both seeds, one backward pass) rather than
run_channel_loop's per-channel loop.
test_compute_symmetrix_mace.py: added test_vjp (CPU + Kokkos, checks the
3-column VJP output against finite differences of v . summed descriptor)
and test_descriptor_group_subset (regression test for the earlier
groupbit fix). build_lammps() gained group_subset_ids/vjp_seed params;
the seed vector is supplied via `compute reduce ave` over atom-style
constant variables, avoiding a new test-only C++ compute.
Also swapped the kokkos test id's skip condition from an nvidia-smi
hardware probe to _kokkos_gpu_available(), which checks the LAMMPS
build's own accelerator_config for a GPU-capable Kokkos backend
(cuda/hip/sycl) -- correctly skips on CI runners with a CPU-only/host
Kokkos build regardless of whether a GPU chip happens to be present.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…use_vjp is parsed size_peratom_cols was computed from use_vjp before the "Parse optional args" block that actually sets use_vjp (it defaults false as a class member), so compute symmetrix/maced/atom ... vjp <id> always allocated the full 3*(2*num_channels)-column array and only ever wrote into its first 3 columns -- caught by the new test_vjp[cpu] test, which failed with array.shape == (3, 576) instead of (3, 3). Moved the arg-parsing block above the sizing block; no other logic changed. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
|
@wcwitt PR is ready - the reason the tests are failing is because extract_mace_data.py now also needs to extract With the new version of extract_mace_data.py, all tests should pass smoothly (I hope!) |
|
Just to note what this PR does - To extract it into python (in the correct order):
where the global descriptor Provided the user can supply a LAMMPS compute id for Syntax is: Extract in Python with: Both jac and vjp are covered with finite difference tests against the separate python version - CPU and GPU. |
A useful feature for
symmetrixwould be the ability to use a LAMMPScomputeto directly extract the per-atom descriptors and descriptor gradients, following the precedent set by compute sna/atom .To be clear, by 'descriptor' what I mean is the invariant part of the final layer, the part which is typically used as the input to the MLP readout.
This PR will aim to
compute sna/atom).compute snad/atom).Before merging, would be good to:
pair_symmetrix/for ease of use withinstall.sh).compute symmetrix/mace/atomandcompute symmetrix/maced/atom)