diff --git a/CHANGELOG.md b/CHANGELOG.md index 9173ce8..2078217 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,43 @@ Notable changes of grace-runtime, grace-runtime-cuda12 and grace-runtime-rocm6 (they share one version). Each release has a section "Changed defaults" when a default behaves differently. +## [0.1.1](https://github.com/ICAMS/grace-runtime/releases/tag/0.1.1) (2026-10-29) + +Bug fixes and documentation. + +### Fixed + +- NVIDIA and AMD GPUs: GRACE-1L with 33.5 million atoms (2^25) or more stopped the whole process instead of running. + It now runs up to the GPU memory (A100 80 GB: 41 million atoms) and beyond that raises the usual "does not fit the + GPU memory" error. +- The PyPI page of `grace-runtime` shows the README. +- macOS: `pip install "grace-runtime[torchsim]"` installed PyTorch from PyPI, whose OpenMP runtime clashes with the + one of grace-runtime (`OMP: Error #15`). On macOS the extra now installs nothing; install PyTorch from conda-forge + (see the installation guide). +- macOS: `grace-runtime info` says that the GPU modules are Linux-only, instead of listing where it looked for the + NVIDIA driver. + +### LAMMPS (`grace-runtime` branch of the LAMMPS fork) + +- CPU styles: `GRACE_BLAS=AUTO` uses a BLAS only if it is MKL, OpenBLAS, BLIS or Accelerate, otherwise the built-in + kernel. Other BLAS libraries (for example ATLAS) made GRACE-3L about 4x slower, and on HPE Cray systems the BLAS + linked by the compiler made the link fail. +- CMake warns when the CPU styles are built with GCC older than 12: with GCC 11 GRACE-3L runs about 37 % slower, + with GCC 9 GRACE-1L about 4.6x and GRACE-3L about 6.6x slower. + +### Documentation + +- LAMMPS guide: build the fork for CPU, NVIDIA and AMD GPUs (with an example for LUMI), get a model, and run a first + input that also prints the extrapolation grade. +- Windows through WSL2; GCC 12 or newer for source builds. +- Benchmarks: time per atom and the largest system per GPU (L40S, A100, MI250X), and time per atom on a CPU. + +### Changed defaults + +- LAMMPS `GRACE_BLAS=AUTO` (see above): with a BLAS other than MKL, OpenBLAS, BLIS or Accelerate the CPU styles use + the built-in kernel. `-D GRACE_BLAS=ON` uses any BLAS that provides `cblas_sgemm`. +- macOS: the `torchsim` extra installs no packages. + ## [0.1.0](https://github.com/ICAMS/grace-runtime/releases/tag/0.1.0) (2026-10-29) First release on PyPI: `grace-runtime` (CPU), `grace-runtime-cuda12` (NVIDIA) and `grace-runtime-rocm6` (AMD). It diff --git a/README.md b/README.md index b749ba6..12d744b 100644 --- a/README.md +++ b/README.md @@ -3,13 +3,13 @@ Fast, efficient execution of GRACE-1L, GRACE-2L and GRACE-3L interatomic potentials without TensorFlow, PyTorch or JAX. CPU execution is standalone; GPU execution uses Kokkos for NVIDIA CUDA and AMD HIP. Use ASE-compatible calculators, batch processing, an optional TorchSim adapter or the shared kernels in LAMMPS with MPI. -PACE (ACE) potentials fitted with pacemaker run from their `.yaml` ([PACE models](docs/pace.md)). +PACE (ACE) potentials fitted with pacemaker run from their `.yaml` ([PACE models](https://github.com/ICAMS/grace-runtime/blob/master/docs/pace.md)). -**[Read the manual](docs/index.md)** · [Installation](docs/installation.md) · [API reference](docs/api.md) +**[Read the manual](https://github.com/ICAMS/grace-runtime/blob/master/docs/index.md)** · [Installation](https://github.com/ICAMS/grace-runtime/blob/master/docs/installation.md) · [API reference](https://github.com/ICAMS/grace-runtime/blob/master/docs/api.md) ## Install -Install one package from PyPI ([source builds](docs/installation.md#cpu-from-source) for other targets): +Install one package from PyPI ([source builds](https://github.com/ICAMS/grace-runtime/blob/master/docs/installation.md#cpu-from-source) for other targets): | Backend | Install | | --- | --- | @@ -18,7 +18,7 @@ Install one package from PyPI ([source builds](docs/installation.md#cpu-from-sou | AMD | `pip install grace-runtime-rocm6` | GPU packages install the matching CPU package automatically. -See the manual for [GPU requirements](docs/backends.md) and [environment setup](docs/installation.md#python-environment). +See the manual for [GPU requirements](https://github.com/ICAMS/grace-runtime/blob/master/docs/backends.md) and [environment setup](https://github.com/ICAMS/grace-runtime/blob/master/docs/installation.md#python-environment). ## First calculation @@ -38,22 +38,22 @@ Named models download on first use; local `kokkos.npz` files work too. ## Guides -The [manual](docs/index.md#guides) covers [batching and MD](docs/usage.md), [relaxation](docs/relax.md), -[properties](docs/properties.md), [dispersion](docs/dispersion.md), [PACE models](docs/pace.md), [TorchSim](docs/torchsim.md) and -[development](docs/development.md). An interactive tour is in the [ASE notebook](python/examples/ase/grace_runtime_ase.ipynb). +The [manual](https://github.com/ICAMS/grace-runtime/blob/master/docs/index.md#guides) covers [batching and MD](https://github.com/ICAMS/grace-runtime/blob/master/docs/usage.md), [relaxation](https://github.com/ICAMS/grace-runtime/blob/master/docs/relax.md), +[properties](https://github.com/ICAMS/grace-runtime/blob/master/docs/properties.md), [dispersion](https://github.com/ICAMS/grace-runtime/blob/master/docs/dispersion.md), [PACE models](https://github.com/ICAMS/grace-runtime/blob/master/docs/pace.md), [TorchSim](https://github.com/ICAMS/grace-runtime/blob/master/docs/torchsim.md) and +[development](https://github.com/ICAMS/grace-runtime/blob/master/docs/development.md). An interactive tour is in the [ASE notebook](https://github.com/ICAMS/grace-runtime/blob/master/python/examples/ase/grace_runtime_ase.ipynb). ## LAMMPS The [LAMMPS fork](https://github.com/yury-lysogorskiy/lammps/tree/grace-runtime) supplies pair styles `grace/runtime` (CPU) and `grace/runtime/kk` (GPU), which take the GRACE-1L / 2L / 3L model from the `kokkos.npz` given to `pair_coeff`, and `pace/runtime(/kk)` for PACE models. Installing the Python package does not install LAMMPS: -[build and run it](docs/lammps.md). +[build and run it](https://github.com/ICAMS/grace-runtime/blob/master/docs/lammps.md). ## Licence -Academic Software Licence (ASL, [LICENSE.md](LICENSE.md)): academic non-commercial use. For commercial use contact +Academic Software Licence (ASL, [LICENSE.md](https://github.com/ICAMS/grace-runtime/blob/master/LICENSE.md)): academic non-commercial use. For commercial use contact ACEworks GmbH, contact@aceworks.works. Bundled third-party code keeps its own licence -([THIRD_PARTY_NOTICES.md](THIRD_PARTY_NOTICES.md)). Changes: [CHANGELOG.md](CHANGELOG.md). +([THIRD_PARTY_NOTICES.md](https://github.com/ICAMS/grace-runtime/blob/master/THIRD_PARTY_NOTICES.md)). Changes: [CHANGELOG.md](https://github.com/ICAMS/grace-runtime/blob/master/CHANGELOG.md). ## Reference diff --git a/docs/benchmarks.md b/docs/benchmarks.md new file mode 100644 index 0000000..2ac49e0 --- /dev/null +++ b/docs/benchmarks.md @@ -0,0 +1,22 @@ +# Benchmarks + +How fast grace-runtime evaluates the GRACE foundation models, and how many atoms fit on one GPU. + +![Time per atom and largest system on one GPU, for GRACE-1L, 2L and 3L on NVIDIA L40S, NVIDIA A100 and AMD MI250X](images/benchmarks-gpu-light.png#only-light) +![Time per atom and largest system on one GPU, for GRACE-1L, 2L and 3L on NVIDIA L40S, NVIDIA A100 and AMD MI250X](images/benchmarks-gpu-dark.png#only-dark) + +![Time per atom on one CPU with 1 and 16 cores, for GRACE-1L, 2L and 3L](images/benchmarks-cpu-light.png#only-light) +![Time per atom on one CPU with 1 and 16 cores, for GRACE-1L, 2L and 3L](images/benchmarks-cpu-dark.png#only-dark) + +## How we measured + +| | | +|---|---| +| Models | GRACE-1L-OMAT-large-base, GRACE-2L-OMAT-large-base, GRACE-3L-OMAT-large (cutoff 6 Å) | +| System | bcc crystal (a = 3.165 Å) with 25 elements (Al to Os) in random order, atoms displaced by about 0.02 Å; about 58 neighbours per atom | +| Calculation | `GRACERuntimeCalculator` with default settings; each call computes energy, forces and stress | +| Time per atom | median of 3 calls after the first; the lowest value over system sizes from about 1000 atoms up. It changes little above about 16000 atoms (GPU) or 2000 atoms (CPU) | +| Largest system | double the number of atoms until a structure does not fit the GPU memory, then bisect; the largest structure that runs, each size in a new process | +| GPUs | NVIDIA L40S 48 GB, NVIDIA A100 80 GB PCIe, AMD Instinct MI250X: one GCD (half of the card) with 64 GB | +| CPU | Intel Xeon Gold 6526Y (16 cores per socket); 1 thread, or 16 threads on one socket (`OMP_NUM_THREADS`) | +| Software | grace-runtime 0.1 wheels from PyPI (`grace-runtime`, `grace-runtime-cuda12`, `grace-runtime-rocm6`) | diff --git a/docs/images/benchmarks-cpu-dark.png b/docs/images/benchmarks-cpu-dark.png new file mode 100644 index 0000000..53206a2 Binary files /dev/null and b/docs/images/benchmarks-cpu-dark.png differ diff --git a/docs/images/benchmarks-cpu-light.png b/docs/images/benchmarks-cpu-light.png new file mode 100644 index 0000000..0f768a7 Binary files /dev/null and b/docs/images/benchmarks-cpu-light.png differ diff --git a/docs/images/benchmarks-gpu-dark.png b/docs/images/benchmarks-gpu-dark.png new file mode 100644 index 0000000..2c04a27 Binary files /dev/null and b/docs/images/benchmarks-gpu-dark.png differ diff --git a/docs/images/benchmarks-gpu-light.png b/docs/images/benchmarks-gpu-light.png new file mode 100644 index 0000000..668b058 Binary files /dev/null and b/docs/images/benchmarks-gpu-light.png differ diff --git a/docs/index.md b/docs/index.md index 9d1427e..edf4ffb 100644 --- a/docs/index.md +++ b/docs/index.md @@ -50,6 +50,7 @@ print([r["energy"] for r in results]) | Use TorchSim integrators and optimizers | [TorchSim](torchsim.md) | | Build LAMMPS and run GRACE or PACE in it | [LAMMPS](lammps.md) | | Manage memory, threads and fallbacks | [Performance and diagnostics](advanced.md) | +| See speed and the largest systems per GPU | [Benchmarks](benchmarks.md) | | Look up parameters, defaults, CLI and results | [API reference](api.md) | | Build and test | [Development](development.md) | diff --git a/docs/installation.md b/docs/installation.md index 54f3da0..94a48fd 100644 --- a/docs/installation.md +++ b/docs/installation.md @@ -79,8 +79,13 @@ micromamba install -c conda-forge llvm-openmp ``` The macOS wheel ships its OpenMP runtime and shares it with conda-forge packages in the same process (NumPy, PyTorch). -PyTorch from PyPI ships its own, and the two abort with `OMP: Error #15`: for TorchSim on macOS, use conda-forge -PyTorch, which shares the environment's OpenMP runtime. +PyTorch from PyPI ships its own, and the two abort with `OMP: Error #15`. So on macOS the `torchsim` extra installs +nothing; install PyTorch from conda-forge, which shares the environment's OpenMP runtime, and then TorchSim: + +```bash +micromamba install -c conda-forge pytorch +pip install grace-runtime torch-sim-atomistic +``` ### GPU source builds @@ -93,7 +98,7 @@ For editable installs and library tests, see [development](development.md). | Feature | Install | | --- | --- | | Phonons | `pip install "grace-runtime[phonons]"` | -| TorchSim adapter (Python >= 3.12; on a CUDA 12 driver torch from the CUDA 12 index, see [TorchSim](torchsim.md)) | `pip install "grace-runtime[torchsim]"` | +| TorchSim adapter (Python >= 3.12; on a CUDA 12 driver torch from the CUDA 12 index, see [TorchSim](torchsim.md); on macOS see [macOS](#macos)) | `pip install "grace-runtime[torchsim]"` | | Tests | `pip install "grace-runtime[test]"` | | cuBLAS for the NVIDIA package (GRACE-3L up to ~15 % faster, UQ slightly faster) | `pip install "grace-runtime-cuda12[cublas]"` | diff --git a/docs/lammps.md b/docs/lammps.md index f32be6f..e56ae46 100644 --- a/docs/lammps.md +++ b/docs/lammps.md @@ -30,7 +30,7 @@ cd lammps ## 2. Build The pair styles are part of the `ML-PACE` package. When CMake configures the build, it downloads grace-runtime -0.1.0 from GitHub, so the configure step needs network access. It also downloads the ML-PACE library. Pick one of the +0.1.1 from GitHub, so the configure step needs network access. It also downloads the ML-PACE library. Pick one of the three recipes; run each from the `lammps` folder. ### CPU @@ -45,6 +45,12 @@ cmake --build build-cpu -j 8 binary on the same kind of CPU. `BUILD_OMP=yes` enables the OpenMP threads that the styles use inside each MPI rank; the `OPENMP` package is not needed. +GRACE-3L uses a BLAS for its matrix products only if the BLAS that CMake finds is MKL, OpenBLAS, BLIS or Accelerate; +otherwise it uses a built-in kernel. Other BLAS libraries, such as ATLAS, are slower than the built-in kernel. The +configure step prints the choice in a line `CPU GRACE styles: ...`. To pick a BLAS, add for example +`-D BLA_VENDOR=Intel10_64lp_seq` (sequential MKL) or `-D BLA_VENDOR=OpenBLAS`. `-D GRACE_BLAS=ON` uses any BLAS +that provides `cblas_sgemm` and stops if there is none; `-D GRACE_BLAS=OFF` always uses the built-in kernel. + ### NVIDIA GPU (Kokkos CUDA) You need the CUDA toolkit 12.2 or newer, 12.8 for Blackwell GPUs (`nvcc` on `PATH`). @@ -89,6 +95,33 @@ On HPE Cray systems (`PrgEnv-amd` with `cray-mpich`), also add `-D GRACE_BLAS=OF a BLAS at configure time that `hipcc` does not link, and the link then fails with `undefined symbol: cblas_sgemm`. The setting affects only the CPU styles, which then use their built-in matrix product. +Example for [LUMI](https://docs.lumi-supercomputer.eu/) (MI250X, `PrgEnv-amd`). Configure on a login node, because +the configure step needs network access; build in a batch job and run on a GPU node. On the login node, ROCm's CMake +files warn that `AMDGPU_TARGETS` is not set and no GPU was detected; the warning is harmless, the binary is built +for gfx90a: + +```bash +module --force purge +module load LUMI/25.03 partition/G +module load rocm/6.3.4 buildtools/25.03 PrgEnv-amd cray-mpich + +cmake -S cmake -B build-hip -D CMAKE_BUILD_TYPE=Release -D CMAKE_CXX_COMPILER=hipcc \ + -D PKG_KOKKOS=yes -D Kokkos_ENABLE_HIP=yes -D Kokkos_ARCH_AMD_GFX90A=yes \ + -D GPU_TARGETS=gfx90a -D CMAKE_CXX_FLAGS="-munsafe-fp-atomics" -D BUILD_MPI=yes -D PKG_ML-PACE=yes \ + -D GRACE_BLAS=OFF +sbatch -A project_ -p small-g -N 1 -n 1 -c 8 --gpus=1 -t 00:30:00 --wrap "cmake --build build-hip -j 8" +``` + +The build takes about 12 minutes; its output goes to `slurm-.out`. + +One GCD (half of an MI250X) per run, with the input of [step 4](#4-run): + +```bash +srun -A project_ -p small-g -N 1 -n 1 -c 7 --gpus=1 -t 00:30:00 \ + build-hip/lmp -k on g 1 -sf kk -pk kokkos newton on neigh half -in in.lmp \ + -var model ~/.cache/grace/GRACE-2L-OMAT-medium-ft-E/kokkos.npz +``` + Other AMD GPUs ([GPU backends](backends.md#amd-rocm)): | GPU | Flags | @@ -98,10 +131,10 @@ Other AMD GPUs ([GPU backends](backends.md#amd-rocm)): ### Without network access -Download grace-runtime 0.1.0 on a machine with network access, copy it over and pass its path: +Download grace-runtime 0.1.1 on a machine with network access, copy it over and pass its path: ```bash -git clone --depth 1 -b 0.1.0 https://github.com/ICAMS/grace-runtime.git +git clone --depth 1 -b 0.1.1 https://github.com/ICAMS/grace-runtime.git cmake -S cmake -B build-cpu ... -D LOCAL_GRACE_RUNTIME=$PWD/grace-runtime ``` @@ -130,7 +163,8 @@ set). ## 4. Run Save this input as `in.lmp`. It is a bcc W/Mo alloy of 1024 atoms; `pair_coeff` maps LAMMPS atom types 1 and 2 to W -and Mo, and the model file is given on the command line as the variable `model`: +and Mo, and the model file is given on the command line as the variable `model`. The input also prints the largest +extrapolation grade γ of all atoms (the model's uncertainty) and writes γ per atom to `uq.dump`: ```text units metal @@ -146,12 +180,23 @@ mass 2 95.95 pair_style grace/runtime pair_coeff * * ${model} W Mo +fix gamma all pair 20 grace/runtime gamma 1 +compute max_gamma all reduce max f_gamma + velocity all create 300 4928459 fix 1 all nvt temp 300 300 0.1 +thermo_style custom step temp pe press c_max_gamma thermo 20 +dump uq all custom 100 uq.dump id type x y z f_gamma run 100 ``` +`fix pair` asks the pair style for γ every 20 steps, the same interval as `thermo`. γ needs a model with an +uncertainty head (GRACE-2L-OMAT-medium-ft-E has one; see [uncertainty quantification](usage.md#uncertainty-quantification)). +For a model without one, leave out the `fix pair`, `compute` and `dump` lines and `c_max_gamma`. The fields +`atomic_sigma` and `gmm_cluster` are available in the same way; see the +[`pair_style grace/runtime` documentation](https://github.com/yury-lysogorskiy/lammps/blob/grace-runtime/doc/src/pair_grace_runtime.rst). + CPU, 4 MPI ranks with 4 OpenMP threads each: ```bash @@ -178,10 +223,14 @@ The style prints which model kind it loaded. On the CPU (`grace/runtime`): ```text [grace/runtime] loading the GRACE-2L model .../GRACE-2L-OMAT-medium-ft-E/kokkos.npz [GRACE-2L] lmax 4 | L1 42 ch, A1/AA/AAA/AAAA 25/141/107/378 | I 4x32 | L2 32 / 64 ch, ... | spline N 1024 - Step Temp E_pair E_mol TotEng Press - 0 300 -12248.756 0 -12209.086 40140.971 + Step Temp PotEng Press c_max_gamma + 0 300 -12248.756 40140.972 0.42767153 + 20 184.01952 -12233.328 43814.062 0.58249744 ``` +γ is the distance of an atom's environment from the training data, divided by a threshold: below 1 the atom is +inside the region the model was trained on, above 1 it extrapolates. Here all atoms stay below 1. + On the GPU (`grace/runtime/kk`), also the element mapping and the chunk size chosen from the free GPU memory: ```text @@ -189,12 +238,11 @@ On the GPU (`grace/runtime/kk`), also the element mapping and the chunk size cho [grace/runtime/kk] LAMMPS type 1 -> W (model index 83) [grace/runtime/kk] LAMMPS type 2 -> Mo (model index 43) [GRACE-2L/KK] chunksize auto: 1126 atoms per chunk (...) - Step Temp E_pair E_mol TotEng Press - 0 300 -12248.756 0 -12209.086 40140.93 ``` -The energy at step 0 is the same on the CPU and the GPU. Both evaluate the network in fp32, in a different order, so -the last digits and later steps of the trajectory can differ slightly. +With `fix pair` in the input, LAMMPS also warns that it switches to legacy (host) communication and sorting for +Kokkos; this is expected and harmless. The thermo output follows, as on the CPU. The energy at step 0 is the same on the CPU and the GPU. Both evaluate the +network in fp32, in a different order, so the last digits and later steps of the trajectory can differ slightly. Wrong input stops LAMMPS with an error that names the problem: a PACE `.yaml` in `grace/runtime`, a GRACE model in `pace/runtime`, a keyword that does not fit the model, or units other than `metal`. diff --git a/docs/torchsim.md b/docs/torchsim.md index 8e7170d..7e56851 100644 --- a/docs/torchsim.md +++ b/docs/torchsim.md @@ -2,7 +2,8 @@ [Installation](installation.md#optional-dependencies) · [API reference](api.md#torchsim-adapter) -Install the `torchsim` extra with Python >= 3.12. This complete example uses the default device selection: +Install the `torchsim` extra with Python >= 3.12 (on macOS, PyTorch from conda-forge instead: see +[macOS](installation.md#macos)). This complete example uses the default device selection: ```python import numpy # load before torch diff --git a/libgrace/include/grace/kk/csr_images.h b/libgrace/include/grace/kk/csr_images.h index 9463be3..11ffda3 100644 --- a/libgrace/include/grace/kk/csr_images.h +++ b/libgrace/include/grace/kk/csr_images.h @@ -153,11 +153,11 @@ template struct CsrImagesKK { auto jj = nl.j; auto SS = nl.S; // the bonds to images (S != 0) counted first: the sort keys are one per such bond, not one per bond (a large cell - // has few: 1L at 33M atoms, 64 bonds per atom, needed 15.75 GiB for keys of all bonds) - int bad = 0; - int64_t nimg = 0; - Kokkos::parallel_reduce("csr:check_S", Kokkos::RangePolicy(0, nbd), - KOKKOS_LAMBDA(const int64_t b, int &v, int64_t &c) { + // has few: 1L at 33M atoms, 64 bonds per atom, needed 15.75 GiB for keys of all bonds). Ranges over bonds and + // image bonds are Range64: GRACE-1L at 2^25 atoms has 2^31 bonds. + int64_t bad = 0, nimg = 0; + Kokkos::parallel_reduce("csr:check_S", Range64(0, nbd), + KOKKOS_LAMBDA(const int64_t b, int64_t &v, int64_t &c) { for (int k = 0; k < 3; k++) if (SS(b, k) <= -512 || SS(b, k) >= 512) v++; if (SS(b, 0) != 0 || SS(b, 1) != 0 || SS(b, 2) != 0) c++; @@ -165,7 +165,7 @@ template struct CsrImagesKK { if (bad) throw grace::Error("CSR images: lattice translation beyond +-511 cells"); Kokkos::View keys("csr:keys", std::max(nimg, 1)); int64_t nk = 0; - Kokkos::parallel_scan("csr:keys", Kokkos::RangePolicy(0, nbd), + Kokkos::parallel_scan("csr:keys", Range64(0, nbd), KOKKOS_LAMBDA(const int64_t b, int64_t &acc, const bool fin) { if (SS(b, 0) != 0 || SS(b, 1) != 0 || SS(b, 2) != 0) { if (fin) keys(acc) = pack_key(jj(b), SS(b, 0), SS(b, 1), SS(b, 2)); @@ -175,13 +175,14 @@ template struct CsrImagesKK { if (nk > 1) Kokkos::sort(Kokkos::subview(keys, std::make_pair((int64_t) 0, nk))); Kokkos::View uniq("csr:uniq", std::max(nk, 1)); int64_t ng = 0; - Kokkos::parallel_scan("csr:unique", Kokkos::RangePolicy(0, nk), + Kokkos::parallel_scan("csr:unique", Range64(0, nk), KOKKOS_LAMBDA(const int64_t i, int64_t &acc, const bool fin) { if (i == 0 || keys(i) != keys(i - 1)) { if (fin) uniq(acc) = keys(i); acc++; } }, ng); + if (ng > INT32_MAX - n) throw grace::Error("CSR images: more than 2^31 - 1 atoms and images"); nall = n + (int) ng; x = decltype(x)("csr:x", std::max(nall, 1)); type = typename Traits::t_int_1d("csr:type", std::max(nall, 1)); diff --git a/libgrace/include/grace/kk/support.h b/libgrace/include/grace/kk/support.h index 5f985aa..fb7e11b 100644 --- a/libgrace/include/grace/kk/support.h +++ b/libgrace/include/grace/kk/support.h @@ -266,6 +266,10 @@ template KOKKOS_INLINE_FUNCTION BondSlot bond_slot(const Member &t return {team.team_rank() + ts * (lr % teams_per_slot), lr / teams_per_slot}; } +// A RangePolicy with 64-bit indices, for ranges that can pass 2^31 (bonds of a whole system). Kokkos >= 5.2 indexes a +// RangePolicy without IndexType with int on CUDA (before: unsigned int) and aborts the process on a larger bound. +template using Range64 = Kokkos::RangePolicy>; + // A View whose element access computes the offset in 64 bits. Kokkos 5 View::operator() computes the offset in the // common type of the indices (an optimisation of Kokkos 5), so v(i, j, k) with int indices wraps once the view holds // more than 2^31 elements: the GRACE-3L [nall][64][50] message rows from nall = 671089. The views with a diff --git a/mkdocs.yml b/mkdocs.yml index 7fc7d7f..a3150a7 100644 --- a/mkdocs.yml +++ b/mkdocs.yml @@ -50,6 +50,7 @@ nav: - LAMMPS: lammps.md - Reference: - Performance and diagnostics: advanced.md + - Benchmarks: benchmarks.md - API reference: api.md - Development: development.md diff --git a/pyproject.toml b/pyproject.toml index d4f236b..7711667 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -4,8 +4,9 @@ build-backend = "scikit_build_core.build" [project] name = "grace-runtime" -version = "0.1.0" +version = "0.1.1" description = "Fast CPU and GPU inference of GRACE and PACE interatomic potentials with ASE, batching and TorchSim" +readme = "README.md" requires-python = ">=3.10" license = { text = "Academic Software License (ASL); bundled third-party code: see THIRD_PARTY_NOTICES.md" } classifiers = ["License :: Other/Proprietary License"] @@ -19,7 +20,8 @@ Homepage = "https://github.com/ICAMS/grace-runtime" Issues = "https://github.com/ICAMS/grace-runtime/issues" [project.optional-dependencies] -torchsim = ["torch>=2; python_version >= '3.12'", "torch-sim-atomistic>=0.6; python_version >= '3.12'"] +torchsim = ["torch>=2; python_version >= '3.12' and sys_platform != 'darwin'", + "torch-sim-atomistic>=0.6; python_version >= '3.12' and sys_platform != 'darwin'"] phonons = ["phonopy>=2.20"] test = ["pytest", "pyyaml"] diff --git a/python-cuda/pyproject.toml b/python-cuda/pyproject.toml index 67409ec..c391723 100644 --- a/python-cuda/pyproject.toml +++ b/python-cuda/pyproject.toml @@ -4,13 +4,13 @@ build-backend = "scikit_build_core.build" [project] name = "grace-runtime-cuda12" -version = "0.1.0" # = grace-runtime (scripts/build_cuda_wheels.sh sets both from the repository pyproject) +version = "0.1.1" # = grace-runtime (scripts/build_cuda_wheels.sh sets both from the repository pyproject) description = "NVIDIA GPU backend (CUDA 12) of grace-runtime: fused GRACE-1L/2L/3L Kokkos kernels for V100, RTX 20xx/30xx/40xx/50xx, A100/A30, L40S, H100/H200, B200/B300" readme = "README.md" requires-python = ">=3.10" license = { text = "Academic Software License (ASL); bundled third-party code (Kokkos: Apache-2.0 WITH LLVM-exception): see THIRD_PARTY_NOTICES.md" } classifiers = ["License :: Other/Proprietary License"] -dependencies = ["grace-runtime==0.1.0"] +dependencies = ["grace-runtime==0.1.1"] [project.urls] Homepage = "https://github.com/ICAMS/grace-runtime" diff --git a/python-rocm6/pyproject.toml b/python-rocm6/pyproject.toml index 0140398..85e5125 100644 --- a/python-rocm6/pyproject.toml +++ b/python-rocm6/pyproject.toml @@ -4,13 +4,13 @@ build-backend = "scikit_build_core.build" [project] name = "grace-runtime-rocm6" -version = "0.1.0" # = grace-runtime (scripts/build_rocm_wheels.sh sets both from the repository pyproject) +version = "0.1.1" # = grace-runtime (scripts/build_rocm_wheels.sh sets both from the repository pyproject) description = "AMD GPU backend (ROCm 6) of grace-runtime: fused GRACE-1L/2L/3L Kokkos kernels for MI210/MI250/MI250X and MI300A/MI300X" readme = "README.md" requires-python = ">=3.10" license = { text = "Academic Software License (ASL); bundled third-party code (Kokkos: Apache-2.0 WITH LLVM-exception): see THIRD_PARTY_NOTICES.md" } classifiers = ["License :: Other/Proprietary License"] -dependencies = ["grace-runtime==0.1.0"] +dependencies = ["grace-runtime==0.1.1"] [project.urls] Homepage = "https://github.com/ICAMS/grace-runtime" diff --git a/python/grace_runtime/cli.py b/python/grace_runtime/cli.py index 3e1a37b..199d727 100644 --- a/python/grace_runtime/cli.py +++ b/python/grace_runtime/cli.py @@ -98,6 +98,9 @@ def info(): "HSA_OVERRIDE_GFX_VERSION") if k in os.environ} print("settings:", ", ".join(f"{k}={v}" for k, v in flags.items()) or "defaults") + if sys.platform != "linux": + print("GPU: none (GPU modules are Linux-only); device='auto' runs on the CPU") + return 0 try: _cuda.compute_capability(0) except (_cuda.NoGPU, ValueError) as e: diff --git a/python/tests/test_rocm_plugin.py b/python/tests/test_rocm_plugin.py index 2428e7c..6594ef9 100644 --- a/python/tests/test_rocm_plugin.py +++ b/python/tests/test_rocm_plugin.py @@ -110,6 +110,7 @@ def cdll(name, *a, **k): return real(name, *a, **k) no_env.setattr(ctypes, "CDLL", cdll) no_env.setattr(_rocm, "plugins", lambda: []) + no_env.setattr(sys, "platform", "linux") # info off Linux skips the GPU probes return no_env @@ -371,6 +372,16 @@ def test_info_amd_native(amd_node, capsys): assert "rocBLAS:" in out +def test_info_on_macos_does_not_probe_gpus(no_env, capsys): + def probe(*a, **k): + raise AssertionError("GPU probed on macOS") + no_env.setattr(sys, "platform", "darwin") + no_env.setattr(_cuda, "compute_capability", probe) + no_env.setattr(_rocm, "kfd_present", probe) + assert cli.main(["info"]) == 0 + assert "GPU: none (GPU modules are Linux-only); device='auto' runs on the CPU" in capsys.readouterr().out + + def test_info_zero_visible_cuda_devices(amd_node, capsys): """cuInit(0) succeeds with 0 visible CUDA devices (ValueError from compute_capability(0)): info falls back to AMD or CPU instead of crashing.""" diff --git a/scripts/check_cuda_wheel.py b/scripts/check_cuda_wheel.py index bf65232..5fa22bf 100644 --- a/scripts/check_cuda_wheel.py +++ b/scripts/check_cuda_wheel.py @@ -2,7 +2,8 @@ """check_cuda_wheel.py [archs]: a grace-runtime-cuda12 wheel holds one module per architecture (default sm70, sm80, sm90, sm100, sm120) for its CPython tag, the plug-in package with ABI and the grace_runtime.backends entry point; no CUDA driver or toolkit library except the redistributable cudart; every module exports only PyInit_* (Kokkos and libstdc++ hidden); -auditwheel accepts it as manylinux_2_17 x86_64. Prints the size. Exit status 1 on a failed check.""" +it is manylinux_2_17 x86_64 but for libcuda.so.1, which the driver provides (wheel_policy.py). Prints the size. Exit +status 1 on a failed check.""" import re import subprocess import sys @@ -10,6 +11,8 @@ import zipfile from pathlib import Path +import wheel_policy + wheel = Path(sys.argv[1]) archs = (sys.argv[2] if len(sys.argv) > 2 else "sm70,sm80,sm90,sm100,sm120").split(",") bad = [] @@ -37,17 +40,15 @@ if re.match(r"lib(cuda|cublas|cublasLt|nvrtc|cufft|cusparse)[.-]", base): bad.append(f"bundles {base} (driver / toolkit library)") with tempfile.TemporaryDirectory() as d: - for n in names: - if "/_kk_" in n and n.endswith(".so"): - z.extract(n, d) - out = subprocess.run(["nm", "-D", "--defined-only", str(Path(d) / n)], capture_output=True, - text=True).stdout + elfs = {n: Path(z.extract(n, d)) for n in names if re.search(r"\.so(\.\d+)*$", n)} + for n, p in elfs.items(): + if "/_kk_" in n: + out = subprocess.run(["nm", "-D", "--defined-only", str(p)], capture_output=True, text=True).stdout extra = [l.split()[-1] for l in out.splitlines() if l.split() and not l.split()[-1].startswith("PyInit_")] if extra: bad.append(f"{n} exports {len(extra)} symbols besides PyInit_*: {extra[:5]}") -show = subprocess.run([sys.executable, "-m", "auditwheel", "show", str(wheel)], capture_output=True, text=True) -if "manylinux_2_17_x86_64" not in show.stdout: - bad.append("auditwheel: not manylinux_2_17_x86_64 compatible:\n" + show.stdout[-800:] + show.stderr[-400:]) + needs = {n: wheel_policy.needed(p) for n, p in elfs.items()} + bad += wheel_policy.check(wheel, "manylinux_2_17_x86_64", needs, ["libcuda.so.1"]) print(f"{wheel.name}: {wheel.stat().st_size / 2**20:.1f} MB, bundled libs: " f"{', '.join(n.rsplit('/', 1)[-1] for n in libs) or 'none'}") for b in bad: diff --git a/scripts/check_rocm_wheel.py b/scripts/check_rocm_wheel.py index 30cd78f..e5c31fb 100644 --- a/scripts/check_rocm_wheel.py +++ b/scripts/check_rocm_wheel.py @@ -4,8 +4,9 @@ grace_runtime.backends entry point; it bundles no ROCm library (the HIP runtime comes from the system ROCm; rocBLAS, hipBLAS, rocSOLVER and their Tensile kernels never: rocBLAS is loaded at run time) and no other library than libgomp; the modules need libamdhip64.so.6 and no BLAS library (DT_NEEDED), carry no absolute RUNPATH and export only -PyInit_*; its glibc symbol versions (auditwheel) allow the platform tag (default manylinux_2_28_x86_64; '-': no platform -checks), and the file name carries that tag alone. Prints the size. Exit status 1 on a failed check.""" +PyInit_*; it is of the platform tag but for libamdhip64.so.6, which the system ROCm provides (wheel_policy.py; default +manylinux_2_28_x86_64; '-': no platform checks), and the file name carries that tag alone. Prints the size. Exit +status 1 on a failed check.""" import re import subprocess import sys @@ -13,6 +14,8 @@ import zipfile from pathlib import Path +import wheel_policy + wheel = Path(sys.argv[1]) archs = (sys.argv[2] if len(sys.argv) > 2 else "gfx90a,gfx942").split(",") plat = sys.argv[3] if len(sys.argv) > 3 else "manylinux_2_28_x86_64" @@ -53,12 +56,13 @@ def elf(path, what): if not re.match(r"libgomp[.-]", base): bad.append(f"bundles {base} (only libgomp may be: the ROCm libraries come from the system)") with tempfile.TemporaryDirectory() as d: + needs = {} for n in names: if not re.search(r"\.so(\.\d+)*$", n): continue p = Path(z.extract(n, d)) dyn = elf(p, "-d") - needed = re.findall(r"\(NEEDED\)\s+Shared library: \[([^]]+)\]", dyn) + needed = needs[n] = wheel_policy.needed(p) blas = [x for x in needed if re.match(r"lib(rocblas|hipblas|hipblaslt|rocsolver|cublas)", x)] if blas: bad.append(f"{n} needs {blas} (BLAS must be loaded at run time, not linked)") @@ -73,22 +77,10 @@ def elf(path, what): extra = [l.split()[-1] for l in out.splitlines() if l.split() and not l.split()[-1].startswith("PyInit_")] if extra: bad.append(f"{n} exports {len(extra)} symbols besides PyInit_*: {extra[:5]}") - -def glibc(tag): - """(major, minor) of a manylinux tag (manylinux2014 = 2.17), None for others""" - m = re.match(r"manylinux_(\d+)_(\d+)_", tag) - return (int(m.group(1)), int(m.group(2))) if m else (2, 17) if tag.startswith("manylinux2014") else None - - -if plat.startswith("manylinux"): # '-': no platform checks (a wheel for this machine) - show = subprocess.run([sys.executable, "-m", "auditwheel", "show", str(wheel)], capture_output=True, text=True) - # the excluded HIP runtime keeps the verdict at linux_x86_64; the glibc symbol versions give the constraint - m = re.search(r'constrains the platform tag to\s*"([^"]+)"', show.stdout) or \ - re.search(r'consistent with the following platform tag:\s*"([^"]+)"', show.stdout) - if not m or not glibc(m.group(1)) or glibc(m.group(1)) > glibc(plat): - bad.append(f"auditwheel: not {plat} compatible:\n" + show.stdout[-800:] + show.stderr[-400:]) - if not wheel.name.endswith(f"-{plat}.whl"): - bad.append(f"platform tag of {wheel.name} is not {plat} alone") + if plat.startswith("manylinux"): # '-': no platform checks (a wheel for this machine) + bad += wheel_policy.check(wheel, plat, needs, ["libamdhip64.so.6"]) + if not wheel.name.endswith(f"-{plat}.whl"): + bad.append(f"platform tag of {wheel.name} is not {plat} alone") print(f"{wheel.name}: {wheel.stat().st_size / 2**20:.1f} MB, bundled libs: " f"{', '.join(n.rsplit('/', 1)[-1] for n in libs) or 'none'}") for b in bad: diff --git a/scripts/wheel_policy.py b/scripts/wheel_policy.py new file mode 100644 index 0000000..896e8ee --- /dev/null +++ b/scripts/wheel_policy.py @@ -0,0 +1,64 @@ +# SPDX-License-Identifier: LicenseRef-ASL (Academic Software Licence, see LICENSE.md) +"""The manylinux check of the GPU wheels (check_cuda_wheel.py, check_rocm_wheel.py). auditwheel calls such a wheel +linux_x86_64, because its modules need the GPU runtime the build leaves out on purpose (libcuda.so.1, +libamdhip64.so.6). So a wheel passes when auditwheel's tag fits the platform tag, or when the tag is linux_x86_64 for +missing libraries alone (its glibc symbol versions fit, no blacklisted symbol, no other policy failure) and every +library its ELF files need is in the policy's whitelist, vendored in the wheel, or one of the left-out GPU runtime +libraries.""" +import importlib.resources +import json +import re +import subprocess +import sys + +LOADER = "ld-linux-x86-64.so.2" # the dynamic loader: needed by some objects, never in a whitelist + + +def glibc(tag): + """(major, minor) of a manylinux tag (manylinux2014 = 2.17), None for others""" + m = re.match(r"manylinux_(\d+)_(\d+)_", tag or "") + return (int(m.group(1)), int(m.group(2))) if m else (2, 17) if (tag or "").startswith("manylinux2014") else None + + +def needed(path): + """the DT_NEEDED libraries of an ELF file (CalledProcessError when readelf cannot read it)""" + out = subprocess.run(["readelf", "-W", "-d", str(path)], capture_output=True, text=True, check=True).stdout + return re.findall(r"\(NEEDED\)\s+Shared library: \[([^]]+)\]", out) + + +def check(wheel, plat, needs, left_out): + """failures (list of str) of the manylinux check of wheel for plat (manylinux___x86_64). needs: {name in the + wheel: its DT_NEEDED} of its ELF files (modules and vendored libraries); left_out: the GPU runtime libraries the + build excludes""" + g = glibc(plat) + if g is None: + return [f"unknown platform tag {plat}"] + canon = f"manylinux_{g[0]}_{g[1]}" + policies = json.loads(importlib.resources.files("auditwheel.policy").joinpath("manylinux-policy.json").read_text()) + policy = next((p for p in policies if canon == p["name"] or canon in p["aliases"]), None) + if policy is None: + return [f"auditwheel has no policy {canon}"] + bad = [] + allowed = set(policy["lib_whitelist"]) | {n.rsplit("/", 1)[-1] for n in needs} | set(left_out) | {LOADER} + for n, libs in needs.items(): + foreign = [x for x in libs if x not in allowed] + if foreign: + bad.append(f"{n} needs {foreign}: not in the {plat} policy, not vendored") + show = subprocess.run([sys.executable, "-m", "auditwheel", "show", "--json", str(wheel)], capture_output=True, + text=True) + try: + r = json.loads(show.stdout) + except ValueError: + return bad + [f"auditwheel show --json failed: {show.stdout[-400:]}{show.stderr[-400:]}"] + + def fits(tag): + return glibc(tag) is not None and glibc(tag) <= g + + if fits(r["overall_tag"]): + return bad + other = [k for k in ("pyfpe", "ucs2", "unsupported_isa") if r.get(k)] + blacklisted = r["policy_upgrades"].get(f"{canon}_x86_64", {}).get("blacklisted_symbols") + if r["overall_tag"] != "linux_x86_64" or not fits(r["sym_tag"]) or other or blacklisted: + bad.append(f"auditwheel: not {plat}: tag {r['overall_tag']}, symbol versions {r['sym_tag']}, " + f"failed {other}, blacklisted symbols {blacklisted}") + return bad diff --git a/tests/kk/test_kk_large.cpp b/tests/kk/test_kk_large.cpp index 225bdcb..d613611 100644 --- a/tests/kk/test_kk_large.cpp +++ b/tests/kk/test_kk_large.cpp @@ -53,6 +53,9 @@ void check_guards() const int c = grace::kk::auto_chunk(1000000, 1, 1, grow, [] { return 0.0; }, grace::Logger{}, "test"); std::printf("guards: refused a 2^32-element view; chunksize auto %d for 10000 values per atom\n", c); CHECK((double) c * 10000 <= (double) INT32_MAX); + // a range over 2^31 bonds (GRACE-1L at 2^25 atoms): a policy with an int index aborts on this bound when built + const grace::kk::Range64 bonds(0, int64_t(1) << 31); + CHECK(bonds.end() == int64_t(1) << 31); } template void check_copies(Eng &eng, const char *tag, const Cell &w, const Cell &big)