Struct WgpuBackend 

pub struct WgpuBackend { /* private fields */ }

A real wgpu backend for vyre.

Implementations

impl WgpuBackend

pub fn adapter_info(&self) -> &AdapterInfo

Adapter information selected for this backend instance.

pub fn device_limits(&self) -> &Limits

Device limits for this backend instance.

pub fn acquire() -> Result<Self, BackendError>

Acquire the backend, probing adapters and returning a structured error when no compatible GPU is found.

pub fn acquire_adapter(index: usize) -> Result<Self, BackendError>

Acquire a backend bound to a specific enumerable adapter index.

pub fn device_queue(&self) -> Arc<(Device, Queue)>

Consumer-visible snapshot of the live wgpu device + queue.

pub fn force_device_lost(&self) -> Result<(), BackendError>

Test-only hook that marks the backend device as lost and invalidates caches tied to the current device generation.

pub fn invalidate_impacted_pipeline_cache( &self, intervention_mask: &[u32], rule_adj: &[u32], state: &[u32], join_rules: &[u32], n: u32, max_iterations: u32, pipeline_lineage_cell: &[u32], pipeline_keys: &[[u8; 32]], ) -> Result<(), BackendError>

Invalidate compiled pipeline artifacts selected by a rule-impact mask.

pub fn invalidate_pipeline_cache_for_changed_op( &self, changed_op_handle: u32, pipeline_lineage_cell: &[u32], pipeline_keys: &[[u8; 32]], ) -> Result<(), BackendError>

Convenience wrapper around Self::invalidate_impacted_pipeline_cache

pub fn invalidate_impacted_disk_cache( &self, intervention_mask: &[u32], rule_adj: &[u32], state: &[u32], join_rules: &[u32], n: u32, max_iterations: u32, pipeline_lineage_cell: &[u32], cache_keys: &[String], ) -> Result<(), BackendError>

Invalidate disk-cached pipeline artifacts selected by a rule-impact mask.

pub fn new() -> Result<Self, BackendError>

Create the backend if a GPU adapter is available.

pub fn shared() -> Result<Arc<Self>, BackendError>

Process-wide shared backend handle.

pub fn dispatch_borrowed_for_each_mapped_output<F>( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, visitor: F, ) -> Result<(), BackendError>

where F: FnMut(usize, &[u8]) -> Result<(), BackendError>,

Dispatch borrowed inputs and visit each mapped output byte slice.

pub fn dispatch_borrowed_for_each_pod_output<T, F>( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, visitor: F, ) -> Result<(), BackendError>

where T: Pod, F: FnMut(usize, &[T]) -> Result<(), BackendError>,

Dispatch borrowed inputs and visit each mapped output as a typed POD slice.

