API reference — cuda

API reference — cuda#

CUDA-graph capture and launch machinery, the GPU scan / reduction engines, plus the RAII stream primitives.

Namespace `eagle::cuda`#

Note

The native graph-node protocol is a C++20 concept, which the documentation toolchain cannot render; it is summarised here and documented inline in eagle/cuda/NativeNode.h. A type N satisfies eagle::cuda::NativeNode when it provides reserveScratch(ScratchArena&) and buildInto(Graph&, deps) -> idx_t. Graph::addNative(node, deps) (constrained on the concept, hence also not listed below) registers such a node; Graph::finalizeNatives() commits the shared scratch arena and builds every registered node in add-order.

namespace cuda#

Functions

inline std::vector<cudaGraphNode_t> captureSnapshotNodes(cudaStream_t stream)#

Snapshot the node handles of the graph currently being captured into, as seen from stream: the mechanism GraphPipeline.build() uses for per-member node attribution.

Queries cudaStreamGetCaptureInfo_v3 for the in-construction cudaGraph_t and lists its current nodes via cudaGraphGetNodes — both calls are legal WHILE a ThreadLocal-mode capture is active (proven on this driver by probing the call mid-capture: the observed node count strictly grows, one per launch, as more kernels are captured). Every stream touched by one active capture session — the origin AND any CaptureFork branch pulled in via its fork event — reports the SAME underlying cudaGraph_t (the mechanism CaptureConditional::begin() already relies on for a fork-branch origin), so calling this immediately before and after one pipeline member’s launches — on whichever stream that member captures on — and taking the SET DIFFERENCE of the two snapshots yields exactly the nodes that member contributed, with no cross-member ambiguity: capture itself is single-threaded host-side sequencing, so one member’s launches fully enter the graph before the next member’s capture begins even when the members are siblings under a CaptureFork.

Deliberately node-HANDLE-only: never calls cudaGraphKernelNodeGetParams (the templated-kernel trap KernelNodeRecord’s own comment documents, Graph.h) — attribution only needs to know WHICH nodes a member owns and, via isNodeToggleable, their TYPE; never their launch parameters.

STOP-THE-LINE hardening (post-incident, round 2): The naive count-then-fill idiom (query count with nodes=nullptr, allocate exactly that many, fill) trusts that the count is still exact by the second call. CUDA’s own docs describe only the OVER-capacity case (numNodes higher than actual: excess entries set to NULL, actual count returned) and are silent on UNDER-capacity. Measured directly on this driver (deliberately under-reporting capacity by 5 against a 300-node graph): the fill call returns cudaSuccess and reports back EXACTLY the requested capacity, not the true (larger) count — i.e. it does not overflow the buffer, but it also gives no positive signal that truncation happened. To close that gap without trusting either call in isolation, this re-queries the count (nullptr) immediately AFTER the fill and compares against what was actually allocated: if the graph grew during the fill, retry with the new, larger size. Bounded so a pathologically still-growing graph cannot loop forever (never observed; capture is single-threaded host-side sequencing per call in every consumer this facility has).

Parameters:

stream – [in] A stream that is part of an ACTIVE Global-mode capture (the pipeline’s main stream, or a CaptureFork branch).

Throws:

aether::Error – If the capture-info or node-listing call fails, or the node count fails to stabilize (see below).
std::runtime_error – If stream is not actively capturing.

Returns:

Node handles in cudaGraphGetNodes order (not necessarily capture order; stable within one capture session).

inline bool isNodeToggleable(cudaGraphNode_t node)#

Whether node's type is one setNodeEnabled may safely toggle.

cudaGraphNodeSetEnabled’s own documentation (CUDA 12.6 cuda_runtime_api.h) states that only kernel, memset and memcpy nodes are currently supported — this predicate names exactly that set, so a member whose captured node set includes anything else (most pertinently a CaptureConditional IF node, cudaGraphNodeTypeConditional — the conditional/skippable facility) is reported non-toggleable at GraphPipeline.build() time rather than failing later, deep inside a raw driver call at the first toggle attempt.

inline void setNodeEnabled(cudaGraphExec_t exec, cudaGraphNode_t node, bool enabled)#

Enable or disable one node of an INSTANTIATED exec graph (the member-enable facility’s “mode=enabled” node toggling).

A disabled node becomes a driver-level no-op on every subsequent cudaGraphLaunch of exec — NO recapture, NO change to graph structure. Legal BETWEEN replays; toggling mid-replay is undefined and ordering it correctly against a replay is the caller’s responsibility (e.g. via GraphPipeline.launch()’s existing stream-ordering fix).

node must be a handle from the SAME (non-executable) cudaGraph_t that exec was instantiated from — e.g. one returned by captureSnapshotNodes at build time, before instantiation. This holds for eagle::cuda::Graph’s captured-graph constructor: launcher() instantiates the adopted graph verbatim (no clone), so node handles captured mid-build remain valid identifiers into it.

