C1 — Capture a launch once, replay it many times

Record a stream of kernel launches once, then replay the recorded graph as many times as you like – with no per-launch API overhead.

Time: ~3 min · Runs on: CUDA GPU only (graph capture has no pure-C++ analogue) · You need: a CUDA toolchain, the build command above

The kernel we want to run:

// A trivial elementwise kernel — the "work" we want to capture and replay.
__global__ void addOne(int* buf, int n)
{
    const int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < n)
        buf[tid] += 1;
}

Record the launch into a graph. Everything issued on the stream between begin() and end() is captured rather than executed eagerly, then handed to the Graph as a node:

    // A stream to record on, and a Graph that will own the captured work.
    eagle::cuda::Stream stream;
    eagle::cuda::Graph graph;
    graph.stream(stream.cuda());

    // Everything launched on the stream between begin() and end() is recorded
    // into a cudaGraph_t instead of executing eagerly.
    eagle::cuda::StreamCapturer capturer(stream.cuda());
    capturer.begin();
    addOne<<<1, N, 0, stream.cuda()>>>(dBuf, N);
    graph.addNode(capturer.end());

Instantiate the graph once and replay it – each launch() is a single cudaGraphLaunch:

    // Instantiate the graph once, then replay it REPLAYS times. Each launch()
    // is a single cudaGraphLaunch — no kernel-config marshalling per call.
    eagle::cuda::Launcher launcher = graph.launcher();
    for (int r = 0; r < REPLAYS; ++r)
        launcher.launch();
    launcher.synchronize();

What it printed

01_capture_replay: 10 replays -> buf[0]=10 (expected 10) : OK

What just happened

• StreamCapturer recorded addOne’s launch instead of running it: the cudaGraph_t it produced between begin()/end() became one node on graph.

• graph.launcher() instantiated that graph exactly once into a Launcher; every later launch() replays the same instantiated graph as a single driver call, not a fresh launch of addOne.

• 10 replays of a kernel that adds 1 left every element at 10 – the proof that a replay really re-issues the captured work and not a no-op.

Try this

Raise REPLAYS to 1000 and rebuild: the loop issuing launcher.launch() still costs one driver call per iteration, but nothing about the capture above changes – the graph itself does not grow.

Next

C2 — Run one loop on the CPU or the GPU, unchanged – the same replay idea, but for work that also has a pure-C++ / OpenMP side with no CUDA at all.