Examples

Examples#

A gallery of complete, self-contained EAGLE programs. Each lives under docs/examples/ and is compiled and run by the exampletest gate (docs/examples/run_examples.sh) in every mode it supports, so the sources here are guaranteed to build and produce the output the Tutorials show.

Example

Mode

Modules

What it shows

01_capture_replay.cu

CUDA

graph

Capture a launch, build a graph, instantiate a Launcher, replay it.

02_host_dispatch.{cu,cpp}

CUDA + C++

host

One task per index through the OpenMP + SIMD cpu::Host launcher.

03_reduction.{cu,cpp}

CUDA + C++

reduce

Sum an aether SoA array with the GPU or the OpenMP reduction.

Full sources#

01_capture_replay.cu
 1// Copyright 2026 Alessandro Masat
 2// SPDX-License-Identifier: Apache-2.0
 3
 4// 01_capture_replay.cu
 5//
 6// The flagship EAGLE pattern: record a stream of kernel launches once, assemble
 7// them into a CUDA graph, instantiate a Launcher, and replay the instantiated
 8// graph many times with no per-launch API overhead.
 9//
10// This is a CUDA-only example (graph capture has no pure-C++ analogue). Build:
11//   nvcc -std=c++20 -arch=sm_61 -I<prefix>/include -DEAGLE_BLOCKSIZE=256 \
12//        -DCUDA_API_PER_THREAD_DEFAULT_STREAM=1 -Xcompiler -fPIE \
13//        01_capture_replay.cu -o 01_capture_replay
14#include <cstdio>
15#include <vector>
16
17#include <eagle/eagle.h>
18
19// [cell:kernel]
20// A trivial elementwise kernel — the "work" we want to capture and replay.
21__global__ void addOne(int* buf, int n)
22{
23    const int tid = threadIdx.x + blockIdx.x * blockDim.x;
24    if (tid < n)
25        buf[tid] += 1;
26}
27// [cell:kernel:end]
28
29int main()
30{
31    constexpr int N       = 64;
32    constexpr int REPLAYS = 10;
33
34    int* dBuf = nullptr;
35    if (cudaMalloc(&dBuf, N * sizeof(int)) != cudaSuccess)
36        return 1;
37    cudaMemset(dBuf, 0, N * sizeof(int));
38
39    // [cell:capture]
40    // A stream to record on, and a Graph that will own the captured work.
41    eagle::cuda::Stream stream;
42    eagle::cuda::Graph graph;
43    graph.stream(stream.cuda());
44
45    // Everything launched on the stream between begin() and end() is recorded
46    // into a cudaGraph_t instead of executing eagerly.
47    eagle::cuda::StreamCapturer capturer(stream.cuda());
48    capturer.begin();
49    addOne<<<1, N, 0, stream.cuda()>>>(dBuf, N);
50    graph.addNode(capturer.end());
51    // [cell:capture:end]
52
53    // [cell:replay]
54    // Instantiate the graph once, then replay it REPLAYS times. Each launch()
55    // is a single cudaGraphLaunch — no kernel-config marshalling per call.
56    eagle::cuda::Launcher launcher = graph.launcher();
57    for (int r = 0; r < REPLAYS; ++r)
58        launcher.launch();
59    launcher.synchronize();
60    // [cell:replay:end]
61
62    std::vector<int> host(N, -1);
63    cudaMemcpy(host.data(), dBuf, N * sizeof(int), cudaMemcpyDeviceToHost);
64    cudaFree(dBuf);
65
66    bool ok = true;
67    for (int i = 0; i < N; ++i)
68        ok &= (host[i] == REPLAYS);
69
70    std::printf("01_capture_replay: %d replays -> buf[0]=%d (expected %d) : %s\n",
71        REPLAYS, host[0], REPLAYS, ok ? "OK" : "FAIL");
72    return ok ? 0 : 1;
73}
02_host_dispatch.cu (CUDA-mode twin)
 1// Copyright 2026 Alessandro Masat
 2// SPDX-License-Identifier: Apache-2.0
 3
 4// 02_host_dispatch.cu
 5//
 6// The pure-C++ path: dispatch one task per index through EAGLE's OpenMP + SIMD
 7// host launcher. `eagle::cpu::Host::launch` batches consecutive indices into
 8// aether SIMD packets — the same call compiles and runs identically in CUDA
 9// mode (this file) and in pure-C++ mode (the .cpp twin), which is EAGLE's