Throws:: aether::Error – If the node’s type does not support enable/disable, the node is not part of this exec graph, or any other driver rejection (call isNodeToggleable ahead of time to check the type without risking this call).

inline int currentSMCount()#: Cached single-device SM count.

inline int currentFp64PerfRatio()#

Cached FP32:FP64 throughput ratio of the current device.

cudaDevAttrSingleToDoublePrecisionPerfRatio — e.g. 2 on data-center parts (GV100/A100/H100), 32 on FP64-throttled workstation parts (sm_61 Pascal, Turing), 64 on consumer Ampere/Ada. Consumers use it to select between native-FP64 and emulated (SoftDouble) kernel sets at graph construction time: emulation pays off when the ratio is large.

inline void computeBlocks(idx_t numStates, idx_t &nBlocks, idx_t &blockSize, idx_t idealBlockSize = 256, idx_t blocksPerSM = 4)#

Block-size policy for an N-sample kernel.

Computes (blockSize, nBlocks) over numStates work items given the kernel’s __launch_bounds__-derived idealBlockSize cap and a target blocksPerSM. Targets nSMs * blocksPerSM blocks, rounds up to WARP=32, clamps to [WARP, idealBlockSize].

The caller supplies blocksPerSM rather than calling cudaOccupancyMaxActiveBlocksPerMultiprocessor because that fails with cudaErrorInvalidResourceHandle when the consumer DSO links a statically-linked cudart distinct from the kernel stub’s.

template<typename LaunchTraits> inline void computeBlocks(idx_t numStates, idx_t &nBlocks, idx_t &blockSize)#

Traits-taking overload: block-size policy from a kernel’s compile-time launch traits — eagle::launch::Traits or any type exposing maxBlockSize/minBlocksPerSM constants (duck-typed so this header carries no dependency on eagle/launch/Traits.h).

Explicit opt-in by design: pre-existing call sites keep passing values, because auto-threading a trait’s minBlocksPerSM into a site that defaulted blocksPerSM=4 would change its runtime grid silently.

inline DeviceProps deviceProps(int device = 0)#: Query cudaGetDeviceProperties for device and derive the roofline fields through :func:eagle::deriveFields. The one live-hardware source of an :class:eagle::DeviceProps — eagle::cpu:: deviceProps() (eagle/cpu/DeviceProps.h) is the GPU-free twin.

Note

Available only when EAGLE_CPU_ONLY is not defined.

template<typename T> inline T shflDown(unsigned int mask, T val, unsigned int offset)#: __shfl_down_sync for any reduction operand type.

template<typename T> inline T shflUp(unsigned int mask, T val, unsigned int offset)#: __shfl_up_sync twin (inclusive-scan warp stage).

class CaptureConditional#

#include <CaptureConditional.h>

Weave a device-evaluated conditional node — an IF (skippable region) or a WHILE (device-side loop) — into an in-flight stream capture: the CaptureFork sibling for gated regions.

GraphPipeline-style capture records ordinary launches as a linear chain; there is no capture-time entry point for a cudaGraphAddNode conditional today. CaptureConditional supplies exactly that, using the capture-interop mechanism proven under cudaStreamCaptureModeThreadLocal (eagle’s real capture mode, StreamCapturer::begin): query the in-flight graph and dependency set via cudaStreamGetCaptureInfo_v3, create a conditional handle bound to it, add the setter kernel node (setCond for an IF — the same CountGuard evaluation ConditionalGroup uses — or the loop head for a WHILE) and the conditional node explicitly, splice them back into the capture stream’s dependency set via cudaStreamUpdateCaptureDependencies, then open a NESTED capture of the body on a separate, pre-created stream.

See also

ConditionalGroup — the explicit-builder sibling (IF only in v1: same predicate, same node shape, non-capture assembly).

See also

CaptureFork — the sibling this class’s lifecycle mirrors.

The two kinds

IF (the two-argument constructor): the body runs at most once per replay, iff the guard holds when the setter node runs.
WHILE (the four-argument constructor): the body runs up to loopCap times per replay, while the guard holds. The loop state is ONE caller-owned device unsigned int[2] cell [remaining, ran]: the head setter (the node BEFORE the WHILE node) resets it every replay (remaining = cap, ran = 0) and sets handle = guard && cap > 0; the tail setter, launched by end() on the body stream as the body’s LAST node, sets handle = guard && --remaining > 0 and ++ran. No memset node, no cudaGraphCondAssignDefault, no host trip. ran is the device-visible iteration index (0 during the first iteration), readable by body kernels through the same cell. An empty loop body is refused like an empty IF body — it would be an infinite loop.

Lifecycle mirrors exactly

Construction is the only phase that acquires resources (the body stream): resource creation is illegal mid-capture on the capturing thread, so it must happen BEFORE StreamCapturer::begin(). begin() performs the mid-capture weave and returns the body stream to launch the guarded region’s work on; end() closes the body capture. end() is idempotent and the destructor calls it automatically (never throws from the destructor — same rationale as CaptureFork: a failure here is already being reported by the capture that is about to fail).

