Contents

class

SushiRuntime::API::Graph

A chainable graph of tasks built fluently, then run.

Declared in
include/SushiRuntime/api/graph/graph.hpp

Every add() overload appends one task and returns *this so calls chain. The concept-constrained overloads give a readable compiler error when a callable has the wrong signature — something a macro cannot do. Obtain one from Runtime::graph(); it is moveable but not copyable.

Static public attributes

static constexpr std::uint32_t no_registration = BackingGraph::NO_REGISTRATION

Stands for "this rule has no index yet".

Public member functions

Graph(Graph &&other)=default

Takes over another graph's recorded tasks.

Parameters

other

Graph to move from; left with no tasks.

Graph & operator=(Graph &&other)=default

Replaces this graph's contents with another's.

Parameters

other

Graph to move from; left with no tasks.

Returns

*this, now owning other's tasks.

Graph(const Graph &)=delete
Graph & operator=(const Graph &)=delete
template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add_untracked(std::size_t n, F &&fn)

Adds a per-element task that records no regions and is ordered against nothing.

Parameters

fn

Callable invoked as fn(i) for each index i.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_range_kernel_v< F >, int > = 0>
Graph & add_untracked(std::size_t n, F &&fn)

Adds a range task that records no regions; fn receives bounds and handler.

Parameters

n

Upper bound; fn is called as fn(0, n, handler).

template <class F, std::enable_if_t< is_raw_kernel_v< F >, int > = 0>
Graph & add_untracked(F &&fn)

Adds a raw kernel that records no regions and drives the handler itself.

Parameters

fn

Callable invoked as fn(sycl::handler&).

template <class F, std::enable_if_t< is_graph_emitter_v< F >, int > = 0>
Graph & add_emitter(F &&emit)

Adds an emitter that records its own nodes on the backing graph.

Parameters

emit

Callable invoked as emit(BackingGraph&) at every build of the plan.

Precondition

emit only adds nodes: it does not compile, drain or keep the graph.

See also

include/SushiRuntime/api/graph/README.md

template <int N, class F, std::enable_if_t< std::is_invocable_v< F, sycl::id< N > >, int > = 0>
Graph & add_untracked(Extent< N > e, F &&fn)

Adds a multidimensional task over an Extent that records no regions.

Template parameters

N

Iteration rank (1, 2, or 3).

Parameters

fn

Callable invoked as fn(sycl::id<N>) per work-item.

See also

include/SushiRuntime/api/graph/README.md

template <int N, class F, std::enable_if_t< std::is_invocable_v< F, sycl::nd_item< N > >, int > = 0>
Graph & add_untracked(NdRange< N > r, F &&fn)

Adds a tiled task over an NdRange that records no regions.

Template parameters

N

Iteration rank (1, 2, or 3).

Parameters

r

Global and local ranges; global must tile by local.

fn

Callable invoked as fn(sycl::nd_item<N>) per work-item.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(const Reads &reads, const Writes &writes, std::size_t n, F &&fn)

Adds a 1-D per-element task ordered by the handles in Reads and Writes.

Parameters

n

Iteration range; fn is called for each i in [0, n).

See also

include/SushiRuntime/api/graph/README.md

template <int N, class F, std::enable_if_t< std::is_invocable_v< F, sycl::id< N > >, int > = 0>
Graph & add(const Reads &reads, const Writes &writes, Extent< N > e, F &&fn)

Adds an N-D per-element task over an Extent, ordered by Reads and Writes.

Template parameters

N

Iteration rank (1, 2, or 3).

Parameters

fn

Callable invoked as fn(sycl::id<N>) per work-item.

template <int N, class F, std::enable_if_t< std::is_invocable_v< F, sycl::nd_item< N > >, int > = 0>
Graph & add(const Reads &reads, const Writes &writes, NdRange< N > r, F &&fn)

Adds a tiled NdRange task ordered by the regions it declares.

Template parameters

N

Iteration rank (1, 2, or 3).

Parameters

r

Global and local ranges; global must tile by local.

fn

Callable invoked as fn(sycl::nd_item<N>) per work-item.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_range_kernel_v< F >, int > = 0>
Graph & add(const Reads &reads, const Writes &writes, std::size_t n, F &&fn)

Adds a range task ordered by its declared regions; fn drives the handler over [0, n).

Parameters

fn

Callable invoked as fn(begin, end, sycl::handler&).

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_raw_kernel_v< F >, int > = 0>
Graph & add(const Reads &reads, const Writes &writes, F &&fn)