10// dual-mode promise.
11//
12// IMPORTANT: `launch` only WORK-SHARES across OpenMP threads when the caller
13// is already inside an existing `#pragma omp parallel` region; called bare,
14// as below, it runs serially (SIMD-vectorised, single-threaded). This
15// standalone example has no enclosing region, so it is a *correctness* demo,
16// not a *threading* demo — see `eagle::cpu::Host::launch`'s doc comment
17// (eagle/cpu/Host.h) for the full contract: wrapping several `launch` calls
18// in your own `#pragma omp parallel` region is what gets them threaded in
19// production code.
20#include <cstdio>
21#include <vector>
22
23#include <eagle/eagle.h>
24
25int main()
26{
27    constexpr eagle::idx_t N = 1024;
28    std::vector<double> x(N, 0.0);
29
30    // [cell:launch]
31    // One host task per index i; i.global() is this task's flat index. We write
32    // the sequence of odd numbers 1, 3, 5, ... whose prefix sums are N*N.
33    eagle::cpu::Host::launch(N, [&](const auto& i) {
34        const eagle::idx_t k = i.global();
35        x[k] = 2.0 * static_cast<double>(k) + 1.0;
36    });
37    // [cell:launch:end]
38
39    double sum = 0.0;
40    for (eagle::idx_t k = 0; k < N; ++k)
41        sum += x[k];
42
43    const double expected = static_cast<double>(N) * static_cast<double>(N);
44    const bool ok         = (sum == expected);
45    std::printf("02_host_dispatch: N=%ld sum=%.0f (expected %.0f) : %s\n",
46        static_cast<long>(N), sum, expected, ok ? "OK" : "FAIL");
47    return ok ? 0 : 1;
48}
03_reduction.cu / 03_reduction.cpp
 1// Copyright 2026 Alessandro Masat
 2// SPDX-License-Identifier: Apache-2.0
 3
 4// 03_reduction.cu  (CUDA mode)
 5//
 6// Sum an aether SoA array with EAGLE's multi-level GPU reduction. The pure-C++
 7// twin (03_reduction.cpp) computes the same result with the cache-padded
 8// OpenMP reduction — same array, same operator, one dual-mode API.
 9#include <cstdio>
10
11#include <eagle/eagle.h>
12
13int main()
14{
15    using T   = double;
16    using Sum = aether::SumOp<T>;
17    constexpr eagle::idx_t N = 1000;
18
19    // Fill a device-backed array with 1, 2, ..., N on the host, then upload.
20    auto arr = eagle::makeArray<T>(N);
21    {
22        auto a = arr.hostView();
23        for (eagle::idx_t i = 0; i < N; ++i)
24            a(i) = static_cast<T>(i + 1);
25    }
26    arr.upload();
27
28    cudaStream_t stream;
29    cudaStreamCreate(&stream);
30
31    // [cell:reduce]
32    // Blocking multi-level reduction on the GPU: returns the scalar sum.
33    const T total = eagle::cuda::Reduction::reduceBlocking<T, Sum>(arr, T(0), stream);
34    // [cell:reduce:end]
35
36    cudaStreamDestroy(stream);
37
38    const T expected = static_cast<T>(N) * static_cast<T>(N + 1) / T(2);
39    const bool ok    = (total == expected);
40    std::printf("03_reduction (CUDA): sum(1..%ld)=%.0f (expected %.0f) : %s\n",
41        static_cast<long>(N), total, expected, ok ? "OK" : "FAIL");
42    return ok ? 0 : 1;
43}

 1// Copyright 2026 Alessandro Masat
 2// SPDX-License-Identifier: Apache-2.0
 3
 4// 03_reduction.cpp  (pure-C++ / OpenMP mode)
 5//
 6// The C++-mode twin of 03_reduction.cu: the same aether SoA array reduced with
 7// EAGLE's cache-padded OpenMP reduction. Built when eagle is configured with
 8// EAGLE_CPP_MODE (which defines EAGLE_CPU_ONLY and removes the CUDA path).
 9#include <cstdio>
10
11#include <eagle/eagle.h>
12
13int main()
14{
15    using T   = double;
16    using Sum = aether::SumOp<T>;
17    constexpr eagle::idx_t N = 1000;
18
19    auto arr = eagle::makeArray<T>(N);
20    {
21        auto a = arr.hostView();
22        for (eagle::idx_t i = 0; i < N; ++i)
23            a(i) = static_cast<T>(i + 1);
24    }
25
26    // [cell:reduce]
27    // Cache-padded OpenMP reduction over the host-resident array.
28    const T total = eagle::cpu::Reduction<T, Sum>::reduce(arr.hostView().as_const(), T(0));
29    // [cell:reduce:end]
30
31    const T expected = static_cast<T>(N) * static_cast<T>(N + 1) / T(2);
32    const bool ok    = (total == expected);
33    std::printf("03_reduction (C++): sum(1..%ld)=%.0f (expected %.0f) : %s\n",
34        static_cast<long>(N), total, expected, ok ? "OK" : "FAIL");
35    return ok ? 0 : 1;
36}

