eagle

Runs kernels on CPU and GPU, with precise control. Part of the RAPTOR family.

eagle: run a million per-sample kernels on the GPU or the CPU

Each sample stops when it is done. eagle records the loop on the GPU once, replays it, and skips the samples that already finished.

The RAPTOR family

eagle is the execution layer of the RAPTOR family: it captures and replays kernels as CUDA graphs to keep work efficient on the GPU, and connects them to the code you already have. See where RAPTOR fits for the full four-part story.

Linux x86_64 wheels (Python 3.9-3.14). On macOS or Windows, open the notebook on Colab or Kaggle (the Kaggle badge above).

pip install raptor-eagle raptor-hawk                           # CPU threads, no NVIDIA packages
pip install "raptor-eagle[cuda12]" "raptor-hawk[cuda12]"       # NVIDIA GPU ([cuda13] for a CUDA 13 driver)

import cupy as cp

import eagle
import hawk
from hawk import Mutable, Param, Scalar, Terminated


@hawk.kernel
def oscillator(omega: Scalar, t_end: Param, dt: Param, terminated: Terminated,
               x: Mutable[Scalar], v: Mutable[Scalar], t: Mutable[Scalar]):
    x0, v0 = x, v
    x = x0 + dt * v0
    v = v0 - dt * omega * omega * x0
    t += dt
    terminated = t >= t_end


n = 1_000_000
result = eagle.simulate(
    oscillator,
    omega=cp.linspace(1.0, 3.0, n), t_end=1.0, dt=1e-3,
    x=cp.ones(n), v=cp.zeros(n), t=cp.zeros(n),
    max_steps=10_000,
)
print(result.status, result.steps, result.x[:3])
# finished 1000 [0.54057281 0.54057112 0.54056944]
# (steps counts whole check intervals, so it varies with the launch strategy eagle picks; x does not)

The same call runs on CPU threads with NumPy arrays. No GPU? Change the first line of the example to:

import numpy as cp   # CPU route: the rest runs unchanged (steps then prints 1024, not 1000: it counts whole check intervals, which differ by route)

Errors: no CUDA GPU visible means there is no GPU (or CUDA_VISIBLE_DEVICES hides it); cudaErrorInsufficientDriver means the installed extra does not match the driver, so install raptor-eagle[cuda12] for a CUDA 12 driver or [cuda13] for CUDA 13.

With the [torch] extra, a hawk kernel works inside PyTorch autograd:

import torch
from eagle.frameworks import torch as eagle_torch

step = eagle_torch.function(decay, wrt=("k", "x"))   # decay: a hawk kernel
k = torch.full((1000,), 0.5, dtype=torch.float64, requires_grad=True)
x_next = step(k, torch.ones(1000, dtype=torch.float64), torch.zeros(1000, dtype=torch.bool))
x_next.sum().backward()      # fills k.grad

Quickstart · Install · Performance · C++ API

_{NVIDIA H100 80GB HBM3 (card of 2026-10-13). NVIDIA Warp's per-thread kernel is the fastest GPU arm in 2 of 12 RK4 cells; on the dense batch (every sample running all 1000 steps) at N = 1,000,000 it takes 1.71 ms and eagle.simulate 1.75 ms. See every number, every arm.}

Threading

Free-threaded CPython (3.13t/3.14t) is supported: the compiled modules declare GIL-free operation and the GIL stays disabled after import. Any number of threads may call module-level functions, build, compile, plan and cache concurrently. Distinct objects may be used from distinct threads without synchronisation. One stateful object (a stream, capture, launcher, graph, composer, plan, pipeline, active set, host kernel or arg block) shared by several threads is memory-safe — its calls serialise and a consumed object raises — but the ORDER of those calls is the caller's responsibility, exactly as for a NumPy array or a CuPy stream. CUDA adds two rules of its own: a stream capture is begun, filled and ended by one thread, and while any capture is open no thread may synchronise the whole device (stream-level synchronisation is fine). GPU routes are supported on free-threaded 3.14; on 3.13t the CPU route runs. CUDA graph capture is per thread: a nested capture on one thread raises, while two threads may capture concurrently.

Install

pip install raptor-eagle raptor-hawk                      # CPU threads: no NVIDIA packages, nothing else to install
pip install "raptor-eagle[cuda12]" "raptor-hawk[cuda12]"  # NVIDIA GPU; [cuda13] for CUDA 13; add [torch] for PyTorch interop

Extras, platforms, driver and Python versions: installation. Building from source: build from source. The C++ library is a separate CMake install; the Python package from source is under Python package below.

# Configure (CUDA/C++ mode, with tests)
cmake -DCMAKE_PREFIX_PATH=${PREFIX} -DEAGLE_BUILD_TESTS=ON -B build .

# Configure (pure C++ mode)
cmake -DCMAKE_PREFIX_PATH=${PREFIX} -DEAGLE_CPP_MODE=ON -DEAGLE_BUILD_TESTS=ON -B build .

# Build
cmake --build build

# Test — the same command CI runs (there is no ctest wiring)
tests/check_gate.sh cuda build/tests/eagle_tests     # or: cpp, for a CPP_MODE build

# Install
cmake --install build --prefix ${PREFIX}

cmake -DCMAKE_PREFIX_PATH=${PREFIX} -B build-hdr . && cmake --install build-hdr --prefix ${PREFIX}
pip install raptor-core         # or a clone of the raptor repository: pip install ../raptor
export CMAKE_PREFIX_PATH=${PREFIX}
# append ;-DCMAKE_CUDA_HOST_COMPILER=<g++-13> if your default compiler is newer than nvcc accepts
export SKBUILD_CMAKE_ARGS="-DCMAKE_CUDA_ARCHITECTURES=<your GPU's arch, e.g. 70>"
pip install -e ./python
pytest python/tests -m "not gpu and not interop_matrix"   # drop the -m filter on a GPU machine with cupy and torch

C++ API

#include <eagle/eagle.h>

__global__ void addOne(int* buf, int n) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < n) buf[tid] += 1;
}