Origin

origin may be the pipeline’s main capturing stream, a CaptureFork branch stream, OR another CaptureConditional’s body stream (bodyStream()): weaving on any actively capturing stream is the same mechanism (an eagle-owned kernel node followed by an explicitly-added conditional node spliced back into whichever stream is capturing). A skippable region inside a loop body is exactly an IF weave whose origin is the WHILE weave’s body stream.

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Functions

inline CaptureConditional(const cudaStream_t &origin, CountGuard guard)#

Create the body stream for a later IF weave.

Must be called BEFORE StreamCapturer::begin() — stream creation is not permitted while a capture is in flight.

Parameters:

origin – [in] Stream the skippable region would have been captured on directly (the pipeline’s main stream, a CaptureFork branch, or another weave’s body stream) — the stream begin() will weave the conditional into.
guard – [in] The declarative predicate; retained by value for the lifetime of this object.

inline CaptureConditional(const cudaStream_t &origin, CountGuard guard, unsigned int loopCap, unsigned int *loopCounter)#

Create the body stream for a later WHILE weave (device-side loop). Same pre-capture rule as the IF constructor.

Parameters:

loopCap – [in] Maximum iterations per replay, >= 1; baked into the head setter as a kernel argument (a new cap is a new build). 0 is refused: a loop that can never run is a caller error, not a valid degenerate loop.
loopCounter – [in] Device unsigned int[2] cell [remaining, ran] owned by the caller for the graph’s whole lifetime (never nullptr). eagle writes both words every replay.

CaptureConditional(const CaptureConditional&) = delete#: Not copyable — it owns the body stream.

CaptureConditional &operator=(const CaptureConditional&) = delete#: Not copyable — it owns the body stream.

CaptureConditional(CaptureConditional&&) = delete#: Not movable — its lifetime is the begin/end scope.

CaptureConditional &operator=(CaptureConditional&&) = delete#: Not movable — its lifetime is the begin/end scope.

inline bool isLoop() const#: True for a WHILE weave, false for an IF weave.

inline cudaStream_t bodyStream() const#: The pre-created body stream, valid from construction: the origin a NESTED weave (a skippable region inside this loop’s body) must be constructed against, BEFORE capture begins — begin() returns this same handle later, mid-capture, but a nested weave’s own body stream cannot be created by then.

inline cudaStream_t begin()#

Weave the setter kernel + conditional node at the origin’s current capture tip and open capture of the body on the pre-created body stream.

Call with a capture already in flight on origin (the stream passed to the constructor). Every CUDA API call here is a cold, once-per-build control-plane call, checked with EAGLE_CHECK_ALWAYS (never the debug-only EAGLE_CHECK) — same convention as StreamCapturer / CaptureFork.

Returns:: The body stream: launch the guarded region’s work on it (e.g. via cp.cuda.ExternalStream on the Python side).

inline void end()#

Close the body capture. Idempotent — calling it again (or without a preceding begin()) is a no-op, not an error.

For a WHILE weave this first checks the body is non-empty (counted MID-capture, before anything eagle adds itself) and then launches the loop tail setter on the body stream, so it is captured as the body’s last node — after every node the caller launched, including a nested weave’s conditional node — before the body capture ends.

Exits the process-wide capture-depth scope (STOP-THE-LINE fix) UNCONDITIONALLY before the error check — the nested capture session is over the moment cudaStreamEndCapture returns, whether or not it reports success.

inline ~CaptureConditional()#: Close the body capture automatically if the caller did not. Never throws — a destructor that escaped during unwinding would terminate, and a failure here is already being reported by the capture that is about to fail (identical rationale to CaptureFork::~CaptureFork).

class CapturedGraph#

#include <CapturedGraph.h>

Move-only RAII owner of a captured cudaGraph_t plus per-kernel idealBlockSize caps.

Mirrors the ownership shape of the rest of the eagle graph layer (Graph, Launcher, Graph::Storage): non-copyable; move-only; destructor releases the wrapped cudaGraph_t. The type itself enforces that constructing a CapturedGraph from a borrowed handle is a programming error — the handle handed in must always represent a transfer of ownership.

Pair with Graph::addNode(CapturedGraph) (the owned-input overload): the by-value parameter is move-constructed from the argument, cudaGraphAddChildGraphNode clones into the parent, and the parameter’s destructor at the end of the addNode call releases the source. The Graph::addNode(cudaGraph_t) overload — which does not take ownership — is the right path for borrowed handles, e.g. a raw cudaGraph_t exposed by a consumer-owned subgraph builder.

idealBlockSizes[k] is the cap for the k-th kernel in cudaGraphGetNodes order. 0 means grid-only mutation. Vector size must equal the kernel count, or be empty (= grid-only for every kernel).

Public Functions

CapturedGraph() = default#: Default constructor: empty wrapper (no handle, no caps).

