Contents

class

SushiRuntime::API::Region

Records one mutable partition of a DynamicGraph through Graph's add() surface.

Declared in
include/SushiRuntime/api/graph/region.hpp
Inherits
SushiRuntime::API::Graph

See also

include/SushiRuntime/api/graph/README.md

Public member functions

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_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< 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

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_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 <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_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 <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

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

std::size_t size() const noexcept

Reports how much work has been recorded.

Returns

Number of update rules (tasks) added so far.

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