vyre_driver/backend/dispatch_config.rs
1//! Immutable dispatch policy supplied by callers before backend execution.
2
3use std::time::Duration;
4
5/// Immutable execution policy supplied by the caller before dispatch.
6///
7/// `DispatchConfig` is an additive, non-exhaustive struct so that new backend
8/// options (conformance profiles, adapter hints, etc.) can be added without
9/// breaking the frozen `VyreBackend::dispatch` signature. Backends must treat
10/// every field as read-only policy and must not assume the presence of any
11/// particular option.
12///
13/// # Examples
14///
15/// ```
16/// use vyre::DispatchConfig;
17///
18/// // DispatchConfig is `#[non_exhaustive]`; construct it through
19/// // `default()` and overwrite the fields you want to change.
20/// let mut config = DispatchConfig::default();
21/// config.profile = Some("stress".to_string());
22/// config.ulp_budget = None;
23/// ```
24#[derive(Clone, Debug, Default, PartialEq, Eq)]
25#[non_exhaustive]
26pub struct DispatchConfig {
27 /// Optional stable profile identifier such as `default`, `stress`, or a
28 /// backend-defined conformance mode.
29 pub profile: Option<String>,
30 /// Optional maximum ULP error budget for approximate transcendental lowering.
31 ///
32 /// `None` and `Some(0)` require the strict target-text intrinsic path. A positive
33 /// budget allows backends to select fast approximate intrinsic wrappers only
34 /// when the wrapper contract is bounded by the supplied ULP ceiling.
35 pub ulp_budget: Option<u8>,
36 /// Optional timeout for the dispatch.
37 pub timeout: Option<Duration>,
38 /// Optional label for the dispatch (for debugging/profiling).
39 pub label: Option<String>,
40 /// Optional maximum output byte limit.
41 pub max_output_bytes: Option<usize>,
42 /// Optional workgroup size override.
43 ///
44 /// When `Some`, the backend uses the supplied `[x, y, z]` workgroup size
45 /// instead of the one declared on the [`vyre_foundation::ir::Program`].
46 /// This lets callers tune workgroup sizing at dispatch time without
47 /// cloning the program metadata. When `None` (the default), the backend
48 /// falls back to `Program::workgroup_size`.
49 pub workgroup_override: Option<[u32; 3]>,
50 /// Optional grid size override (number of workgroups).
51 ///
52 /// When set, the backend launches the supplied workgroup count instead of
53 /// the one inferred from the program's output buffer size. This is
54 /// required for megakernels where the work queue length is managed through
55 /// storage buffers rather than the primary output slot.
56 pub grid_override: Option<[u32; 3]>,
57 /// True per-invocation element/byte coverage count for an element-grid
58 /// dispatch (e.g. a one-lane-per-byte scan: `Some(haystack_len)`).
59 ///
60 /// This exists SEPARATELY from [`grid_override`](Self::grid_override) because
61 /// that field is OVERLOADED: for an element-grid dispatch it is the workgroup
62 /// count derived from the input size, but for a MEGAKERNEL it is a work-queue
63 /// length managed through storage buffers, the two cannot be told apart from
64 /// the `[u32; 3]` alone. Backends that infer their dispatch coverage from
65 /// buffer SHAPES rather than from a real GPU grid (the CPU reference
66 /// interpreter, [`CpuRefBackend`](../../../vyre_driver_reference/index.html))
67 /// cannot see the runtime scan length, so a byte-scan program would be
68 /// under-dispatched to `haystack_len / 4` invocations and SILENTLY skip high
69 /// positions (a Law-10 recall regression). An element-grid caller sets this to
70 /// the true coverage so such a backend dispatches exactly what the GPU would;
71 /// `None` (the default, and every megakernel) means "infer from buffer shapes"
72 ///: so a megakernel is never over-run by a byte count that is not its grid.
73 pub dispatch_elements: Option<u32>,
74 /// True per-workgroup-axis dispatch grid `[x, y, z]` for a multi-dimensional
75 /// element dispatch.
76 ///
77 /// This is the N-dimensional counterpart of
78 /// [`dispatch_elements`](Self::dispatch_elements) (a 1-D floor). A backend that
79 /// infers its coverage from buffer SHAPES rather than a real GPU grid (the CPU
80 /// reference interpreter,
81 /// [`CpuRefBackend`](../../../vyre_driver_reference/index.html)) distributes the
82 /// dispatch only across workgroup axes whose size is greater than one, so a
83 /// program that fans a `[256, 1, 1]` workgroup across `grid.y` (batched
84 /// persistent-BFS runs one query per `grid.y` block) would collapse to
85 /// `grid.y == 1` and SILENTLY compute only the first query (a Law-10
86 /// under-coverage). A caller that knows the real grid, e.g.
87 /// `persistent_bfs_batch_dispatch_grid(node_count, query_count)`, sets it here so
88 /// the interpreter covers every workgroup the GPU would. `None` (the default)
89 /// keeps buffer-shape inference. When both this and `dispatch_elements` are set,
90 /// this wins because it fully specifies the grid.
91 pub dispatch_grid: Option<[u32; 3]>,
92 /// Maximum back-to-back dispatch iterations the backend should run on
93 /// the same persistent input/output handles before reading back the
94 /// final outputs.
95 ///
96 /// `None` means one iteration. `Some(0)` is invalid: backends must reject
97 /// it instead of silently rewriting caller policy.
98 pub fixpoint_iterations: Option<u32>,
99 /// Optional speculation policy.
100 pub speculation: Option<crate::speculate::SpeculationMode>,
101 /// Optional persistent-thread dispatch policy.
102 pub persistent_thread: Option<crate::persistent::PersistentThreadMode>,
103 /// Whether the backend should launch through its cooperative-grid API.
104 ///
105 /// A backend MUST reject `cooperative = true` with `UnsupportedFeature`
106 /// when its `VyreBackend::supports_grid_sync()` returns `false`.
107 pub cooperative: bool,
108}
109
110impl DispatchConfig {
111 /// Construct a `DispatchConfig` from explicit fields in one call.
112 /// Complement to `DispatchConfig::default()` for external crates
113 /// that want all optional fields set up front.
114 #[must_use]
115 pub fn new(
116 profile: Option<String>,
117 ulp_budget: Option<u8>,
118 timeout: Option<Duration>,
119 label: Option<String>,
120 ) -> Self {
121 Self {
122 profile,
123 ulp_budget,
124 timeout,
125 label,
126 max_output_bytes: None,
127 workgroup_override: None,
128 grid_override: None,
129 dispatch_elements: None,
130 dispatch_grid: None,
131 fixpoint_iterations: None,
132 speculation: None,
133 persistent_thread: None,
134 cooperative: false,
135 }
136 }
137}