Adds a raw kernel ordered by its declared regions; fn owns the whole command group.

Parameters

fn

Callable invoked as fn(sycl::handler&).

See also

include/SushiRuntime/api/graph/README.md

template <int N, class Cap0, class... Args, std::enable_if_t< detail::is_capture_v< Cap0 > &&(sizeof...(Args) >=1), int > = 0>
Graph & add(Extent< N > e, Cap0 c0, Args &&... args)

Adds an N-D per-element task that registers its captures and passes their pointers.

Template parameters

N

Iteration rank (1, 2, or 3).

Parameters

args

Remaining captured arguments, with the callable last.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class F, std::enable_if_t< std::is_invocable_v< F, std::size_t, const T *, T * >, int > = 0>
Graph & add(State< T > &s, F &&fn)

Adds the per-step update rule of a State: reads current, writes next.

Parameters

s

The state this rule evolves; must outlive the graph.

fn

Callable invoked as fn(i, const T* cur, T* next) for each i in [0, s.size()).

See also

include/SushiRuntime/api/graph/README.md

template <typename T, int N, class F, std::enable_if_t< std::is_invocable_v< F, sycl::id< N >, const T *, T * >, int > = 0>
Graph & add(State< T > &s, Extent< N > e, F &&fn)

Adds a multidimensional per-step update rule for a State over an Extent.

Parameters

s

The state this rule evolves; must outlive the graph.

fn

Callable invoked as fn(sycl::id<N>, const T* cur, T* next).

See also

include/SushiRuntime/api/graph/README.md

template <typename T, typename A, class... Rest, std::enable_if_t<(sizeof...(Rest) >=1), int > = 0>
Graph & add(State< T > &out, State< A > &in0, Rest &&... rest)

Adds a coupled update rule that writes out's next field from several current fields.

Parameters

out

The field this rule writes; must outlive the graph.

in0

First read field; must outlive the graph.

rest

Further read fields, then the callable fn(i, cur, next, a, ...) last.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(const std::vector< void * > &reads, const std::vector< void * > &writes, std::size_t n, F &&fn)

Adds a per-element task ordered by the raw base pointers it reads and writes.

Parameters

n

Iteration range; fn is called for each i in [0, n).

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(Dynamic dyn, const Reads &reads, const Writes &writes, std::size_t capacity, F &&fn)

Adds a late-bound per-element task whose size and enablement vary per step.

Parameters

capacity

Pre-sized upper bound; the live size must not exceed it.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(Dynamic dyn, const std::vector< void * > &reads, const std::vector< void * > &writes, std::size_t capacity, F &&fn)

Adds a late-bound per-element task that names its data by raw base pointer.

Parameters

capacity

Pre-sized upper bound; the live size must not exceed it.

template <typename T, class F, std::enable_if_t< std::is_invocable_v< F, std::size_t, const T *, T * >, int > = 0>
Graph & add(Dynamic dyn, State< T > &s, F &&fn)

Adds a late-bound per-step update rule for a State.

Parameters

s

The state this rule evolves; must outlive the graph.

fn

Callable invoked as fn(i, const T* cur, T* next).

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< std::is_invocable_v< F >, int > = 0>
Graph & add_host(const Reads &reads, const Writes &writes, F &&fn)

Adds a pure-CPU task ordered by the data it touches.

Parameters

fn

Nullary callable run once; must not submit SYCL work.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< std::is_invocable_v< F >, int > = 0>
Graph & add_host(F &&fn)

Adds an untracked pure-CPU task that runs every step.

Parameters

fn

Nullary callable run once per step; must not submit SYCL work.

See also

include/SushiRuntime/api/graph/README.md

Graph & name_last(const char *name) noexcept

Names the task most recently added to this graph.

Parameters

name

Persistent label, stored not copied; ignored when no task has been added.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_reduce(const Buffer< T > &input, const Buffer< T > &output, std::size_t n, Combine op, T identity)

Adds a reduction whose combination order never varies.

Template parameters

Combine

Binary functor; must be trivially copyable.

Parameters

input

Values to reduce; must hold at least n elements.

Exceptions

Core::SushiException

If n exceeds input's size, if output is empty, or if scratch cannot be allocated.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_reduce(const Buffer< T > &input, const Buffer< T > &output, std::size_t n, Combine op)

Adds a fixed-order reduction using the combiner's static identity().

Exceptions

Core::SushiException

As the five-argument overload.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_reduce(const Buffer< T > &input, const Buffer< T > &output, LateCount n, Combine op, T identity)