inline explicit CapturedGraph(cudaGraph_t graph, std::vector<idx_t> idealBlockSizes = {}) noexcept#

Take ownership of an existing cudaGraph_t and attach a per-kernel idealBlockSize table.

Conventional uses: wrap the result of StreamCapturer::end(), cudaGraphCreate, or cudaGraphClone. Do NOT pass a borrowed handle — that violates the type’s ownership contract and will lead to a double-free when the source’s original owner also calls cudaGraphDestroy.

CapturedGraph(const CapturedGraph&) = delete#: Copy is forbidden — ownership is exclusive.

inline CapturedGraph(CapturedGraph &&other) noexcept#: Move constructor. Source becomes empty.

inline CapturedGraph &operator=(CapturedGraph &&other) noexcept#

Move assignment. Releases any handle currently owned before taking the moved-from one.

STOP-THE-LINE fix: the release is handed to eagle::util::CaptureGuardState::destroyOrDefer rather than called directly — see Stream::destroy_()’s comment for the full rationale (a cudaGraphDestroy firing mid-capture, e.g. from a Python-GC-triggered move/destroy of an unrelated CapturedGraph, invalidates whatever capture happens to be active elsewhere in the process at that moment). Timing is unchanged when no capture is active (runs immediately, same as before).

inline ~CapturedGraph()#: Destructor releases the owned cudaGraph_t if any. STOP-THE-LINE fix: see the move-assignment operator’s note above — same deferred-if-capturing release.

inline cudaGraph_t graph() const noexcept#: Read-only access to the wrapped cudaGraph_t. Borrowed view — caller must not destroy.

inline const std::vector<idx_t> &idealBlockSizes() const noexcept#: Read-only access to the per-kernel idealBlockSize caps.

inline cudaGraph_t release() noexcept#: Relinquish ownership of the wrapped handle and return it. After the call *this is empty (destructor will not destroy anything). Use only when ownership is genuinely transferred to another owner.

inline explicit operator bool() const noexcept#: Whether the wrapper currently owns a handle.

class CaptureFork#

#include <CaptureFork.h>

Fork an in-flight stream capture into mutually independent sibling branches, then join them back.

A single StreamCapturer session records everything launched on its origin stream as one linear chain of graph nodes: each node depends on the one before it, because that is what stream order means. Work that is genuinely independent is therefore recorded as if it were sequential, and the resulting graph forbids the scheduler from ever overlapping it.

CaptureFork restores the missing structure. Between fork() and join() each branch stream is its own capture-order chain, so the captured DAG contains sibling nodes with no edge between them:

Stream origin(true);
CaptureFork fork(origin.cuda(), 2);   // pre-capture: streams + events

StreamCapturer capturer(origin.cuda());
capturer.begin();
stepA<<<..., origin.cuda()>>>(...);   // origin chain

fork.fork();
stepB1<<<..., fork.branch(0)>>>(...); // sibling of stepB2
stepB2<<<..., fork.branch(1)>>>(...); // sibling of stepB1
fork.join();

stepC<<<..., origin.cuda()>>>(...);   // depends on BOTH branches
cudaGraph_t g = capturer.end();

Forking grants the scheduler PERMISSION to co-execute the branches. It is never an obligation. Running every branch one after another, in any order, is always a conforming schedule — a device with one free SM, a device already saturated by another branch, and a future driver that simply chooses not to overlap are all behaving correctly. Code must never depend on co-execution actually happening, and no caller, test, or benchmark may treat serialized execution as a defect.

THE CONTRACT — what a fork does and does not promise

Consequently the relative order among branches is unobservable. There is no ordering between sibling branches, so nothing may rely on one branch’s effects being visible to another, and nothing may rely on a particular interleaving. The only ordering guarantees are: everything before fork() happens-before every branch, and every branch happens-before everything after join().

The caller asserts mutual independence. Branches must not race: no branch may write memory another branch reads or writes. This class cannot check that and does not try. Violating it produces a graph whose results depend on the schedule — which, per the paragraph above, is free to change.

Any benefit is bounded by how much of the device a single branch already occupies. State that boundary parametrically — in terms of the device’s SM count, the per-branch occupancy, and the branch count (e.g. a branch that already fills every SM leaves nothing for a sibling; roughly SM_count / branch_count SMs are available per branch when all branches are resident) — never as a constant measured on one GPU. Numbers from one card are evidence about that card, not a property of this primitive.

The mechanism is deliberately private and swappable. Events are how the fork is expressed today; that is an implementation detail, not part of the contract. The contract above ships forever; the events may not.

Construction is the only phase that acquires resources: all branch streams and all events are created up front, before the capture begins, because resource creation is illegal mid-capture. No phase ever allocates device memory. fork() and join() are pure capture-time signalling.

Lifecycle