pub fn dispatch_speculative_prefilter_confirm<F>( &self, speculator: &AdaptiveSpeculator, plan: SpeculativeDispatchPlan<'_>, inputs: &[&[u8]], config: &DispatchConfig, confirm_serial: F, ) -> Result<SpeculativeDispatchOutcome, BackendError>

where F: FnMut(OutputBuffers) -> Result<OutputBuffers, BackendError>,

Dispatch a real prefilter/confirm scan through the adaptive speculative path.

pub fn dispatch_borrowed_batch( &self, jobs: &[(&Program, &[&[u8]], &DispatchConfig)], ) -> Result<Vec<Result<OutputBuffers, BackendError>>, BackendError>

Dispatch a batch of borrowed (Program, inputs, config) triples.

pub fn dispatch_borrowed_batch_into( &self, jobs: &[(&Program, &[&[u8]], &DispatchConfig)], outputs: &mut [OutputBuffers], ) -> Result<Vec<Result<(), BackendError>>, BackendError>

Dispatch a borrowed batch and write each job’s outputs into caller-owned per-job output buffers.

pub fn dispatch_batch( &self, jobs: &[(Program, Vec<Vec<u8>>, DispatchConfig)], ) -> Result<Vec<Result<OutputBuffers, BackendError>>, BackendError>

Dispatch an owned batch of (Program, inputs, config) triples.

pub fn compile_streaming( &self, program: &Program, config: DispatchConfig, ) -> Result<HostIngressStream, BackendError>

Compile a program into a host-ingress wgpu stream.

pub fn compile_persistent( &self, program: &Program, config: &DispatchConfig, ) -> Result<Arc<WgpuPipeline>, BackendError>

Compile a program into a persistent pipeline.

impl WgpuBackend

pub fn allocate_wgpu_device_buffer( &self, byte_len: usize, ) -> Result<Box<dyn DeviceBuffer>, BackendError>

Allocate a new GPU-resident buffer of byte_len bytes. The buffer is created with STORAGE | COPY_SRC | COPY_DST so it can participate in dispatch as either input or output and round- trip through upload_device_buffer / download_device_buffer.

Errors

Returns a backend error if the underlying wgpu allocation fails (e.g. byte_len exceeds device limits).

pub fn upload_wgpu_device_buffer( &self, buffer: &mut dyn DeviceBuffer, bytes: &[u8], ) -> Result<(), BackendError>

Upload bytes into a previously-allocated wgpu DeviceBuffer. Bytes shorter than the allocation are written at offset 0; the remainder of the buffer is left as-is. Bytes longer than the allocation are an error.

Errors

Returns a backend error when the buffer was not allocated by this backend, when bytes exceeds the buffer’s allocation, or when the wgpu queue write fails.

pub fn download_wgpu_device_buffer( &self, buffer: &dyn DeviceBuffer, ) -> Result<Vec<u8>, BackendError>

Download the full byte_len of a previously-allocated wgpu DeviceBuffer into a fresh Vec<u8>.

Errors

Returns a backend error when the buffer was not allocated by this backend or when the readback fails (typically: buffer missing COPY_SRC, which the standard allocator path includes).

pub fn free_wgpu_device_buffer( &self, buffer: Box<dyn DeviceBuffer>, ) -> Result<(), BackendError>

Free a previously-allocated wgpu DeviceBuffer. The wgpu allocation is released when the underlying Arcwgpu::Buffer reaches zero references - dropping the box here is sufficient.

Errors

Returns a backend error when the buffer was not allocated by this backend.

impl WgpuBackend

pub fn lower_to_backend_ir( &self, program: &Program, ) -> Result<WgpuProgram, LoweringError>

Lower core IR into the backend-owned wgpu IR.

pub fn lower_to_target<'a>(&self, bir: &'a WgpuProgram) -> &'a Module

Borrow the Naga module produced by lowering (WgpuProgram::from_program).

This avoids cloning the entire naga::Module; callers that need an owned copy can call .clone() explicitly.

impl WgpuBackend

pub fn compile(&self, program: &Program) -> Result<WgpuIR, BackendError>

Compile a program once for repeated dispatch.

