#[non_exhaustive]pub struct DispatchConfig {Show 13 fields
pub profile: Option<String>,
pub ulp_budget: Option<u8>,
pub timeout: Option<Duration>,
pub label: Option<String>,
pub max_output_bytes: Option<usize>,
pub workgroup_override: Option<[u32; 3]>,
pub grid_override: Option<[u32; 3]>,
pub dispatch_elements: Option<u32>,
pub dispatch_grid: Option<[u32; 3]>,
pub fixpoint_iterations: Option<u32>,
pub speculation: Option<SpeculationMode>,
pub persistent_thread: Option<PersistentThreadMode>,
pub cooperative: bool,
}Expand description
Immutable execution policy supplied by the caller before dispatch.
DispatchConfig is an additive, non-exhaustive struct so that new backend
options (conformance profiles, adapter hints, etc.) can be added without
breaking the frozen VyreBackend::dispatch signature. Backends must treat
every field as read-only policy and must not assume the presence of any
particular option.
§Examples
use vyre::DispatchConfig;
// DispatchConfig is `#[non_exhaustive]`; construct it through
// `default()` and overwrite the fields you want to change.
let mut config = DispatchConfig::default();
config.profile = Some("stress".to_string());
config.ulp_budget = None;Fields (Non-exhaustive)§
This struct is marked as non-exhaustive
Struct { .. } syntax; cannot be matched against without a wildcard ..; and struct update syntax will not work.profile: Option<String>Optional stable profile identifier such as default, stress, or a
backend-defined conformance mode.
ulp_budget: Option<u8>Optional maximum ULP error budget for approximate transcendental lowering.
None and Some(0) require the strict target-text intrinsic path. A positive
budget allows backends to select fast approximate intrinsic wrappers only
when the wrapper contract is bounded by the supplied ULP ceiling.
timeout: Option<Duration>Optional timeout for the dispatch.
label: Option<String>Optional label for the dispatch (for debugging/profiling).
max_output_bytes: Option<usize>Optional maximum output byte limit.
workgroup_override: Option<[u32; 3]>Optional workgroup size override.
When Some, the backend uses the supplied [x, y, z] workgroup size
instead of the one declared on the vyre_foundation::ir::Program.
This lets callers tune workgroup sizing at dispatch time without
cloning the program metadata. When None (the default), the backend
falls back to Program::workgroup_size.
grid_override: Option<[u32; 3]>Optional grid size override (number of workgroups).
When set, the backend launches the supplied workgroup count instead of the one inferred from the program’s output buffer size. This is required for megakernels where the work queue length is managed through storage buffers rather than the primary output slot.
dispatch_elements: Option<u32>True per-invocation element/byte coverage count for an element-grid
dispatch (e.g. a one-lane-per-byte scan: Some(haystack_len)).
This exists SEPARATELY from grid_override because
that field is OVERLOADED: for an element-grid dispatch it is the workgroup
count derived from the input size, but for a MEGAKERNEL it is a work-queue
length managed through storage buffers, the two cannot be told apart from
the [u32; 3] alone. Backends that infer their dispatch coverage from
buffer SHAPES rather than from a real GPU grid (the CPU reference
interpreter, CpuRefBackend)
cannot see the runtime scan length, so a byte-scan program would be
under-dispatched to haystack_len / 4 invocations and SILENTLY skip high
positions (a Law-10 recall regression). An element-grid caller sets this to
the true coverage so such a backend dispatches exactly what the GPU would;
None (the default, and every megakernel) means “infer from buffer shapes”
: so a megakernel is never over-run by a byte count that is not its grid.
dispatch_grid: Option<[u32; 3]>True per-workgroup-axis dispatch grid [x, y, z] for a multi-dimensional
element dispatch.
This is the N-dimensional counterpart of
dispatch_elements (a 1-D floor). A backend that
infers its coverage from buffer SHAPES rather than a real GPU grid (the CPU
reference interpreter,
CpuRefBackend) distributes the
dispatch only across workgroup axes whose size is greater than one, so a
program that fans a [256, 1, 1] workgroup across grid.y (batched
persistent-BFS runs one query per grid.y block) would collapse to
grid.y == 1 and SILENTLY compute only the first query (a Law-10
under-coverage). A caller that knows the real grid, e.g.
persistent_bfs_batch_dispatch_grid(node_count, query_count), sets it here so
the interpreter covers every workgroup the GPU would. None (the default)
keeps buffer-shape inference. When both this and dispatch_elements are set,
this wins because it fully specifies the grid.
fixpoint_iterations: Option<u32>Maximum back-to-back dispatch iterations the backend should run on the same persistent input/output handles before reading back the final outputs.
None means one iteration. Some(0) is invalid: backends must reject
it instead of silently rewriting caller policy.
speculation: Option<SpeculationMode>Optional speculation policy.
persistent_thread: Option<PersistentThreadMode>Optional persistent-thread dispatch policy.
cooperative: boolWhether the backend should launch through its cooperative-grid API.
A backend MUST reject cooperative = true with UnsupportedFeature
when its VyreBackend::supports_grid_sync() returns false.
Implementations§
Trait Implementations§
Source§impl Clone for DispatchConfig
impl Clone for DispatchConfig
Source§fn clone(&self) -> DispatchConfig
fn clone(&self) -> DispatchConfig
1.0.0 (const: unstable) · Source§fn clone_from(&mut self, source: &Self)
fn clone_from(&mut self, source: &Self)
source. Read moreSource§impl Debug for DispatchConfig
impl Debug for DispatchConfig
Source§impl Default for DispatchConfig
impl Default for DispatchConfig
Source§fn default() -> DispatchConfig
fn default() -> DispatchConfig
impl Eq for DispatchConfig
Source§impl PartialEq for DispatchConfig
impl PartialEq for DispatchConfig
impl StructuralPartialEq for DispatchConfig
Auto Trait Implementations§
impl Freeze for DispatchConfig
impl RefUnwindSafe for DispatchConfig
impl Send for DispatchConfig
impl Sync for DispatchConfig
impl Unpin for DispatchConfig
impl UnsafeUnpin for DispatchConfig
impl UnwindSafe for DispatchConfig
Blanket Implementations§
Source§impl<T> BorrowMut<T> for Twhere
T: ?Sized,
impl<T> BorrowMut<T> for Twhere
T: ?Sized,
Source§fn borrow_mut(&mut self) -> &mut T
fn borrow_mut(&mut self) -> &mut T
Source§impl<T> CloneToUninit for Twhere
T: Clone,
impl<T> CloneToUninit for Twhere
T: Clone,
Source§impl<Q, K> Equivalent<K> for Q
impl<Q, K> Equivalent<K> for Q
Source§impl<Q, K> Equivalent<K> for Q
impl<Q, K> Equivalent<K> for Q
Source§fn equivalent(&self, key: &K) -> bool
fn equivalent(&self, key: &K) -> bool
key and return true if they are equal.