join() is idempotent — calling it again is a no-op, not an error — and the destructor joins automatically if the caller did not, so a scoped CaptureFork is correct by construction. An unjoined branch is not a benign leak: it makes cudaStreamEndCapture fail with cudaErrorStreamCaptureUnjoined and discard the whole graph.

See also

StreamCapturer — owns the capture session this class forks within.

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Functions

inline CaptureFork(const cudaStream_t &origin, const std::size_t &branchCount)#

Create the branch streams and events for a later fork.

Must be called before StreamCapturer::begin(): stream and event creation is not permitted while a capture is in flight. Branch streams are non-blocking, so they carry no implicit dependency on the legacy default stream; events disable timing, since they are used purely for ordering.

Parameters:

origin – [in] Stream being captured — the one that forks and that the branches rejoin.
branchCount – [in] Number of independent branches.

Throws:

aether::Error – If stream or event creation fails.

CaptureFork(const CaptureFork&) = delete#: Not copyable — it owns streams and events.

CaptureFork &operator=(const CaptureFork&) = delete#: Not copyable — it owns streams and events.

CaptureFork(CaptureFork&&) = delete#: Not movable — its lifetime is the fork/join scope.

CaptureFork &operator=(CaptureFork&&) = delete#: Not movable — its lifetime is the fork/join scope.

inline void fork()#

Open the fork: every branch becomes a sibling of every other.

Call at the fork point, with a capture in flight on the origin stream. Records the fork event on the origin and makes each branch wait on it, which both pulls the branches into the capture and makes every node recorded so far a predecessor of every branch. Calling twice is a no-op.

Throws:: aether::Error – If the record or a wait fails.

inline void join()#

Close the fork: the origin waits for every branch.

Must be called before StreamCapturer::end() — an unjoined branch makes the capture fail wholesale. Each branch records its join event and the origin waits on all of them, so every node launched after the join depends on all branches.

Idempotent: calling it twice, or without a preceding fork(), is a no-op rather than an error.

Throws:: aether::Error – If a record or wait fails.

inline const cudaStream_t &branch(const std::size_t &index) const#

Stream to launch branch index on.

Only meaningful between fork() and join(); outside that window it is an ordinary idle stream.

Parameters:: index – [in] Branch index in [0, size()).
Returns:: The branch’s raw stream handle, for use as a launch argument.

inline std::size_t size() const#: Number of branches.

inline bool forked() const#: True once fork() has run and join() has not.

inline const cudaStream_t &origin() const#: The origin stream this fork branches from and rejoins.

inline ~CaptureFork()#

Join automatically if the caller did not.

Makes scoped use correct by construction: destroying an open fork before StreamCapturer::end() closes it rather than poisoning the capture. Never throws — a destructor that escaped during unwinding would terminate, and a join failure here is already being reported by the capture that is about to fail.

class Event#

#include <Event.h>

Event manager for non-default streams.

Public Functions

inline Event(const unsigned int &flags = cudaEventDefault)#: Default constructor.

Event(Event &other) = delete#: Copy constructor.

inline Event(Event &&other)#: Move constructor.

Event &operator=(Event &other) = delete#: Copy constructor is forbidden.

inline Event &operator=(Event &&other) noexcept#: Move assignment operator.

inline void record(const Stream &stream)#: record the event on the current stream

inline void synchronize() const#: Synchronize the host wrt this event.

inline const cudaEvent_t &cuda() const#: Get the low-level stream.

inline void waitQuery() const#: Wait event query.

inline ~Event()#: Default constructor.

class HostCallback#

#include <HostCallback.h>

Host-callback creator.

Public Functions

inline explicit HostCallback(std::function<void()> lambda)#: Construct with the given lambda function.

inline void launch(const Stream &stream)#: Enqueue the callback in the given stream.

template<typename IdxT> struct IsMultipleDependency#: #include <traits.h>

Manage multiple vs single dependencies.

template<typename T> struct IsMultipleDependency<std::deque<T>>#

template<typename T> struct IsMultipleDependency<std::initializer_list<T>>#

template<typename T> struct IsMultipleDependency<std::list<T>>#

template<typename T> struct IsMultipleDependency<std::vector<T>>#

class Launcher#

#include <Launcher.h>

Instantiated CUDA graph executor.

Holds a cudaGraphExec_t (instantiated from a cudaGraph_t) and provides launch() / synchronize() / setLogicalSize() for repeated execution on the associated stream. Constructed by Graph::launcher() only.

Lifetime: shares ownership of the producing Graph::Storage via std::shared_ptr<Graph::Storage> storage_. The source graph and every captured-child original it owns stay alive for as long as any Launcher (or the producing Graph) holds a reference. That guarantee keeps the kernel-node handles and kernelParams pointers in kernelNodes_ valid for any subsequent setLogicalSize call.

Multi-launcher fan-out: any number of Launchers can share the same Storage. Each instantiates its own independent cudaGraphExec_t.

Copy construction and copy assignment are disabled; use move semantics.

See also

Graph — builds the graph and produces a Launcher.

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Functions

