Quick start

Quick start#

This page walks through a minimal EAGLE program: capture a kernel launch into a CUDA graph, instantiate a launcher, and replay it. It mirrors the shipped eagle_demo.cu smoke test.

Note

New to EAGLE’s vocabulary? Foundations — the words this documentation uses introduces Stream, Graph, StreamCapturer, and Launcher one at a time before this page puts them together. This page’s code block is a working confirmation of that walkthrough, not a first introduction.

A minimal capture-and-launch (CUDA)#

The four objects below are introduced in full in Foundations — the words this documentation uses; in short, here’s what each does in the program that follows:

  • Stream — an RAII wrapper around a cudaStream_t, the ordered queue of GPU work a kernel launch goes onto.

  • StreamCapturer — turns launch recording on (begin()) and off (end()), handing back a cudaGraph_t.

  • Graph — collects captured recordings into a runnable dependency graph.

  • Launcher — the instantiated, replayable graph; launch() submits it again, synchronize() waits for it to finish.

#include <cstdio>
#include <vector>
#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;
}

int main()
{
    constexpr int N = 64;
    int* dBuf = nullptr;
    cudaMalloc(&dBuf, N * sizeof(int));
    cudaMemset(dBuf, 0, N * sizeof(int));

    // RAII stream + graph builder
    eagle::cuda::Stream stream;
    eagle::cuda::Graph  graph;
    graph.stream(stream.cuda());

    // Capture the kernel launch into the graph
    eagle::cuda::StreamCapturer capturer(stream.cuda());
    capturer.begin();
    addOne<<<1, N, 0, stream.cuda()>>>(dBuf, N);
    graph.addNode(capturer.end());

    // Instantiate and replay
    eagle::cuda::Launcher launcher = graph.launcher();
    launcher.launch();
    launcher.synchronize();

    std::vector<int> host(N, -1);
    cudaMemcpy(host.data(), dBuf, N * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(dBuf);
    std::printf("eagle_demo: %s\n", host[0] == 1 ? "OK" : "FAIL");
    return 0;
}

Building the example#

find_package(eagle CONFIG REQUIRED)
add_executable(demo demo.cu)
target_link_libraries(demo PRIVATE eagle::eagle)

The same source compiles unchanged in EAGLE_CPP_MODE when written against the eagle::cpu::Host dispatcher instead of the CUDA graph launchers — see Module: cpu.

Next steps#

See also

Installation — install eagle and its aether dependency.