pub fn dispatch_compiled( &self, compiled: &WgpuIR, inputs: &[MemoryRef<'_>], config: &DispatchConfig, ) -> Result<Vec<Memory>, BackendError>

Dispatch a previously compiled program artifact.

impl WgpuBackend

pub fn dispatch_wgsl( &self, wgsl: &str, input: &[u8], output_size: usize, workgroup_size: u32, ) -> Result<Vec<u8>, String>

Dispatch a raw WGSL compute shader.

Errors

Returns an actionable error when shader compilation, staging-buffer creation, command submission, or readback fails.

impl WgpuBackend

pub fn probe_op(&self, op: UnOp, input: &[u8]) -> Result<Vec<u8>, BackendError>

Dispatch a canonical one-op f32 unary probe and return raw output bytes.

Errors

Returns a backend error when input is not one f32, the op is not a supported f32 unary probe, or the WGSL dispatch/readback fails.

pub fn probe_op_many( &self, op: UnOp, inputs: &[f32], ) -> Result<Vec<f32>, BackendError>

Dispatch a canonical f32 unary probe over a batch of inputs.

This keeps parity tests from paying one GPU submission and readback per scalar sample. The generated WGSL is keyed by operation, so the backend pipeline cache reuses it across calls.

Errors

Returns a backend error when the operation is unsupported, the batch is too large for WebGPU dispatch dimensions, or dispatch/readback fails.

impl WgpuBackend

pub fn adapter_caps(&self) -> AdapterCaps

Optimizer-facing capability snapshot for this live backend.

Unlike adapter-only probes, this reflects the features that were actually enabled on the device after backend construction.

pub fn device_profile(&self) -> DeviceProfile

Driver-neutral capability profile for this live backend.

pub fn stats(&self) -> WgpuBackendStats

Observability snapshot - pipeline cache size, buffer-pool stats, and adapter identity. SRE-friendly: consumers feed the returned numbers into prometheus / OpenTelemetry / Datadog pipelines for dashboards and alerting.

Reads use atomic cache counters and the lock-free persistent-pool pointer, so the call is safe for metrics-scrape loops.

Trait Implementations

impl BackendValidationCapabilities for WgpuBackend

fn backend_name(&self) -> &'static str

Stable backend name used in diagnostics.

fn supports_cast_target(&self, target: &DataType) -> bool

Return true when the backend can lower a cast whose destination is target.

fn supports_subgroup_ops(&self) -> bool

Return true when the backend supports subgroup operations.

fn supports_indirect_dispatch(&self) -> bool

Return true when the backend supports indirect dispatch.

fn supports_specialization_constants(&self) -> bool

Return true when the backend supports specialization constants.

fn supports_distributed_collectives(&self) -> bool

Return true when the backend supports distributed collective nodes.

fn backend_capabilities(&self) -> BackendCapabilities

Export backend capabilities in a version-stable value object.

impl Clone for WgpuBackend

fn clone(&self) -> WgpuBackend

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for WgpuBackend

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Executable for WgpuBackend

fn dispatch( &self, program: &Program, inputs: &[MemoryRef<'_>], config: &DispatchConfig, ) -> Result<Vec<Memory>, BackendError>

Dispatch a validated program.

impl OptimizerDispatcher for WgpuBackend

fn dispatch( &self, program: &Program, inputs: &[Vec<u8>], grid_override: Option<[u32; 3]>, ) -> Result<Vec<Vec<u8>>, DispatchError>

Dispatch program with the given byte inputs (one Vec<u8> per declared input buffer in canonical buffer order). Returns the declared outputs in the same canonical order.

fn supports_persistent(&self) -> bool

Whether this dispatcher supports the persistent-resident path. Default: false. CUDA backend overrides to true. The orchestrator uses this to decide whether to take the persistent fast-path (encode arena once → upload once → dispatch many → readback once) or use the non-resident per-call GPU dispatch path.

fn device_feature_cache_key(&self) -> u64

Device/lowering feature bits that affect reusable plan identity.

fn alloc_resident(&self, _byte_len: usize) -> Result<u64, DispatchError>

Allocate a backend-resident buffer. Returns an opaque u64 handle. Callers must free_resident to release.

fn alloc_resident_many( &self, byte_lens: &[usize], ) -> Result<Vec<u64>, DispatchError>

Allocate a logical group of resident buffers and roll back partial state if any allocation fails.

fn upload_resident( &self, _handle: u64, _bytes: &[u8], ) -> Result<(), DispatchError>

Upload host bytes into a resident buffer.

fn upload_resident_many( &self, uploads: &[(u64, &[u8])], ) -> Result<(), DispatchError>

Upload several resident buffers with one backend fence when supported.

fn acquire_resident_static_uploads( &self, _cache_domain: u64, payloads: &[&[u8]], ) -> Result<ResidentStaticBufferSet, DispatchError>

Acquire resident handles for immutable payloads.

fn release_resident_static_uploads( &self, set: ResidentStaticBufferSet, ) -> Result<(), DispatchError>

Release a static resident buffer set acquired from Self::acquire_resident_static_uploads.

fn read_resident(&self, _handle: u64) -> Result<Vec<u8>, DispatchError>

Download a resident buffer’s current contents to host bytes.

fn read_resident_many( &self, handles: &[u64], ) -> Result<Vec<Vec<u8>>, DispatchError>

Download several resident buffers with one backend fence when supported.

fn read_resident_ranges( &self, ranges: &[ResidentReadRange], ) -> Result<Vec<Vec<u8>>, DispatchError>

Download selected byte ranges from resident buffers.

fn read_resident_ranges_into( &self, ranges: &[ResidentReadRange], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Download selected byte ranges from resident buffers into caller-owned byte slots.

fn free_resident(&self, _handle: u64) -> Result<(), DispatchError>

Free a resident buffer previously returned by alloc_resident.

fn dispatch_resident( &self, _program: &Program, _handles: &[u64], _grid_override: Option<[u32; 3]>, ) -> Result<(), DispatchError>

Dispatch a Program against resident-buffer handles. Each handle is referenced from the Program’s declared buffer in the same canonical buffer order. RW buffers are not read back - caller invokes read_resident once at end of pipeline.

fn dispatch_resident_sequence( &self, steps: &[ResidentDispatchStep<'_>], ) -> Result<(), DispatchError>

Dispatch an ordered sequence of resident-buffer Programs.

fn dispatch_resident_sequence_read_many( &self, steps: &[ResidentDispatchStep<'_>], read_handles: &[u64], ) -> Result<Vec<Vec<u8>>, DispatchError>

Dispatch an ordered resident sequence and read selected resident buffers.

fn dispatch_resident_sequence_read_ranges( &self, steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange], ) -> Result<Vec<Vec<u8>>, DispatchError>

Dispatch an ordered resident sequence and read selected byte ranges.

fn upload_resident_many_sequence_read_many( &self, uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_handles: &[u64], ) -> Result<Vec<Vec<u8>>, DispatchError>

Upload resident buffers, dispatch an ordered resident sequence, then read selected resident buffers.

fn upload_resident_many_sequence_read_ranges( &self, uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange], ) -> Result<Vec<Vec<u8>>, DispatchError>

Upload resident buffers, dispatch an ordered resident sequence, then read selected byte ranges.

fn upload_resident_many_sequence_read_many_into( &self, uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_handles: &[u64], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Same contract as Self::upload_resident_many_sequence_read_many, but writes readbacks into caller-owned byte slots.

fn clear_upload_resident_many_sequence_read_many_into( &self, clears: &[(u64, usize)], uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_handles: &[u64], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Same contract as Self::upload_resident_many_sequence_read_many_into, but first clears full resident buffers to zero.

fn fill_upload_resident_many_sequence_read_many_into( &self, fills: &[(u64, usize, u8)], uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_handles: &[u64], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Same contract as Self::clear_upload_resident_many_sequence_read_many_into, but fills each resident buffer with an arbitrary byte value.

fn fill_upload_resident_many_sequence_read_ranges_into( &self, fills: &[(u64, usize, u8)], uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Same contract as Self::upload_resident_many_sequence_read_ranges_into, but fills resident buffers first. CUDA overrides this to use device memset and compact D2H range copies on the same stream.

fn upload_resident_many_sequence_read_ranges_into( &self, uploads: &[(u64, &[u8])], steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange], outputs: &mut Vec<Vec<u8>>, ) -> Result<(), DispatchError>

Same contract as Self::upload_resident_many_sequence_read_ranges, but writes compact readbacks into caller-owned byte slots.

impl Sealed for WgpuBackend

impl VyreBackend for WgpuBackend

fn id(&self) -> &'static str

Stable backend identifier used for logging, certificates, and adapter selection.

fn version(&self) -> &'static str

Backend implementation version string used for certificates and regression tracking.

fn supported_ops(&self) -> &HashSet<OpId>

Operation ids this backend can execute without further lowering.

fn dispatch( &self, program: &Program, inputs: &[Vec<u8>], config: &DispatchConfig, ) -> Result<Vec<Vec<u8>>, BackendError>

Executes the program with the given input buffers and returns the output buffers.

fn dispatch_borrowed( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, ) -> Result<Vec<Vec<u8>>, BackendError>

Executes the program with borrowed input buffers.

fn dispatch_borrowed_into( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, outputs: &mut OutputBuffers, ) -> Result<(), BackendError>

Executes the program with borrowed input buffers and writes outputs into caller-owned storage.

fn dispatch_borrowed_timed( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, ) -> Result<TimedDispatchResult, BackendError>

Executes a borrowed-input dispatch and returns backend-owned timing.

fn dispatch_async( &self, program: &Program, inputs: &[Vec<u8>], config: &DispatchConfig, ) -> Result<Box<dyn PendingDispatch>, BackendError>

Non-blocking dispatch primitive.

fn dispatch_borrowed_async( &self, program: &Program, inputs: &[&[u8]], config: &DispatchConfig, ) -> Result<Box<dyn PendingDispatch>, BackendError>

Non-blocking dispatch with borrowed input buffers.

fn compile_native( &self, program: &Program, config: &DispatchConfig, ) -> Result<Option<Arc<dyn CompiledPipeline>>, BackendError>

Optional pre-compilation hook for the pipeline-mode API.

fn allocate_device_buffer( &self, byte_len: usize, ) -> Result<Box<dyn DeviceBuffer>, BackendError>

Allocate a backend-owned device buffer of byte_len bytes.

fn upload_device_buffer( &self, buffer: &mut dyn DeviceBuffer, bytes: &[u8], ) -> Result<(), BackendError>

Upload host bytes into a previously-allocated device buffer.

fn download_device_buffer( &self, buffer: &dyn DeviceBuffer, ) -> Result<Vec<u8>, BackendError>

Download bytes from a device buffer back to a host Vec<u8>.

fn free_device_buffer( &self, buffer: Box<dyn DeviceBuffer>, ) -> Result<(), BackendError>

Free a device buffer previously returned by Self::allocate_device_buffer. Explicit-free is required because the substrate does not assume reference-counted backend handles; consumers are responsible for calling this when done.

fn allocate_resident(&self, byte_len: usize) -> Result<Resource, BackendError>

Allocate a backend-resident buffer and return a stable resource handle.

fn upload_resident( &self, resource: &Resource, bytes: &[u8], ) -> Result<(), BackendError>

Upload bytes into a backend-resident resource.

fn upload_resident_many( &self, uploads: &[(&Resource, &[u8])], ) -> Result<(), BackendError>

Upload several backend-resident resources as one logical staging operation.

fn upload_resident_at( &self, resource: &Resource, dst_offset_bytes: usize, bytes: &[u8], ) -> Result<(), BackendError>

Upload bytes into a subrange of a backend-resident resource.

fn upload_resident_at_many( &self, uploads: &[(&Resource, usize, &[u8])], ) -> Result<(), BackendError>

Upload several resident subranges as one logical staging operation.

fn download_resident( &self, resource: &Resource, ) -> Result<Vec<u8>, BackendError>

Download a backend-resident resource into a new host buffer.

fn download_resident_into( &self, resource: &Resource, out: &mut Vec<u8>, ) -> Result<(), BackendError>

Download a backend-resident resource into caller-owned storage.

fn download_resident_range( &self, resource: &Resource, byte_offset: usize, byte_len: usize, ) -> Result<Vec<u8>, BackendError>

Download a byte range from a backend-resident resource into a new host buffer.

fn download_resident_range_into( &self, resource: &Resource, byte_offset: usize, byte_len: usize, out: &mut Vec<u8>, ) -> Result<(), BackendError>

Download a byte range from a backend-resident resource into caller-owned storage.

fn download_resident_ranges_into( &self, ranges: &[(&Resource, usize, usize)], outputs: &mut [&mut Vec<u8>], ) -> Result<(), BackendError>

Download several byte ranges from backend-resident resources into caller-owned storage as one logical readback operation.

fn free_resident(&self, resource: Resource) -> Result<(), BackendError>

Free a backend-resident resource previously returned by VyreBackend::allocate_resident.

fn dispatch_resident_timed( &self, program: &Program, resources: &[Resource], config: &DispatchConfig, ) -> Result<TimedDispatchResult, BackendError>

Dispatch using backend-resident resources and return backend-owned timing.

fn dispatch_with_device_buffers( &self, program: &Program, inputs: &[&dyn DeviceBuffer], outputs: &mut [&mut dyn DeviceBuffer], config: &DispatchConfig, ) -> Result<(), BackendError>

Dispatch a Program with backend-owned device buffers as inputs and outputs.

fn pipeline_cache_snapshot(&self) -> Option<PipelineCacheSnapshot>

Optional compiled-pipeline cache counters for compile telemetry.

fn supports_subgroup_ops(&self) -> bool

Whether this backend’s lowering path emits subgroup / wave intrinsics AND the current adapter exposes them.

fn supports_f16(&self) -> bool

Whether this backend lowers IEEE 754 binary16 (DataType::F16) natively rather than emulating through f32.

fn supports_bf16(&self) -> bool

Whether this backend lowers bfloat16 (DataType::BF16) natively.

fn supports_tensor_cores(&self) -> bool

Whether this backend emits tensor-core / matrix-engine intrinsics for supported tensor shapes.

fn supports_async_compute(&self) -> bool

Whether this backend overlaps copies and compute via independent queues or async engines.

fn supports_indirect_dispatch(&self) -> bool

Whether this backend supports indirect dispatch (Node::IndirectDispatch).

fn supports_speculation(&self) -> bool

Whether this backend supports speculative dispatch - a fused prefilter + confirmer kernel with commit-gated output and a counter tail read back by the host.

fn supports_persistent_thread_dispatch(&self) -> bool

Whether this backend supports device-side persistent-thread dispatch (a long-running kernel that polls a work queue).

fn is_distributed(&self) -> bool

Whether this backend partitions a program across more than one physical device / node.

fn max_workgroup_size(&self) -> [u32; 3]

Maximum supported workgroup size per axis [x, y, z].

fn max_compute_workgroups_per_dimension(&self) -> u32

Maximum number of compute workgroups the backend can launch in one dispatch dimension.

fn max_compute_invocations_per_workgroup(&self) -> u32

Maximum total invocations allowed in a single workgroup.

fn subgroup_size(&self) -> Option<u32>

Native subgroup size for the backing device when the backend knows it. Returning None tells the dispatch planner the backend can’t report a subgroup width - the planner falls back to max_workgroup_size for its sizing heuristic.

fn max_storage_buffer_bytes(&self) -> u64

Maximum size in bytes of a single storage buffer the backend accepts. 0 means the backend has not reported a limit, not “unlimited”.

fn device_profile(&self) -> DeviceProfile

Unified backend-neutral device profile.

fn flush(&self) -> Result<(), BackendError>

Flush any queued work to the device and wait for it to complete.

fn device_lost(&self) -> bool

Probe whether the underlying device has been lost since the last successful dispatch.

fn try_recover(&self) -> Result<(), BackendError>

Attempt to recover from device loss by reacquiring the underlying device and invalidating pipeline caches.

fn dispatch_resident_sequence_read_ranges_into( &self, steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange<'_>], outputs: &mut [&mut Vec<u8>], ) -> Result<(), BackendError>

Dispatch an ordered sequence of resident-buffer programs and read selected resident byte ranges into caller-owned storage.

fn dispatch_resident_sequence_read_ranges_timed_into( &self, steps: &[ResidentDispatchStep<'_>], read_ranges: &[ResidentReadRange<'_>], outputs: &mut [&mut Vec<u8>], ) -> Result<ResidentSequenceTiming, BackendError>

Timed variant of VyreBackend::dispatch_resident_sequence_read_ranges_into.

fn dispatch_resident_repeated_sequence_read_ranges_into( &self, prefix_steps: &[ResidentDispatchStep<'_>], repeated_steps: &[ResidentDispatchStep<'_>], repeat_count: u32, read_ranges: &[ResidentReadRange<'_>], outputs: &mut [&mut Vec<u8>], ) -> Result<(), BackendError>

Dispatch a resident prefix, repeat a resident sub-sequence, and read selected resident byte ranges into caller-owned storage.

fn compile_native_shared( &self, program: Arc<Program>, config: &DispatchConfig, ) -> Result<Option<Arc<dyn CompiledPipeline>>, BackendError>

Optional pre-compilation hook for callers that already own a shared program allocation.

fn backend_metric_snapshot(&self) -> Vec<(&'static str, u64)>

Optional backend-specific numeric telemetry for release evidence.

fn supports_grid_sync(&self) -> bool

Whether this backend can satisfy Node::Barrier { ordering: MemoryOrdering::GridSync } inside a single dispatch - i.e. every thread in the entire grid waits at the barrier and every prior write is globally visible afterwards. Backends that lack a native grid barrier (workgroup-only fences) must return false; registration-based dispatch may lower a GridSync barrier to a host-orchestrated kernel split only when VyreBackend::allows_host_grid_sync_split also returns true.

fn cooperative_grid_sync_fits( &self, _program: &Program, _inputs: &[&[u8]], _config: &DispatchConfig, ) -> Result<bool, BackendError>

Whether a native cooperative grid-sync launch of program with these inputs and config can be made fully resident on this device.

fn allows_host_grid_sync_split(&self) -> bool

Whether the shared registry wrapper may emulate whole-grid synchronization for this backend by splitting one program into multiple host-dispatched kernels.

fn supports_resident_dispatch(&self) -> bool

Whether this backend implements the resident half of the contract (allocate_resident / upload_resident / dispatch_resident_timed / dispatch_resident_repeated_sequence_read_ranges_into / download_resident_* / free_resident) well enough to run a device-resident dispatch sequence.

fn supports_distributed_collectives(&self) -> bool

Whether this backend lowers distributed collective communication nodes (AllReduce, AllGather, ReduceScatter, Broadcast).

fn prepare(&self) -> Result<(), BackendError>

Pre-dispatch warmup. Called before the first dispatch on a new program so the backend can warm caches, compile ahead-of-time, or acquire a device handle without paying that cost on the hot path.

fn shutdown(&self) -> Result<(), BackendError>

Release device resources held by this backend. After shutdown returns the backend is in an unspecified state and may not be used for further dispatches.