Launcher() = default#: Default constructor: empty Launcher (no source storage, no exec instance). Useful as a member that gets move-assigned later.

Launcher(Launcher &other) = delete#: Copy constructor is forbidden.

inline Launcher(Launcher &&other) noexcept#: Move constructor.

Self &operator=(Launcher &other) = delete#: Copy assignment is forbidden.

inline Self &operator=(Launcher &&other) noexcept#: Move assignment.

inline Self &stream(const cudaStream_t &stream)#: Update the stream.

inline const cudaStream_t &stream() const#: Expose the stream.

inline void launch()#: Launch the graph.

inline void synchronize()#: Synchronize.

inline void setLogicalSize(idx_t logicalSize)#

Patch each registered kernel node’s launch params for the new logicalSize.

Per-record: idealBlockSize == 0 keeps capture-time blockDim and shrinks only gridDim (required for fixed-layout reduction-tree kernels); > 0 runs computeBlocks to re-tune both. Microsecond-class — no re-capture.

Never calls cudaGraphKernelNodeGetParams at runtime: the exec graph round-trip rejects templated kernels with InvalidDeviceFunction. Snapshotted capturedParams is copied and mutated instead.

Lifetime: walks our own kernelNodes_ snapshot. Both the node handles and the kernelParams pointers reference internals of storage_->graph; safe because storage_ keeps Storage alive for our lifetime.

inline idx_t kernelNodeCount() const#: Number of registered kernel records (for tests / debug).

inline const std::vector<KernelNodeRecord> &kernelNodes() const#: Read-only accessor to the registered kernel records (tests / diagnostics).

inline ~Launcher()#

Destructor — destroys the exec instance.

storage_ (the shared Graph::Storage reference) is released after; if we hold the last reference the Storage destructor runs here and tears down the source CUDA graph.

inline cudaGraph_t graph() const#: Expose the source graph (lifetime tied to storage_).

inline cudaGraphExec_t execHandle() const#: Raw cudaGraphExec_t handle of the instantiated exec graph — escape hatch for facilities (member-enable toggling, eagle::cuda::setNodeEnabled in CaptureAttribution.h) that must call cudaGraphExec* APIs directly, mirroring Graph::nodeHandle()’s escape-hatch precedent. A bare accessor — no CUDA API call here, so this adds no new version floor to this shared header.

struct MatBinding#

class PluginRegistry#

Public Functions

inline int inject_partition(CUstream stream, const eagle::exec::Partition &part, int block = 256)#

The PARTITIONED injection opening: every enabled plugin, in manifest order, over ONE partition of the run.

inject(stream, n) is exactly this with Partition::whole(n), so the legacy opening and the partitioned one are ONE implementation — the packing rules, the gate order and the manifest order cannot drift between them. A v2 plugin receives the int64 triple after its role args; a v1 plugin receives nothing extra and is REFUSED a non-whole partition.

inline int inject_plugin_partition(CUstream stream, std::size_t index, const eagle::exec::Partition &part, int block = 256)#: Single-plugin twin of :func:inject_partition.

inline int abi_version(std::size_t index) const#: Which ABI generation the plugin at index speaks (1 or 2).

inline CUfunction function(std::size_t index) const#: The resolved CUfunction of the plugin at index — what an execution structure (eagle::exec::DeviceKernel) launches.

inline PackedArgs pack(std::size_t index, std::int64_t nSamples) const#

Pack the plugin at index's role args for a run spanning nSamples samples.

Every role is packed WHOLE: eagle does no pointer arithmetic on plugin buffers, and the body slices per-sample roles by its own triple. The count that reaches the nsamples role and every view mirror’s extent is the TRUE nSamples of the run, never a partition’s count — review Note: a wide/accum plane’s addressing bakes the full sample count, so handing it a partition’s length silently re-targets every column.

struct Plugin#

struct PackedArgs#

struct Reduction#

#include <Reduction.h>

GPU parallel reduction over aether scalar arrays.

All overloads of reduceBlocking synchronise the CUDA stream before returning the result on the host. For graph-based (non-blocking) variants, use the overloads that accept a cuda::Graph&.

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Static Functions

template<typename T> static inline aether::Array<T> makeBuffer(const idx_t &size)#

Generate a work buffer for use with reduceBlocking, of the appropriate size to perform a reduction over arr.

Note: the returned array is only allocated on the host. You will need to call oarr.upload before invoking reduceBlocking.

template<typename T> static inline aether::Array<T> makeBuffer(const aether::Array<T> &arr)#

Generate a work buffer for use with reduceBlocking, of the appropriate size to perform a reduction over arr.

Note: the returned array is only allocated on the host. You will need to call oarr.upload before invoking reduceBlocking.

template<typename T, typename OP> static inline T reduceBlocking(aether::Array<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)#

Perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.
init – Initial value for the accumulation.
stream – CUDA stream to use.