// Record a launch once, replay the captured graph many times.
eagle::cuda::Stream stream;
eagle::cuda::Graph graph;
graph.stream(stream.cuda());

eagle::cuda::StreamCapturer capturer(stream.cuda());
capturer.begin();
addOne<<<1, 64, 0, stream.cuda()>>>(buf, 64);
graph.addNode(capturer.end());

eagle::cuda::Launcher launcher = graph.launcher();
launcher.launch();          // cudaGraphLaunch — no per-call kernel-config overhead
launcher.synchronize();

EAGLE — Extensible Adaptive Graph Launch Engine — is a header-only C++23/CUDA library providing a parallel-dispatch substrate for GPU batch workloads. It bundles the reusable launch machinery in one place:

• CUDA-graph capture & launch — record a stream of kernel launches once (StreamCapturer), assemble them into a dependency graph (Graph), and replay via an instantiated Launcher with microsecond-class per-launch re-tuning (setLogicalSize).
• Stream primitives — RAII Stream, Event, and host-callback wrappers.
• OpenMP + SIMD host dispatch — eagle::cpu::Host packet-batched launchers with masked / flag-aware variants for the pure-C++ path.
• Scan / scatter / compaction — filtering::Scanner, filtering::Slice, and the prefix-sum + scatter kernels behind stream compaction.
• Reduction — eagle::cuda::Reduction (multi-level GPU) and eagle::cpu::Reduction (cache-padded OpenMP).
• PyTorch bridge — kernels plug into loss.backward() and torch.autograd.forward_ad, so both reverse-mode and forward-mode derivatives run through the same dual-mode engine.

Like its sibling libraries it is dual-mode: the same API compiles as CUDA or as pure C++23 (OpenMP), selected by EAGLE_CPP_MODE. The C++ core depends only on aether; the Python package (raptor-eagle) also depends on raptor, whose schema constants it re-exports.

Performance

eagle's committed performance card compares many ways of integrating a batch of damped harmonic oscillators (fixed-step RK4, float64, each sample stopping at its own step count) from 1,000 to 1,000,000 samples: eagle.simulate (the arm above), the CUDA-graph device loop alone, the same loop with active-set compaction, the same loop with compaction plus an occasional reorder, two K-fused-steps variants, the automatic-policy and persistent-launch arms, a plain eager launch loop, a CuPy array implementation, PyTorch masked, CUDA-graph and torch.compile implementations, a JAX jit(vmap(while_loop)) implementation, an NVIDIA Warp per-thread kernel, and a CPU OpenMP arm (no device or transfers needed, and the correctness reference the rest are checked against). The tables below show the NVIDIA H100 card; the five-arm table is the headline cut (both distributions, N = 1,000,000), and the full performance page has every arm, every N.

Where the performance comes from explains why eagle is ahead where samples finish at different times and level where they do not.

The chart shows the H100; the performance page has a tab for every device and CPU measured.

NVIDIA H100 80GB HBM3, card of 2026-10-13

Measured on:

• NVIDIA H100 80GB HBM3
• NVIDIA L4
• NVIDIA A100-SXM4-40GB
• Tesla T4
• Quadro P2000
• NVIDIA RTX PRO 6000 Blackwell Server Edition
• AMD EPYC 9B45
• AMD EPYC 7B13
• Google Axion (Arm Neoverse-V2)
• Intel Xeon Platinum 8581C

On a CPU alone: AMD EPYC 9B45, 8 cores, card of 2026-10-13

On a CPU alone: AMD EPYC 7B13, 8 cores, card of 2026-10-13

On a CPU alone: Google Axion (Arm Neoverse-V2), 16 cores, card of 2026-10-13

On a CPU alone: Intel Xeon Platinum 8581C, 8 cores, card of 2026-10-13

Figures are written by tools/sync_readme_cards.py from the card JSON (the same numbers card_nvidia-h100-80gb-hbm3.md renders): wall time and rates rounded to 3 significant figures, percent-of-peak to 2; a dash means that arm has no FP64-peak baseline recorded (CPU). Re-run the script after a card is re-run or a new perf_card device is added; python/tests/test_readme_cards.py fails if the committed block has drifted from a fresh render.

Where each tool fits. On a dense batch, where every sample runs all 1000 steps, NVIDIA Warp's per-thread kernel is a few percent ahead of eagle.simulate from N = 100,000 up (245 µs vs 260 µs at N = 100,000, 6%;

than Warp's per-thread kernel at N = 1,000,000 (613 µs against Warp's 1.54 ms, 2.5×). JAX compiles cold faster than hawk + eagle (about 242 ms, against eagle.simulate's 875 ms; NVIDIA Warp's cold compile is 574 ms), and JAX's warm recompile (40 ms) is shorter than hawk + eagle's own (70 ms). These are complementary tools: Warp and hawk + eagle both skip nothing, and the difference is who writes the per-thread loop and who decides when to compact. The CPU OpenMP arm needs no GPU, which makes it the tool to reach for on a machine without one; where a GPU is present, the GPU arms are far ahead on this card (199 ms against eagle.simulate's 613 µs on spread, 170 ms against 1.75 ms on uniform, at N = 1,000,000).

Compaction's advantage tracks how early the batch thins, and the device matters. On this H100 the plain graph runs the whole batch so quickly that active-set bookkeeping has little to recover: the plain graph is faster than the compaction arm in every cell of the card, 2.13 ms against 16 ms on the spread distribution at N = 1,000,000 and 2.32 ms against

shows a card where compaction pays on a thinning batch: 586 ms → 374 ms at N = 1,000,000 on spread (1.5×). The eager launch loop takes 15× (spread) and 16× (uniform) the plain graph's time at this N (32.3 ms / 39.1 ms vs 2.13 ms / 2.32 ms), which is what the CUDA graph is for: it keeps the step loop on the device. CuPy masked is the array-style option (368 ms on spread, 74% of peak memory bandwidth on uniform, recorded in the card).

The reorder arm is compaction's index map plus an occasional physical reorder, opt-in through eagle.ActiveSet(mask, reorder=theta). On the H100 it is faster than compaction alone in one cell of the card, N = 1,000,000 on the spread distribution (16 ms → 10.7 ms, 1.4×), and takes longer in the others, for example 23.3 ms against

CPU, no GPU at all: the CPU block above shows each CPU card measured; on the AMD EPYC 9B45 the host termination loop on 8 threads takes 142 ms (spread) / 187 ms (uniform) at N = 1,000,000, against Numba's

N = 1,000,000: 43.6 ms for eagle.simulate on the H100, 3.76 s for the eagle CPU path on 8 EPYC threads, and

where RAPTOR fits for that comparison.

Hardware caveat: the tables and text above are measured on a NVIDIA H100 80GB HBM3, a data-centre card (132 SMs,

carries over unchanged to other GPUs; the absolute numbers, and the crossover points between arms, differ from device to device, and the performance page has a tab for each one measured.

Full write-up, method and the per-N tables: the performance page. Raw numbers: card_nvidia-h100-80gb-hbm3.json / card_nvidia-h100-80gb-hbm3.md.

hawk write it · eagle run it (this repo) · aether the numerics core · raptor the shared contracts — the RAPTOR family

Apache-2.0 (see LICENSE and NOTICE) · cite via "Cite this repository" (CITATION.cff; every tagged release is archived on Zenodo: doi:10.5281/zenodo.23250240) · built to make GPU computing accessible on modest hardware, for research and education. Collaboration is the point, and a citation is the currency — get in touch.