Adds a reduction whose live element count is late-bound.

Parameters

input

Values to reduce; must hold at least n.capacity elements.

Exceptions

Core::SushiException

If n.capacity exceeds input's size, if output is empty, or at run if the live count exceeds n.capacity.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_reduce(const Buffer< T > &input, const Buffer< T > &output, LateCount n, Combine op)

Adds a late-bound reduction using the combiner's static identity().

Parameters

input

Values to reduce; must hold at least n.capacity elements.

Exceptions

Core::SushiException

As the five-argument overload.

template <typename T, class Combine>
Graph & add_segmented_reduce(const Buffer< T > &values, const Buffer< std::uint32_t > &segment_offsets, const Buffer< T > &output, std::size_t segment_count, Combine op, T identity)

Reduces each segment of a ragged array in a fixed order.

Parameters

segment_offsets

segment_count + 1 non-decreasing entries, the last within values.

Exceptions

Core::SushiException

If segment_offsets or output is too small for segment_count.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_segmented_reduce(const Buffer< T > &values, const Buffer< std::uint32_t > &segment_offsets, const Buffer< T > &output, std::size_t segment_count, Combine op)

Reduces each segment using the combiner's static identity().

Parameters

segment_offsets

segment_count + 1 non-decreasing entries.

output

Holds segment_count entries.

Exceptions

Core::SushiException

As the six-argument overload.

template <typename T, class Combine>
Graph & add_segmented_reduce(const Buffer< T > &values, const Buffer< std::uint32_t > &segment_offsets, const Buffer< T > &output, LateCount segment_count, Combine op, T identity)

Reduces each segment with a late-bound segment count.

Parameters

segment_offsets

segment_count.capacity + 1 non-decreasing entries.

output

Holds segment_count.capacity entries.

Exceptions

Core::SushiException

If segment_offsets or output is too small for segment_count.capacity, or at run if the live count exceeds it.

See also

include/SushiRuntime/api/graph/README.md

template <typename T, class Combine>
Graph & add_segmented_reduce(const Buffer< T > &values, const Buffer< std::uint32_t > &segment_offsets, const Buffer< T > &output, LateCount segment_count, Combine op)

Reduces each segment, late-bound, using the combiner's static identity().

Parameters

segment_offsets

segment_count.capacity + 1 non-decreasing entries.

output

Holds segment_count.capacity entries.

Exceptions

Core::SushiException

As the late-bound six-argument overload.

Graph & add_prefix_scan(const Buffer< std::uint32_t > &counts, const Buffer< std::uint32_t > &offsets, const Buffer< std::uint32_t > &total, LateCount n)

Scans uint32 counts into exclusive running offsets, late-bound.

Parameters

offsets

Destination; must hold at least n.capacity + 1 elements.

Exceptions

Core::SushiException

If a buffer is too small, if counts and offsets sit on different devices, or at run if the live count exceeds n.capacity.

See also

include/SushiRuntime/api/graph/README.md

Graph & add_counting_sort(const Buffer< std::uint32_t > &keys, const Buffer< std::uint32_t > &indices, const Buffer< std::uint32_t > &offsets, const Buffer< std::uint32_t > &total, LateCount n, LateCount key_range, const Buffer< std::uint32_t > *counts=nullptr)

Stably argsorts keys by 8-bit digits into bucket offsets.

Exceptions

Core::SushiException

If a buffer is too small, if keys, indices, offsets or counts sit on different devices, if key_range.capacity is zero, or at run if either live count exceeds its capacity.

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(::SushiRuntime::Graph::OpID op, std::size_t n, const std::vector< void * > &reads, const std::vector< void * > &writes, F &&local, std::initializer_list< std::uint64_t > params={})

Adds an offloadable task: a named remote kernel with a local fallback.

Parameters

op

OpID of the kernel, e.g. "vec.add"_op.

local

Local fallback invoked as local(i) for i in [0, n).

See also

include/SushiRuntime/api/graph/README.md

template <class F, std::enable_if_t< is_per_element_v< F >, int > = 0>
Graph & add(::SushiRuntime::Graph::OpID op, std::size_t n, const Reads &reads, const Writes &writes, F &&local, std::initializer_list< std::uint64_t > params={})

Adds an offloadable task that names its data with Reads and Writes.

Parameters

local

Local fallback invoked as local(i) for i in [0, n).

template <typename T, class F, std::enable_if_t< std::is_invocable_v< F, std::size_t, const T *, T * >, int > = 0>
Graph & add(::SushiRuntime::Graph::OpID op, State< T > &s, F &&local, std::initializer_list< std::uint64_t > params={})