template<typename T, typename OP> static inline T reduceBlocking(aether::Array<T> &arr, const T &init = 0, const cudaStream_t &stream = 0)#

Perform a reduction over a GPU-allocated array, allocating a work buffer ad-hoc.

Note: causes a blocking synchronization of the given stream.

Note that this will allocate a new, temporary work buffer on host and device. If performing multiple reductions with same-sized arrays, it may be more efficient to generate the buffer once with makeBuffer and reuse it with the overloaded version of this function.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
init – Initial value for the accumulation.
stream – CUDA stream to use.

template<typename T, typename OP> static inline T reduceBlocking(const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)#

Perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.
init – Initial value for the accumulation.
stream – CUDA stream to use.

template<typename T, typename OP, bool DirectCopy = false> static inline void reduceBlocking(cuda::Graph &graph, T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const idx_t N, const T &init = 0, const cudaStream_t &stream = 0)#

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

result – Pointer to the reduction result
arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.
init – Initial value for the accumulation.
stream – CUDA stream to use.

template<typename T, typename OP, bool DirectCopy = false> static inline void reduceBlocking(cuda::Graph &graph, T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)#

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

result – Pointer to the reduction result
arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.
init – Initial value for the accumulation.
stream – CUDA stream to use.

template<typename T, typename OP> static inline cuda::Graph reduceBlocking(T *result, const CRefArrT<T> &arr, aether::Array<T> &oarr, const T &init = 0, const cudaStream_t &stream = 0)#

Build a graph to perform a reduction over a GPU-allocated array using a pre-allocated work buffer.

Note: causes a blocking synchronization of the given stream.

Template Parameters:

T – Data type of the elements
OP – Reduction operator to use. Must be callable as OP{}(a, b) (aether’s functor shape). Example: aether::SumOp<T>

Parameters:

result – Pointer to the reduction result
arr – Input scalar array. The reduction will be performed over the device copy, but it will not be modified. Must be already allocated on the device.
oarr – Output array. Used for intermediate copies. Must be allocated on host and device. Its values on both host and device will be overwritten. It is recommended to generate it via makeBuffer to ensure it has the appropriate size. Must be distinct from arr.
init – Initial value for the accumulation.
stream – CUDA stream to use.

struct Scan

#include <Scan.h>

Perform the Prefix sum (scan) algorithm.

Public Static Functions

template<typename T, typename OP, bool Inclusive, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline void enqueue(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0)

Enqueue the full scan on stream, without capturing it.

The launch sequence graph captures, issued straight onto stream: nothing is allocated, nothing is synchronised and nothing is read back to the host, so a caller that is ALREADY capturing stream (a graph-building consumer) records it into its own graph, and an uncaptured caller simply runs it asynchronously. blockSums must hold at least arr.samples() elements. The exclusive variant ends with a device-to-device copy; the inclusive one is kernels only.