Running the gate#

# from the repo root, with eagle + aether installed into $CONDA_PREFIX
cd docs && make exampletest        # or: CONDA_PREFIX=<env> ./examples/run_examples.sh

The gate compiles each *.cu with nvcc (CUDA mode) and each *.cpp with g++ (pure-C++ / OpenMP mode), runs the resulting programs, and diffs their stdout against the committed *.expected.txt files. Adding an example is just dropping a new NN_topic.cu (and, for a dual-mode example, its .cpp twin) plus the matching .expected.txt — the gate and this gallery pick it up by its numeric prefix.

C4 — Embed a compiled kernel in your own C++ program#

Load a compiled plugin manifest straight into a C++/CUDA host through EAGLE’s binary ABI — no aether headers, no Python, no code generator at build or run time.

Time: ~5 min · Runs on: CUDA GPU only · You need: a CUDA toolchain, CMake (its own standalone build, separate from C1–C3’s exampletest gate above)

The examples above show EAGLE’s own core facilities (graph, host, reduce). A separate, larger example shows the other side of EAGLE: embedding a plugin deployment in a C++/CUDA host through nothing but EAGLE’s binary ABI (plugin/*.h) — no aether headers, no Python, no code generator at build or run time. It lives under examples/cpp_embed_plugin/ (its own standalone CMake project, not part of the numbered exampletest gate above, because it stages a multi-artifact manifest + fixture PTX rather than compiling one NN_topic.cu) and is the stable C++ embedding entry point for deploying a plugin this way.

embedding_demo.cu loads a 3-plugin manifest through eagle::cuda::PluginRegistry and captures all three as nodes in one CUDA graph opening, proving two capabilities together:

  • a matrix (mat_in) input, bound through the device registry’s bind_matrix (plugin/plugin_registry/registry.h) — a batch out[i] = trace(M[i]) over 3x3 matrices;

  • a derivative (VJP) artifact carrying the optional derivative sidecar block (plugin/sidecar.h) — a primal e = p * exp(-0.5*|x|^2) and its custom VJP, checked against a host central-difference reference (a check independent of the analytic formula the VJP kernel itself evaluates), the same shape of check the code generator’s own test suite runs on the Python producer/consumer side.

See examples/cpp_embed_plugin/README.rst for the full walkthrough (why it is a separate directory from plugin/’s three Driver-API demos, build/run instructions, expected output) and Plugin schema (manifest and sidecar) for the mat_in role and the derivative block’s schema.

embedding_demo.cu
  1// Copyright 2026 Alessandro Masat
  2// SPDX-License-Identifier: Apache-2.0
  3
  4// The C++ embedding reference demo.
  5//
  6// One driver-loaded generated "pure" plugin SET (three artifacts, one manifest),
  7// launched through the SAME independent `eagle::cuda::PluginRegistry` the other
  8// three plugin/ demos use (plugin/plugin_host.cpp, plugin/graph_inject/inject_demo.cu,
  9// plugin/pure_inject/pure_inject_demo.cu -- see README.rst for how this relates to
 10// them and why it lives in a separate directory rather than a fourth copy under
 11// plugin/). It proves the two capabilities this example exists to close:
 12//
 13//   1. a matrix (`mat_in`) input, bound through the device registry's new
 14//      `bind_matrix` (plugin/plugin_registry/registry.h) -- fixtures/mattrace.cu,
 15//      `out[i] = trace(M[i])` for a batch of 3x3 matrices;
 16//   2. a derivative (VJP) artifact -- fixtures/toy_energy.cu (the primal,
 17//      `e = p * exp(-0.5*|x|^2)`) and fixtures/toy_energy_vjp.cu (its custom VJP,
 18//      carrying the optional `derivative` sidecar block, Phase A recompute-only),
 19//      checked against a host CENTRAL-DIFFERENCE reference of the primal (an
 20//      reference independent of the analytic formula the VJP kernel itself uses) --
 21//      the same shape of proof as the code generator's downstream-toy-extension deploy test
 22//      reimplemented by
 23//      hand here so this example never imports it.
 24//
 25// All three plugins are captured as nodes in ONE CUDA graph opening (Runtime-API
 26// capture, Driver-API launch -- the same idiom pure_inject_demo.cu uses) and
 27// replayed twice, proving the matrix + derivative artifacts both survive replay
 28// like any other pure kernel (each is idempotent -- a fresh recompute every
 29// launch, no RMW state -- so both replays must agree with the same golden).
 30//
 31// Usage: embedding_demo [artifact_dir] [N]
 32//   artifact_dir defaults to the fixtures this target compiles to PTX at build
 33//   time (baked in via EMBED_DEMO_ART_DIR); N defaults to 4096.
 34
 35#include <algorithm>
 36#include <cmath>
 37#include <cstdint>
 38#include <cstdio>
 39#include <cstdlib>
 40#include <random>
 41#include <string>
 42#include <vector>
 43
 44#include <cuda.h>           // Driver API -- loads + launches the JIT plugins
 45#include <cuda_runtime.h>  // Runtime API -- capture, graph
 46
 47#include "plugin/plugin_registry/registry.h"
 48
 49#ifndef EMBED_DEMO_ART_DIR
 50#define EMBED_DEMO_ART_DIR "art"   // fallback if built outside this target's CMakeLists.txt
 51#endif
 52
 53namespace {
 54
 55constexpr double TOL     = 1e-12;  // mattrace / primal: exact closed-form agreement
 56constexpr double FD_H    = 1e-6;   // central-difference step
 57constexpr double FD_RTOL = 1e-5;   // VJP-vs-FD relative tolerance
 58constexpr double FD_ATOL = 1e-8;
 59
 60void cu_check(CUresult r, const char* what) {
 61    if (r != CUDA_SUCCESS) {
 62        const char* m = nullptr; cuGetErrorString(r, &m);
 63        std::fprintf(stderr, "Driver error %d (%s) at %s\n", int(r), m ? m : "?", what);
 64        std::exit(2);
 65    }
 66}
 67void rt_check(cudaError_t e, const char* what) {
 68    if (e != cudaSuccess) {
 69        std::fprintf(stderr, "Runtime error %d (%s) at %s\n", int(e),
 70                     cudaGetErrorString(e), what);
 71        std::exit(2);
 72    }
 73}
 74#define CU_CHECK(c) cu_check((c), #c)
 75#define RT_CHECK(c) rt_check((c), #c)
 76
 77// The SAME closed form fixtures/toy_energy.cu computes on device -- an
 78// independent (different language, different hardware) host implementation.
 79double host_energy(double x0, double x1, double x2, double p) {
 80    const double r2 = x0 * x0 + x1 * x1 + x2 * x2;
 81    return p * std::exp(-0.5 * r2);
 82}
 83
 84}  // namespace
 85
 86int run(int argc, char** argv) {
 87    const std::string dir = (argc >= 2) ? (std::string(argv[1]) + "/")
 88                                         : (std::string(EMBED_DEMO_ART_DIR) + "/");
 89    const std::uint32_t N = (argc >= 3)
 90        ? std::uint32_t(std::strtoul(argv[2], nullptr, 10)) : 4096;
 91
 92    // --- init: Runtime creates the primary context; Driver API shares it -------
 93    RT_CHECK(cudaSetDevice(0));
 94    RT_CHECK(cudaFree(0));
 95    CU_CHECK(cuInit(0));
 96    CUcontext ctx = nullptr; CU_CHECK(cuCtxGetCurrent(&ctx));
 97    if (ctx == nullptr) {
 98        CUdevice d; CU_CHECK(cuDeviceGet(&d, 0));
 99        CU_CHECK(cuDevicePrimaryCtxRetain(&ctx, d));
100        CU_CHECK(cuCtxSetCurrent(ctx));
101    }
102
103    // --- load the 3-plugin pure set (mattrace, toy_energy, toy_energy_vjp) ------
104    eagle::cuda::PluginRegistry registry =
105        eagle::cuda::PluginRegistry::from_manifest(dir + "manifest.json");
106
107    // --- host inputs (deterministic RNG; no external data files) ---------------
108    std::mt19937_64 rng(20260719);
109    std::uniform_real_distribution<double> ux(-1.3, 1.3);
110    std::uniform_real_distribution<double> ueb(0.3, 2.0);
111    const double p = 1.7;
112
113    std::vector<double> M(9 * std::size_t(N));   // flat (9, N) SoA, mattrace input
114    for (int d = 0; d < 9; ++d)
115        for (std::uint32_t i = 0; i < N; ++i)
116            M[std::size_t(d) * N + i] = double(d) + 0.25 * double(i);
117
118    std::vector<double> x(3 * std::size_t(N));   // flat (3, N) SoA, shared primal input
119    for (int k = 0; k < 3; ++k)
120        for (std::uint32_t i = 0; i < N; ++i)
121            x[std::size_t(k) * N + i] = ux(rng);
122
123    std::vector<double> e_bar(N);
124    for (std::uint32_t i = 0; i < N; ++i) e_bar[i] = ueb(rng);
125
126    // --- device buffers (all PRE-ALLOCATED; the registry only references them) --
127    double *d_M, *d_out, *d_x, *d_e, *d_e_bar, *d_x_bar, *d_p_bar;
128    RT_CHECK(cudaMalloc(&d_M, M.size() * sizeof(double)));
129    RT_CHECK(cudaMalloc(&d_out, N * sizeof(double)));
130    RT_CHECK(cudaMalloc(&d_x, x.size() * sizeof(double)));
131    RT_CHECK(cudaMalloc(&d_e, N * sizeof(double)));
132    RT_CHECK(cudaMalloc(&d_e_bar, N * sizeof(double)));
133    RT_CHECK(cudaMalloc(&d_x_bar, x.size() * sizeof(double)));
134    RT_CHECK(cudaMalloc(&d_p_bar, N * sizeof(double)));
135    RT_CHECK(cudaMemcpy(d_M, M.data(), M.size() * sizeof(double), cudaMemcpyHostToDevice));
136    RT_CHECK(cudaMemcpy(d_x, x.data(), x.size() * sizeof(double), cudaMemcpyHostToDevice));
137    RT_CHECK(cudaMemcpy(d_e_bar, e_bar.data(), N * sizeof(double), cudaMemcpyHostToDevice));
138
139    // --- bind by name (gap 1: bind_matrix is the new device entry point) -------
140    registry.bind_matrix("M", d_M, N, 3, 3);
141    registry.bind_handle("out", d_out);
142    registry.bind_vector("x", d_x, N);
143    registry.bind_handle("e", d_e);
144    registry.bind_uniform("p", p);
145    registry.bind_handle("e_bar", d_e_bar);
146    registry.bind_vector("x_bar", d_x_bar, N);
147    registry.bind_handle("p_bar", d_p_bar);
148
149    const int block = 256;
150
151    // === capture all 3 pure nodes in one opening, replay twice ==================
152    cudaStream_t stream; RT_CHECK(cudaStreamCreate(&stream));
153    RT_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal));
154    const int injected = registry.inject(reinterpret_cast<CUstream>(stream), int(N), block);
155    cudaGraph_t graph; RT_CHECK(cudaStreamEndCapture(stream, &graph));
156
157    std::size_t n_nodes = 0;
158    RT_CHECK(cudaGraphGetNodes(graph, nullptr, &n_nodes));
159
160    cudaGraphExec_t exec; RT_CHECK(cudaGraphInstantiate(&exec, graph, 0));
161    RT_CHECK(cudaGraphLaunch(exec, stream));
162    RT_CHECK(cudaGraphLaunch(exec, stream));   // replay: every plugin here is idempotent
163    RT_CHECK(cudaStreamSynchronize(stream));
164
165    std::vector<double> out(N), e(N), x_bar(3 * std::size_t(N)), p_bar(N);
166    RT_CHECK(cudaMemcpy(out.data(), d_out, N * sizeof(double), cudaMemcpyDeviceToHost));
167    RT_CHECK(cudaMemcpy(e.data(), d_e, N * sizeof(double), cudaMemcpyDeviceToHost));
168    RT_CHECK(cudaMemcpy(x_bar.data(), d_x_bar, x_bar.size() * sizeof(double),
169                        cudaMemcpyDeviceToHost));
170    RT_CHECK(cudaMemcpy(p_bar.data(), d_p_bar, N * sizeof(double), cudaMemcpyDeviceToHost));
171
172    // --- gap 1 check: mattrace vs the direct host sum ---------------------------
173    double mat_max_rel = 0.0;
174    for (std::uint32_t i = 0; i < N; ++i) {
175        const double ref = M[0 * N + i] + M[4 * N + i] + M[8 * N + i];
176        const double rel = std::abs(out[i] - ref) / std::max(std::abs(ref), 1.0);
177        mat_max_rel = std::max(mat_max_rel, rel);
178    }
179
180    // --- gap 2 check, primal: e vs the closed-form host reference ---------------
181    double primal_max_rel = 0.0;
182    for (std::uint32_t i = 0; i < N; ++i) {
183        const double ref = host_energy(x[0 * N + i], x[1 * N + i], x[2 * N + i], p);
184        const double rel = std::abs(e[i] - ref) / std::max(std::abs(ref), 1e-12);
185        primal_max_rel = std::max(primal_max_rel, rel);
186    }
187
188    // --- gap 2 check, VJP: x_bar/p_bar vs an INDEPENDENT central-difference of the
189    // same host closed form (not the analytic formula the device kernel uses) ----
190    double vjp_x_max_rel = 0.0, vjp_p_max_rel = 0.0;
191    for (std::uint32_t i = 0; i < N; ++i) {
192        const double x0 = x[0 * N + i], x1 = x[1 * N + i], x2 = x[2 * N + i];
193        double want_x[3];
194        want_x[0] = e_bar[i] *
195            (host_energy(x0 + FD_H, x1, x2, p) - host_energy(x0 - FD_H, x1, x2, p)) / (2 * FD_H);
196        want_x[1] = e_bar[i] *
197            (host_energy(x0, x1 + FD_H, x2, p) - host_energy(x0, x1 - FD_H, x2, p)) / (2 * FD_H);
198        want_x[2] = e_bar[i] *
199            (host_energy(x0, x1, x2 + FD_H, p) - host_energy(x0, x1, x2 - FD_H, p)) / (2 * FD_H);
200        const double want_p = e_bar[i] *
201            (host_energy(x0, x1, x2, p + FD_H) - host_energy(x0, x1, x2, p - FD_H)) / (2 * FD_H);
202        for (int k = 0; k < 3; ++k) {
203            const double got = x_bar[std::size_t(k) * N + i];
204            const double rel = std::abs(got - want_x[k]) / std::max(std::abs(want_x[k]), FD_ATOL);
205            vjp_x_max_rel = std::max(vjp_x_max_rel, rel);
206        }
207        const double relp = std::abs(p_bar[i] - want_p) / std::max(std::abs(want_p), FD_ATOL);
208        vjp_p_max_rel = std::max(vjp_p_max_rel, relp);
209    }
210
211    std::printf("N_PLUGINS=%zu\n", registry.size());
212    std::printf("ACTIVE=%d\n", registry.active());
213    std::printf("INJECTED=%d\n", injected);
214    std::printf("TOTAL_NODES=%zu\n", n_nodes);
215    std::printf("MATTRACE_MAX_REL=%.3e\n", mat_max_rel);
216    std::printf("PRIMAL_MAX_REL=%.3e\n", primal_max_rel);
217    std::printf("VJP_X_MAX_REL_VS_FD=%.3e\n", vjp_x_max_rel);
218    std::printf("VJP_P_MAX_REL_VS_FD=%.3e\n", vjp_p_max_rel);
219
220    const bool ok = (injected == 3) && (n_nodes == 3)
221        && (mat_max_rel < TOL) && (primal_max_rel < TOL)
222        && (vjp_x_max_rel < FD_RTOL) && (vjp_p_max_rel < FD_RTOL);
223    std::printf("%s\n", ok ? "PASS" : "FAIL");
224    return ok ? 0 : 1;
225}
226
227int main(int argc, char** argv) {
228    try {
229        return run(argc, argv);
230    } catch (const std::exception& e) {
231        std::fprintf(stderr, "ERROR: %s\n", e.what());
232        return 2;
233    }
234}

What just happened#

  • eagle::cuda::PluginRegistry loaded three independently compiled artifacts from one manifest and captured all three as nodes in a single CUDA graph opening — the host program never recompiles or links against the kernels themselves.

  • The matrix input (mat_in) crossed the ABI boundary through the registry’s own bind_matrix, with no aether header on the C++ side.

  • The VJP artifact’s output matched a central-difference reference built from the host’s own closed form — a check independent of the analytic formula the deployed kernel evaluates, not a second copy of the same formula.

Try this#

Pass a different <artifact_dir> on the command line (the README shows the baked-in default): the same host program loads whatever manifest it finds there, with no rebuild.

Next#

Tutorials — the three smaller C++ patterns (capture/replay, host dispatch, reduction) this example builds on. deeper: Plugin schema (manifest and sidecar) — the full manifest/sidecar schema embedding_demo.cu reads.