Adds an offloadable per-step update rule for a State.

Precondition

s was registered with State::distributed().

Parameters

s

The state this rule evolves; must outlive the graph.

local

Local fallback invoked as local(i, const T* cur, T* next).

See also

include/SushiRuntime/api/graph/README.md

std::size_t size() const noexcept

Reports how much work has been recorded.

Returns

Number of update rules (tasks) added so far.

std::uint32_t rule_registration(std::size_t rule) const noexcept

Returns the registration index the rule at position rule took.

Parameters

rule

Position in add() order, below size().

Returns

The index, or no_registration when rule is out of range or nothing compiled yet.

See also

include/SushiRuntime/api/graph/README.md

std::size_t compile_count() const noexcept

Number of times the graph has been compiled.

Exposes the replay guarantee: a graph that only changes late-bound size or enablement between steps recompiles zero extra times, so this stays at 1 across a whole run after the first build.

Returns

The compile count since this Graph was created.

void set_profiling_requested(bool requested) noexcept

Requests or withdraws per-node device timing and marks the plan dirty.

See also

include/SushiRuntime/api/graph/README.md

bool profiling_requested() const noexcept

Reports whether this graph currently requests profiling.

Returns

The last value set by set_profiling_requested(), or the runtime's RuntimeSettings::profiling default.

Core::RunReport run()

Runs the graph once, blocking until it completes.

Returns

A RunReport summarising the run.

Exceptions

Core::SushiException

on any synchronous or device error.

Core::RunReport run(std::size_t steps)

Runs the given number of successive steps and blocks until they finish.

Returns

A RunReport aggregating the whole run.

Exceptions

Core::SushiException

on any synchronous or device error.

See also

include/SushiRuntime/api/graph/README.md

Core::RunReport & run(Core::RunReport &reuse)

Runs the graph once into a report the caller owns.

Parameters

reuse

The report to overwrite; its previous contents are lost.

Exceptions

Core::SushiException

on any synchronous or device error.

See also

include/SushiRuntime/api/graph/README.md

Core::RunReport & run(std::size_t steps, Core::RunReport &reuse)

Runs steps successive steps into a report the caller owns.

Parameters

steps

Number of steps to run.

reuse

The report to overwrite; its previous contents are lost.

Returns

A reference to reuse.

Exceptions

Core::SushiException

on any synchronous or device error.

template <class Rep, class Period>
Core::RunReport run_for(std::chrono::duration< Rep, Period > budget)

Runs steps until a wall-clock budget elapses.

Parameters

budget

Wall-clock limit; at least one step always runs.

Returns

A RunReport aggregating the steps that ran.

Exceptions

Core::SushiException

on any synchronous or device error.

template <class Predicate>
Core::RunReport run_until(Predicate done)

Runs steps until done returns true.

Parameters

done

Stop condition checked before each step; the run ends the first time it returns true.

Returns

A RunReport aggregating the steps that ran.

Exceptions

Core::SushiException

on any synchronous or device error.

RunHandle< Graph > run_async()

Starts one step on the worker pool and returns without waiting for it.

Precondition

The returned handle does not outlive this graph.

Exceptions

Core::SushiException

If the graph has no tasks, references a freed Buffer/State, the engine does not support asynchronous runs, or a run is already in progress.

See also

include/SushiRuntime/api/graph/README.md

void cancel()

Requests cancellation of an in-progress run, from any thread, without throwing.

See also

include/SushiRuntime/api/graph/README.md

Core::Result< Core::RunReport, Core::Error > try_run() noexcept

Non-throwing single run.

Returns

The RunReport on success, or an Error on failure.

Core::Result< Core::RunReport, Core::Error > try_run(std::size_t steps) noexcept

Non-throwing multi-step run.

Parameters

steps

Number of steps to run.

Returns

The RunReport on success, or an Error on failure.

Core::Result< Core::RunReport *, Core::Error > try_run(Core::RunReport &reuse) noexcept

Non-throwing single run into a report the caller owns.

Parameters

reuse

The report to overwrite; its previous contents are lost.

Returns

A pointer to reuse on success, or an Error on failure.

Core::Result< Core::RunReport *, Core::Error > try_run(std::size_t steps, Core::RunReport &reuse) noexcept

Non-throwing multi-step run into a report the caller owns.

Parameters

steps

Number of steps to run.

reuse

The report to overwrite; its previous contents are lost.

Returns

A pointer to reuse on success, or an Error on failure.