template<typename T, typename OP, bool Inclusive, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline cudaGraph_t graph(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Capture the full scan algorithm (enqueue) into a graph.

template<typename T, typename OP, bool Inclusive, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline void launch(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Instantiate and launch.

template<typename T, typename OP, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline cudaGraph_t inclusiveGraph(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Run a full inclusive scan algorithm.

template<typename T, typename OP, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline cudaGraph_t exclusiveGraph(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Run a full exclusive scan algorithm.

template<typename T, typename OP, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline void inclusive(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Instantiate and launch an inclusive scan.

template<typename T, typename OP, unsigned int blockSize_ = filtering::detail::ScanCore::blockSize> static inline void exclusive(const CRefArrT<T> arr, GRefArrT<T> results, GRefArrT<T> blockSums, const T &init = 0, const cudaStream_t &stream = 0): Instantiate and launch an exclusive scan.

class Stream#

#include <Stream.h>

Stream manager for non-default streams.

Public Functions

inline explicit Stream(bool nonBlocking = false)#

Construct a stream.

Parameters:: nonBlocking – when true the stream is created with cudaStreamNonBlocking (no implicit dependency on the legacy default stream) — the correct choice for CUDA-graph capture that must not entangle the default stream other libraries (cupy, torch) launch on. Default false reproduces cudaStreamCreate (cudaStreamDefault) exactly — but note precisely what that buys: implicit ordering against the LEGACY default stream ONLY. Under -DCUDA_API_PER_THREAD_DEFAULT_STREAM=1 (or the --default-stream per-thread compile flag) stream-0 work runs on the per-thread default stream and has NO implicit ordering edge, in either direction, with a blocking eagle Stream — a default-stream launch can race work on this stream with every CUDA call returning cudaSuccess. Consumers mixing default-stream launches with eagle-Stream work under per-thread default streams must order explicitly (an Event or a sync). “Existing consumers are unchanged” therefore holds only for legacy-default-stream builds.

Stream(Stream &other) = delete#: Copy constructor.

inline Stream(Stream &&other)#: Move constructor.

Stream &operator=(Stream &other) = delete#: Copy constructor is forbidden.

inline Stream &operator=(Stream &&other) noexcept#: Move assignment operator.

inline void waitFor(const Event &event)#: Wait for the given event.

inline void synchronize() const#: Synchronize the host wrt this stream.

inline const cudaStream_t &cuda() const#: Get the low-level stream.

inline ~Stream()#: Default constructor.

class StreamCapturer#

#include <StreamCapturer.h>

RAII helper for recording kernel launches into a CUDA graph via stream capture.

Call begin() before launching kernels on the associated stream, and end() afterwards to obtain a cudaGraph_t that can be added to a Graph with Graph::addNode().

See also

Graph::addNode() — consumes the captured graph.

See also

eagle::util::CaptureGuardState — begin()/end() are a capture ENTRY/EXIT point (STOP-THE-LINE incident fix): captureScope_ brackets the process-wide capture-depth counter so a CUDA-touching teardown elsewhere (e.g. a GC’d Launcher) defers itself instead of invalidating this capture.

Note

Capturing the default stream (stream = 0, the constructor’s default) only works because eagle exports CUDA_API_PER_THREAD_DEFAULT_STREAM=1 as an INTERFACE compile definition on their CMake targets; a consumer compiling bare nvcc without that target’s flags hits cudaErrorStreamCaptureUnsupported (900) at begin().

Note

Available only when EAGLE_CPU_ONLY is not defined.

Public Functions

inline StreamCapturer(const cudaStream_t &stream = 0)#: Construct by assigning the stream.

inline StreamCapturer(const StreamCapturer &other)#: Copy constructor. NOTE: captureScope_ is deliberately NOT copied (default-constructed empty on the copy) — each StreamCapturer object tracks its OWN begin/end pairing, and copying/moving one mid-capture was already fragile before this change (two objects aliasing the same stream_’s capture session); this is not exercised anywhere in this codebase’s actual usage (construct, begin, end, done).

inline StreamCapturer(StreamCapturer &&other)#: Move constructor. See the copy constructor’s note — captureScope_ is not transferred.

inline Self &operator=(const cudaStream_t &str)#: Assignment from stream.

inline Self &operator=(const StreamCapturer &other)#: Copy assignment.

inline Self &operator=(StreamCapturer &&other)#: Move assignment.

inline Self &stream(const cudaStream_t &stream)#: Update the stream.

inline const cudaStream_t &stream() const#: Expose the stream.

inline void begin()#

Begin the capture.

Checked with EAGLE_CHECK_ALWAYS rather than EAGLE_CHECK: this is a cold, once-per-build control-plane call, and a swallowed failure here leaves the stream not capturing while the caller believes it is.

Enters the process-wide capture-depth scope (STOP-THE-LINE fix) only AFTER the begin call has actually succeeded — a failed begin never opened a capture, so it must not raise the depth.

inline cudaGraph_t end()#

End the capture and return the created graph.

Checked with EAGLE_CHECK_ALWAYS. cudaStreamEndCapture reports structural faults in the captured DAG — most importantly cudaErrorStreamCaptureUnjoined, raised when a forked branch stream was never joined back before the capture ended — through its return code, and writes a null cudaGraph_t. Under the debug-only EAGLE_CHECK that error was discarded and the null handle returned to the caller, so the fault only surfaced later during adoption or launch. It now throws aether::Error at the point of failure.

Exits the process-wide capture-depth scope (STOP-THE-LINE fix) UNCONDITIONALLY, before the error check below — the capture SESSION is over the moment cudaStreamEndCapture returns, whether or not its result reports success, so any deferred teardown queued during this capture must be free to drain even on this failure path.

Throws:: aether::Error – If the capture did not end cleanly.

class StreamRef#

#include <StreamRef.h>

Non-owning handle to a CUDA stream.

Where eagle::cuda::Stream is an owning RAII wrapper that creates and destroys a real cudaStream_t, this handle merely references an existing stream — the legacy default stream (0) unless constructed from another. It is the backend-generic eagle::Stream on the CUDA side: cheap to default-construct, trivially copyable, and exposing exactly the surface interface code needs — synchronize() plus a native() accessor returning the underlying cudaStream_t for stream-aware async copies. Downstream code names it only through the eagle::Stream alias so a single #ifdef selects the host variant; no raw cudaStream_t is guarded at any call site.

Public Types

using NativeT = cudaStream_t#: Underlying native stream type.

Public Functions

StreamRef() = default#: Default handle to the legacy default stream (0).

inline StreamRef(cudaStream_t stream)#: Wrap an existing native stream (e.g. an owning Stream’s cuda()), so a caller holding a real stream can pass a lightweight reference without transferring ownership.

inline void synchronize() const#: Block the host until all work on the referenced stream completes.

inline cudaStream_t native() const#: Return the underlying cudaStream_t for async memory copies.

struct VecBinding#

API reference — cuda

Contents

API reference — cuda#

Namespace eagle::cuda#

Namespace `eagle::cuda`#