runmat_accelerate_api/

lib.rs

use anyhow::anyhow;
use once_cell::sync::{Lazy, OnceCell};
use serde::{Deserialize, Serialize};
#[cfg(not(target_arch = "wasm32"))]
use std::cell::Cell;
use std::collections::{HashMap, HashSet};
use std::future::Future;
use std::pin::Pin;
use std::sync::atomic::{AtomicU32, Ordering};
#[cfg(feature = "wgpu")]
use std::sync::Arc;
#[cfg(target_arch = "wasm32")]
use std::sync::Mutex;
use std::sync::RwLock;

mod placement;

pub use placement::*;

type ResidencyMarkFn = fn(&GpuTensorHandle);
type ResidencyClearFn = fn(&GpuTensorHandle);
type SequenceThresholdFn = fn() -> Option<usize>;
type WorkgroupSizeHintFn = fn() -> Option<u32>;

static RESIDENCY_MARK: OnceCell<ResidencyMarkFn> = OnceCell::new();
static RESIDENCY_CLEAR: OnceCell<ResidencyClearFn> = OnceCell::new();
static SEQUENCE_THRESHOLD_PROVIDER: OnceCell<SequenceThresholdFn> = OnceCell::new();
static WORKGROUP_SIZE_HINT_PROVIDER: OnceCell<WorkgroupSizeHintFn> = OnceCell::new();

pub type GpuHandleIdentity = (u32, u64);

static LOGICAL_HANDLES: Lazy<RwLock<HashSet<GpuHandleIdentity>>> =
    Lazy::new(|| RwLock::new(HashSet::new()));
static LOGICAL_HANDLE_HITS: Lazy<RwLock<HashMap<GpuHandleIdentity, u64>>> =
    Lazy::new(|| RwLock::new(HashMap::new()));
static TRANSPOSED_HANDLES: Lazy<RwLock<HashMap<GpuHandleIdentity, TransposeInfo>>> =
    Lazy::new(|| RwLock::new(HashMap::new()));

static HANDLE_CLASS_NAMES: Lazy<RwLock<HashMap<GpuHandleIdentity, String>>> =
    Lazy::new(|| RwLock::new(HashMap::new()));

pub const fn handle_identity(handle: &GpuTensorHandle) -> GpuHandleIdentity {
    (handle.device_id, handle.buffer_id)
}

#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub struct TransposeInfo {
    pub base_rows: usize,
    pub base_cols: usize,
}

/// Register a callback used to mark residency tracking when GPU tensors are
/// created or returned by device-side execution paths.
pub fn register_residency_mark(handler: ResidencyMarkFn) {
    let _ = RESIDENCY_MARK.set(handler);
}

/// Mark residency metadata for the provided GPU tensor handle, if a backend
/// has registered a handler via [`register_residency_mark`].
pub fn mark_residency(handle: &GpuTensorHandle) {
    if let Some(handler) = RESIDENCY_MARK.get() {
        handler(handle);
    }
}

/// Register a callback used to clear residency tracking when GPU tensors are
/// gathered back to the host. Backends that maintain residency metadata should
/// install this hook during initialization.
pub fn register_residency_clear(handler: ResidencyClearFn) {
    let _ = RESIDENCY_CLEAR.set(handler);
}

/// Clear residency metadata for the provided GPU tensor handle, if a backend
/// has registered a handler via [`register_residency_clear`].
pub fn clear_residency(handle: &GpuTensorHandle) {
    if let Some(handler) = RESIDENCY_CLEAR.get() {
        handler(handle);
    }
}

/// Register a callback that exposes the current sequence length threshold
/// derived from the auto-offload planner. Array constructors can use this hint
/// to decide when to prefer GPU residency automatically.
pub fn register_sequence_threshold_provider(provider: SequenceThresholdFn) {
    let _ = SEQUENCE_THRESHOLD_PROVIDER.set(provider);
}

/// Query the currently registered sequence threshold hint, if any.
pub fn sequence_threshold_hint() -> Option<usize> {
    SEQUENCE_THRESHOLD_PROVIDER
        .get()
        .and_then(|provider| provider())
}

/// Register a callback that reports the calibrated workgroup size selected by
/// the active acceleration provider (if any). Plotting kernels can reuse this
/// hint to match backend tuning.
pub fn register_workgroup_size_hint_provider(provider: WorkgroupSizeHintFn) {
    let _ = WORKGROUP_SIZE_HINT_PROVIDER.set(provider);
}

/// Query the current workgroup size hint exposed by the provider.
pub fn workgroup_size_hint() -> Option<u32> {
    WORKGROUP_SIZE_HINT_PROVIDER
        .get()
        .and_then(|provider| provider())
}

/// Export a shared acceleration context (e.g., the active WGPU device) when the
/// current provider exposes one.
pub fn export_context(kind: AccelContextKind) -> Option<AccelContextHandle> {
    provider().and_then(|p| p.export_context(kind))
}

/// Request a provider-owned WGPU buffer for zero-copy consumers. Returns `None`
/// when the active provider does not expose buffers or does not support the
/// supplied handle.
#[cfg(feature = "wgpu")]
pub fn export_wgpu_buffer(handle: &GpuTensorHandle) -> Option<WgpuBufferRef> {
    provider_for_handle(handle).and_then(|p| p.export_wgpu_buffer(handle))
}

/// Return the floating precision encoded by the handle's durable element type.
pub fn handle_precision(handle: &GpuTensorHandle) -> Option<ProviderPrecision> {
    handle
        .descriptor
        .element_type
        .and_then(NumericElementType::precision)
}

/// Record the MATLAB underlying class associated with a GPU tensor handle.
///
/// Precision alone cannot represent integer gpuArray classes, so runtime
/// introspection and validation use this metadata when the upload path knows
/// the exact requested or inferred class.
pub fn set_handle_class_name(handle: &GpuTensorHandle, class_name: impl Into<String>) {
    if let Ok(mut guard) = HANDLE_CLASS_NAMES.write() {
        guard.insert(handle_identity(handle), class_name.into());
    }
}

/// Look up the recorded MATLAB underlying class for a GPU tensor handle.
pub fn handle_class_name(handle: &GpuTensorHandle) -> Option<String> {
    HANDLE_CLASS_NAMES
        .read()
        .ok()
        .and_then(|guard| guard.get(&handle_identity(handle)).cloned())
        .or_else(|| {
            handle
                .descriptor
                .element_type
                .map(|element_type| element_type.class_name().to_string())
        })
}

/// Clear any recorded MATLAB underlying class metadata for a GPU tensor handle.
pub fn clear_handle_class_name(handle: &GpuTensorHandle) {
    if let Ok(mut guard) = HANDLE_CLASS_NAMES.write() {
        guard.remove(&handle_identity(handle));
    }
}

/// Annotate a GPU tensor handle as logically-typed (`logical` in MATLAB terms)
/// or clear the logical flag when `logical` is `false`.
pub fn set_handle_logical(handle: &GpuTensorHandle, logical: bool) {
    let identity = handle_identity(handle);
    if let Ok(mut guard) = LOGICAL_HANDLES.write() {
        if logical {
            guard.insert(identity);
            if let Ok(mut hits) = LOGICAL_HANDLE_HITS.write() {
                *hits.entry(identity).or_insert(0) += 1;
            }
        } else {
            guard.remove(&identity);
            if let Ok(mut hits) = LOGICAL_HANDLE_HITS.write() {
                hits.remove(&identity);
            }
        }
    }
    if logical {
        if let Ok(mut classes) = HANDLE_CLASS_NAMES.write() {
            classes.insert(identity, "logical".into());
        }
    } else if let Ok(mut classes) = HANDLE_CLASS_NAMES.write() {
        if classes
            .get(&identity)
            .is_some_and(|class| class == "logical")
        {
            classes.remove(&identity);
        }
    }
}

/// Convenience helper for clearing logical annotations explicitly.
pub fn clear_handle_logical(handle: &GpuTensorHandle) {
    set_handle_logical(handle, false);
}

/// Returns true when the supplied handle has been marked as logical.
pub fn handle_is_logical(handle: &GpuTensorHandle) -> bool {
    LOGICAL_HANDLES
        .read()
        .map(|guard| guard.contains(&handle_identity(handle)))
        .unwrap_or(false)
}

pub fn handle_logical_hits(buffer_id: u64) -> Option<u64> {
    LOGICAL_HANDLE_HITS.read().ok().and_then(|guard| {
        let mut matches = guard
            .iter()
            .filter_map(|((_, candidate), hits)| (*candidate == buffer_id).then_some(*hits));
        let first = matches.next()?;
        Some(matches.fold(first, u64::saturating_add))
    })
}

pub fn handle_logical_hits_for_handle(handle: &GpuTensorHandle) -> Option<u64> {
    LOGICAL_HANDLE_HITS
        .read()
        .ok()
        .and_then(|guard| guard.get(&handle_identity(handle)).copied())
}

pub fn record_handle_transpose(handle: &GpuTensorHandle, base_rows: usize, base_cols: usize) {
    if let Ok(mut guard) = TRANSPOSED_HANDLES.write() {
        guard.insert(
            handle_identity(handle),
            TransposeInfo {
                base_rows,
                base_cols,
            },
        );
    }
}

pub fn clear_handle_transpose(handle: &GpuTensorHandle) {
    if let Ok(mut guard) = TRANSPOSED_HANDLES.write() {
        guard.remove(&handle_identity(handle));
    }
}

pub fn handle_transpose_info(handle: &GpuTensorHandle) -> Option<TransposeInfo> {
    TRANSPOSED_HANDLES
        .read()
        .ok()
        .and_then(|guard| guard.get(&handle_identity(handle)).copied())
}

pub fn handle_is_transposed(handle: &GpuTensorHandle) -> bool {
    handle_transpose_info(handle).is_some()
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum GpuTensorStorage {
    Real,
    ComplexInterleaved,
}

/// Exact native integer element storage for host/device transfer contracts.
///
/// This intentionally lives in the acceleration API instead of depending on
/// `runmat-builtins`, so providers can move integer buffers without first
/// materializing RunMat's lossy floating compatibility view.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum IntegerElementType {
    I8,
    I16,
    I32,
    I64,
    U8,
    U16,
    U32,
    U64,
}

impl IntegerElementType {
    pub const fn element_size(self) -> usize {
        match self {
            Self::I8 | Self::U8 => 1,
            Self::I16 | Self::U16 => 2,
            Self::I32 | Self::U32 => 4,
            Self::I64 | Self::U64 => 8,
        }
    }
}

/// Look up the exact native integer class stored by a GPU tensor handle.
pub fn handle_integer_type(handle: &GpuTensorHandle) -> Option<IntegerElementType> {
    handle
        .descriptor
        .element_type
        .and_then(NumericElementType::integer_type)
}

impl Default for GpuTensorStorage {
    fn default() -> Self {
        Self::Real
    }
}

/// Durable physical description of a provider-owned numeric buffer.
///
/// Provider-created numeric handles populate element type and storage at
/// allocation time so cloning or serializing a handle cannot discard its
/// physical interpretation. A missing physical field is invalid at numeric
/// provider boundaries. Provenance is populated by ownership-aware runtime
/// boundaries.
#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
pub struct GpuTensorDescriptor {
    pub element_type: Option<NumericElementType>,
    pub storage: Option<GpuTensorStorage>,
    pub provenance: Option<GpuHandleProvenance>,
}

impl GpuTensorDescriptor {
    pub const fn numeric(element_type: NumericElementType, storage: GpuTensorStorage) -> Self {
        Self {
            element_type: Some(element_type),
            storage: Some(storage),
            provenance: None,
        }
    }
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct GpuTensorHandle {
    pub shape: Vec<usize>,
    /// Stable provider/device namespace identifier. Distinct live provider
    /// instances that can own handles must use distinct identifiers because
    /// provider routing is keyed by this value.
    pub device_id: u32,
    pub buffer_id: u64,
    #[serde(default)]
    pub descriptor: GpuTensorDescriptor,
}

impl GpuTensorHandle {
    pub fn new(shape: Vec<usize>, device_id: u32, buffer_id: u64) -> Self {
        Self {
            shape,
            device_id,
            buffer_id,
            descriptor: GpuTensorDescriptor::default(),
        }
    }

    pub fn with_numeric_descriptor(
        mut self,
        element_type: NumericElementType,
        storage: GpuTensorStorage,
    ) -> Self {
        self.descriptor.element_type = Some(element_type);
        self.descriptor.storage = Some(storage);
        self
    }

    pub fn with_provenance(mut self, provenance: GpuHandleProvenance) -> Self {
        self.descriptor.provenance = Some(provenance);
        self
    }
}

/// Semantic origin of a resident handle.
///
/// Automatic residency is an implementation detail and may transparently return
/// to the host. Explicit residency represents user-visible `gpuArray` intent and
/// must be retained by operations that preserve gpuArray semantics.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum GpuHandleProvenance {
    Automatic,
    Explicit,
}

pub fn set_handle_provenance(handle: &mut GpuTensorHandle, provenance: GpuHandleProvenance) {
    handle.descriptor.provenance = Some(provenance);
}

pub fn handle_provenance(handle: &GpuTensorHandle) -> Option<GpuHandleProvenance> {
    handle.descriptor.provenance
}

pub fn mark_handle_explicit(handle: &mut GpuTensorHandle) {
    set_handle_provenance(handle, GpuHandleProvenance::Explicit);
}

pub fn mark_handle_automatic(handle: &mut GpuTensorHandle) {
    set_handle_provenance(handle, GpuHandleProvenance::Automatic);
}

pub fn handle_is_explicit(handle: &GpuTensorHandle) -> bool {
    handle_provenance(handle) == Some(GpuHandleProvenance::Explicit)
}

/// Clear mutable API side annotations for a released or freshly reused handle,
/// including backend residency tracking registered through
/// [`register_residency_clear`]. The handle-owned physical descriptor remains
/// available on existing handle values.
pub fn clear_handle_metadata(handle: &GpuTensorHandle) {
    clear_residency(handle);
    clear_handle_class_name(handle);
    clear_handle_logical(handle);
    clear_handle_transpose(handle);
}

#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderSpectralRange {
    Onesided,
    Twosided,
    Centered,
}

#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderSpectralFrameMode {
    Sliding {
        hop: usize,
    },
    ColumnSliding {
        hop: usize,
        input_rows: usize,
        frames_per_column: usize,
    },
    FoldedColumns {
        input_rows: usize,
    },
}

#[derive(Clone, Debug)]
pub struct ProviderSpectralRequest<'a> {
    pub input: &'a GpuTensorHandle,
    pub input_len: usize,
    pub input_complex: bool,
    pub window: &'a [f64],
    pub nfft: usize,
    pub frame_count: usize,
    pub frame_mode: ProviderSpectralFrameMode,
    pub range: ProviderSpectralRange,
    pub denominator: f64,
}

#[derive(Clone, Debug)]
pub struct ProviderSpectralResult {
    pub s: GpuTensorHandle,
    pub ps: GpuTensorHandle,
    pub rows: usize,
    pub cols: usize,
}

#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderEnvelopeMethod {
    Analytic,
    AnalyticFir { filter_len: usize },
    Rms { window_len: usize },
}

#[derive(Clone, Debug)]
pub struct ProviderEnvelopeRequest<'a> {
    pub input: &'a GpuTensorHandle,
    pub channel_len: usize,
    pub channel_count: usize,
    pub output_shape: &'a [usize],
    pub method: ProviderEnvelopeMethod,
}

#[derive(Clone, Debug)]
pub struct ProviderEnvelopeResult {
    pub upper: GpuTensorHandle,
    pub lower: GpuTensorHandle,
}

#[derive(Clone, Debug)]
pub struct ProviderHilbertRequest<'a> {
    pub input: &'a GpuTensorHandle,
    /// Optional FFT length along `dim`.
    pub length: Option<usize>,
    /// Zero-based transform dimension.
    pub dim: usize,
}

#[derive(Clone, Debug)]
pub struct ProviderModulationRequest<'a> {
    pub input: &'a GpuTensorHandle,
    /// `(real, imag)` pairs interleaved by symbol index.
    pub constellation: &'a [f64],
}

#[derive(Clone, Debug)]
pub struct ProviderBitModulationRequest<'a> {
    pub input: &'a GpuTensorHandle,
    /// Number of bit rows in the input before grouping.
    pub input_rows: usize,
    /// Number of input bits that form one output symbol.
    pub bits_per_symbol: usize,
    /// `(real, imag)` pairs interleaved by symbol index.
    pub constellation: &'a [f64],
}

pub async fn uniform_spectral_estimate(
    request: ProviderSpectralRequest<'_>,
) -> anyhow::Result<ProviderSpectralResult> {
    validate_uniform_spectral_request(&request)?;

    let provider = provider_for_handle(request.input)
        .ok_or_else(|| anyhow!("uniform_spectral_estimate: GPU provider unavailable"))?;
    provider.uniform_spectral_estimate(&request).await
}

fn validate_uniform_spectral_request(request: &ProviderSpectralRequest<'_>) -> anyhow::Result<()> {
    let invalid_frame_mode = matches!(
        request.frame_mode,
        ProviderSpectralFrameMode::Sliding { hop: 0 }
            | ProviderSpectralFrameMode::ColumnSliding { hop: 0, .. }
    );
    let invalid_input_coverage = match request.frame_mode {
        ProviderSpectralFrameMode::Sliding { hop } => {
            let required = request
                .frame_count
                .checked_sub(1)
                .and_then(|frames| frames.checked_mul(hop))
                .and_then(|offset| offset.checked_add(request.window.len()));
            required.is_none_or(|required| required > request.input_len)
        }
        ProviderSpectralFrameMode::ColumnSliding {
            hop,
            input_rows,
            frames_per_column,
        } => {
            if input_rows == 0 || frames_per_column == 0 {
                true
            } else {
                let last_frame = request.frame_count - 1;
                let source_col = last_frame / frames_per_column;
                let segment = last_frame % frames_per_column;
                source_col
                    .checked_mul(input_rows)
                    .and_then(|base| {
                        segment
                            .checked_mul(hop)
                            .and_then(|step| base.checked_add(step))
                    })
                    .and_then(|offset| offset.checked_add(request.window.len()))
                    .is_none_or(|required| required > request.input_len)
            }
        }
        ProviderSpectralFrameMode::FoldedColumns { input_rows } => {
            input_rows == 0
                || input_rows
                    .checked_mul(request.frame_count)
                    .is_none_or(|required| required > request.input_len)
        }
    };
    if request.window.is_empty()
        || request.nfft == 0
        || request.frame_count == 0
        || invalid_frame_mode
        || invalid_input_coverage
        || !request.denominator.is_finite()
        || request.denominator <= 0.0
    {
        return Err(anyhow!("uniform_spectral_estimate: invalid request"));
    }

    Ok(())
}

pub async fn signal_envelope(
    request: ProviderEnvelopeRequest<'_>,
) -> anyhow::Result<ProviderEnvelopeResult> {
    let expected_len = request
        .channel_len
        .checked_mul(request.channel_count)
        .ok_or_else(|| anyhow!("signal_envelope: invalid request"))?;
    let output_len = request
        .output_shape
        .iter()
        .try_fold(1usize, |acc, &dim| acc.checked_mul(dim))
        .ok_or_else(|| anyhow!("signal_envelope: invalid request"))?;
    let input_len = request
        .input
        .shape
        .iter()
        .try_fold(1usize, |acc, &dim| acc.checked_mul(dim))
        .ok_or_else(|| anyhow!("signal_envelope: invalid request"))?;
    if request.channel_len == 0
        || request.channel_count == 0
        || request.output_shape.is_empty()
        || output_len != expected_len
        || input_len != expected_len
        || !provider_envelope_input_shape_matches(
            &request.input.shape,
            request.channel_len,
            request.channel_count,
        )
    {
        return Err(anyhow!("signal_envelope: invalid request"));
    }

    match request.method {
        ProviderEnvelopeMethod::AnalyticFir { filter_len }
        | ProviderEnvelopeMethod::Rms {
            window_len: filter_len,
        } if filter_len == 0 => return Err(anyhow!("signal_envelope: invalid request")),
        _ => {}
    }

    let provider = provider_for_handle(request.input)
        .ok_or_else(|| anyhow!("signal_envelope: GPU provider unavailable"))?;
    provider.signal_envelope(&request).await
}

fn provider_envelope_input_shape_matches(
    shape: &[usize],
    channel_len: usize,
    channel_count: usize,
) -> bool {
    if channel_count == 1 {
        return match shape {
            [len] => *len == channel_len,
            [rows, cols] => {
                (*rows == channel_len && *cols == 1) || (*rows == 1 && *cols == channel_len)
            }
            _ => false,
        };
    }

    matches!(shape, [rows, cols] if *rows == channel_len && *cols == channel_count)
}

pub async fn signal_hilbert(
    request: ProviderHilbertRequest<'_>,
) -> anyhow::Result<GpuTensorHandle> {
    if request.length == Some(0) {
        return Err(anyhow!("signal_hilbert: invalid request"));
    }
    if request.dim >= request.input.shape.len() {
        return Err(anyhow!("signal_hilbert: invalid request"));
    }

    let provider = provider_for_handle(request.input)
        .ok_or_else(|| anyhow!("signal_hilbert: GPU provider unavailable"))?;
    provider.signal_hilbert(&request).await
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ApiDeviceInfo {
    pub device_id: u32,
    pub name: String,
    pub vendor: String,
    pub memory_bytes: Option<u64>,
    pub backend: Option<String>,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ReduceDimResult {
    pub values: GpuTensorHandle,
    pub indices: GpuTensorHandle,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderCumminResult {
    pub values: GpuTensorHandle,
    pub indices: GpuTensorHandle,
}

/// Result payload returned by provider-side `cummax` scans.
///
/// Alias of [`ProviderCumminResult`] because both operations return the same pair of tensors
/// (running values and MATLAB-compatible indices).
pub type ProviderCummaxResult = ProviderCumminResult;

/// Names a shared acceleration context that callers may request (e.g. plotting).
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum AccelContextKind {
    Plotting,
}

/// Handle returned by [`export_context`] that describes a shared GPU context.
#[derive(Clone)]
pub enum AccelContextHandle {
    #[cfg(feature = "wgpu")]
    Wgpu(WgpuContextHandle),
}

impl AccelContextHandle {
    /// Returns the underlying WGPU context when available.
    #[cfg(feature = "wgpu")]
    pub fn as_wgpu(&self) -> Option<&WgpuContextHandle> {
        match self {
            AccelContextHandle::Wgpu(ctx) => Some(ctx),
        }
    }
}

/// Shared WGPU device/queue pair exported by the acceleration provider.
#[cfg(feature = "wgpu")]
#[derive(Clone)]
pub struct WgpuContextHandle {
    pub instance: Arc<wgpu::Instance>,
    pub device: Arc<wgpu::Device>,
    pub queue: Arc<wgpu::Queue>,
    pub adapter: Arc<wgpu::Adapter>,
    pub adapter_info: wgpu::AdapterInfo,
    pub limits: wgpu::Limits,
    pub features: wgpu::Features,
}

/// Borrowed reference to a provider-owned WGPU buffer corresponding to a `GpuTensorHandle`.
///
/// Providers that expose this handle must ensure the buffer was created with
/// `wgpu::BufferUsages::COPY_SRC`, so zero-copy consumers can export or inspect
/// tensor data without depending on provider-specific readback APIs.
#[cfg(feature = "wgpu")]
#[derive(Clone)]
pub struct WgpuBufferRef {
    pub buffer: Arc<wgpu::Buffer>,
    pub len: usize,
    pub shape: Vec<usize>,
    pub element_size: usize,
    pub precision: ProviderPrecision,
}

pub fn handle_storage(handle: &GpuTensorHandle) -> GpuTensorStorage {
    handle.descriptor.storage.unwrap_or_default()
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum PagefunOp {
    Mtimes,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct PagefunRequest {
    pub op: PagefunOp,
    pub inputs: Vec<GpuTensorHandle>,
    pub output_shape: Vec<usize>,
    pub page_dims: Vec<usize>,
    pub input_page_dims: Vec<Vec<usize>>,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum FindDirection {
    First,
    Last,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderFindResult {
    pub linear: GpuTensorHandle,
    pub rows: GpuTensorHandle,
    pub cols: GpuTensorHandle,
    pub values: Option<GpuTensorHandle>,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderBandwidth {
    pub lower: u32,
    pub upper: u32,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderSymmetryKind {
    Symmetric,
    Skew,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderHermitianKind {
    Hermitian,
    Skew,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderLuResult {
    pub combined: GpuTensorHandle,
    pub lower: GpuTensorHandle,
    pub upper: GpuTensorHandle,
    pub perm_matrix: GpuTensorHandle,
    pub perm_vector: GpuTensorHandle,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderCholResult {
    pub factor: GpuTensorHandle,
    /// MATLAB-compatible failure index (0 indicates success).
    pub info: u32,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderQrResult {
    pub q: GpuTensorHandle,
    pub r: GpuTensorHandle,
    pub perm_matrix: GpuTensorHandle,
    pub perm_vector: GpuTensorHandle,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderQrPowerIterResult {
    pub q: GpuTensorHandle,
    pub r: GpuTensorHandle,
    pub perm_matrix: GpuTensorHandle,
    pub perm_vector: GpuTensorHandle,
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize, Default)]
pub struct ProviderLinsolveOptions {
    pub lower: bool,
    pub upper: bool,
    pub rectangular: bool,
    pub transposed: bool,
    pub conjugate: bool,
    pub symmetric: bool,
    pub posdef: bool,
    pub need_rcond: bool,
    pub rcond: Option<f64>,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderLinsolveResult {
    pub solution: GpuTensorHandle,
    pub reciprocal_condition: f64,
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize, Default)]
pub struct ProviderPinvOptions {
    pub tolerance: Option<f64>,
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize)]
pub struct ProviderPolyvalMu {
    pub mean: f64,
    pub scale: f64,
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize, Default)]
pub struct ProviderPolyvalOptions {
    pub mu: Option<ProviderPolyvalMu>,
}

#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderInvOptions {}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderPolyfitResult {
    pub coefficients: Vec<f64>,
    pub r_matrix: Vec<f64>,
    pub normr: f64,
    pub df: f64,
    pub mu: [f64; 2],
}

/// Numerator/denominator payload returned by provider-backed `polyder` quotient rule.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderPolyderQuotient {
    pub numerator: GpuTensorHandle,
    pub denominator: GpuTensorHandle,
}

/// Supported norm specifications for the `cond` builtin.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderCondNorm {
    Two,
    One,
    Inf,
    Fro,
}

/// Supported norm orders for the `norm` builtin.
#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize)]
pub enum ProviderNormOrder {
    Two,
    One,
    Inf,
    NegInf,
    Zero,
    Fro,
    Nuc,
    P(f64),
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderInterp1Method {
    Linear,
    Nearest,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderInterp1Extrapolation {
    Nan,
    Extrapolate,
    Value,
}

#[derive(Debug, Clone)]
pub struct ProviderInterp1Request<'a> {
    /// Strictly increasing, finite sample coordinates validated by the runtime
    /// before dispatch. Providers may assume this monotonic domain invariant.
    pub x: &'a GpuTensorHandle,
    pub y: &'a GpuTensorHandle,
    pub xq: &'a GpuTensorHandle,
    pub sample_len: usize,
    pub series_count: usize,
    pub query_len: usize,
    pub output_shape: &'a [usize],
    pub method: ProviderInterp1Method,
    pub extrapolation: ProviderInterp1Extrapolation,
    pub extrapolation_value: f64,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderEigResult {
    pub eigenvalues: GpuTensorHandle,
    pub diagonal: GpuTensorHandle,
    pub right: GpuTensorHandle,
    pub left: Option<GpuTensorHandle>,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderQrPivot {
    Matrix,
    Vector,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderQrOptions {
    pub economy: bool,
    pub pivot: ProviderQrPivot,
}

impl Default for ProviderQrOptions {
    fn default() -> Self {
        Self {
            economy: false,
            pivot: ProviderQrPivot::Matrix,
        }
    }
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderPrecision {
    F32,
    F64,
}

/// Declares how provider-owned GPU handles may cross async spawn boundaries.
///
/// This is a runtime/provider policy surface (not a semantic type fact) used by
/// VM/runtime spawn handling to prevent unsynchronized device-handle races.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum SpawnHandleConcurrency {
    /// Provider supports immutable sharing of handle-backed values across spawned tasks.
    ImmutableShare,
    /// Provider supports copy-on-write semantics when spawned and parent tasks diverge.
    CopyOnWrite,
    /// Provider supports synchronized mutation for shared handles.
    SynchronizedMutation,
    /// Provider rejects spawned sharing of raw handles.
    Reject,
}

impl SpawnHandleConcurrency {
    pub fn as_str(self) -> &'static str {
        match self {
            SpawnHandleConcurrency::ImmutableShare => "immutable_share",
            SpawnHandleConcurrency::CopyOnWrite => "copy_on_write",
            SpawnHandleConcurrency::SynchronizedMutation => "synchronized_mutation",
            SpawnHandleConcurrency::Reject => "reject",
        }
    }
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ReductionTwoPassMode {
    Auto,
    ForceOn,
    ForceOff,
}

impl ReductionTwoPassMode {
    pub fn as_str(self) -> &'static str {
        match self {
            ReductionTwoPassMode::Auto => "auto",
            ReductionTwoPassMode::ForceOn => "force_on",
            ReductionTwoPassMode::ForceOff => "force_off",
        }
    }
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize)]
pub enum ReductionFlavor {
    Sum,
    Mean,
    CustomScale(f64),
}

impl ReductionFlavor {
    pub fn is_mean(self) -> bool {
        matches!(self, ReductionFlavor::Mean)
    }

    pub fn scale(self, reduce_len: usize) -> f64 {
        match self {
            ReductionFlavor::Sum => 1.0,
            ReductionFlavor::Mean => {
                if reduce_len == 0 {
                    1.0
                } else {
                    1.0 / reduce_len as f64
                }
            }
            ReductionFlavor::CustomScale(scale) => scale,
        }
    }
}

/// Normalisation mode for correlation coefficients.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum CorrcoefNormalization {
    Unbiased,
    Biased,
}

/// Row-selection strategy for correlation coefficients.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum CorrcoefRows {
    All,
    Complete,
    Pairwise,
}

/// Options controlling provider-backed correlation coefficient computation.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub struct CorrcoefOptions {
    pub normalization: CorrcoefNormalization,
    pub rows: CorrcoefRows,
}

impl Default for CorrcoefOptions {
    fn default() -> Self {
        Self {
            normalization: CorrcoefNormalization::Unbiased,
            rows: CorrcoefRows::All,
        }
    }
}

/// Normalisation mode used by covariance computations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum CovNormalization {
    Unbiased,
    Biased,
}

/// Row handling strategy for covariance computations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum CovRows {
    All,
    OmitRows,
    PartialRows,
}

/// Options controlling provider-backed covariance computation.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub struct CovarianceOptions {
    pub normalization: CovNormalization,
    pub rows: CovRows,
    pub has_weight_vector: bool,
}

impl Default for CovarianceOptions {
    fn default() -> Self {
        Self {
            normalization: CovNormalization::Unbiased,
            rows: CovRows::All,
            has_weight_vector: false,
        }
    }
}

/// Normalization strategy used by provider-backed standard deviation reductions.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderStdNormalization {
    Sample,
    Population,
}

/// NaN handling mode for provider-backed reductions.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderNanMode {
    Include,
    Omit,
}

/// Moving-window reduction operation executed by acceleration providers.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderMovingWindowOp {
    Sum,
    Mean,
    Prod,
    Min,
    Max,
    Median,
    Std,
    Var,
}

/// Endpoint handling for provider-backed moving-window reductions.
#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize)]
pub enum ProviderMovingWindowEndpoints {
    Shrink,
    Discard,
    Fill(f64),
}

/// Request for provider-backed count-window moving reductions.
#[derive(Debug, Clone, Copy)]
pub struct ProviderMovingWindowRequest<'a> {
    pub input: &'a GpuTensorHandle,
    pub output_shape: &'a [usize],
    /// Zero-based dimension along which the moving window is applied.
    pub dim: usize,
    pub before: usize,
    pub after: usize,
    pub op: ProviderMovingWindowOp,
    pub endpoints: ProviderMovingWindowEndpoints,
    pub nan_mode: ProviderNanMode,
    pub normalization: ProviderStdNormalization,
}

/// Dimension selection for provider-backed `mode` reductions.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderModeAxes {
    /// First non-singleton dimension selected by MATLAB default rules.
    Default,
    /// Zero-based dimension supplied by `mode(A, dim)`.
    Dim(usize),
    /// Collapse every element, as in `mode(A, "all")`.
    All,
}

/// Request for provider-backed modal reductions.
#[derive(Debug, Clone, Copy)]
pub struct ProviderModeRequest<'a> {
    pub input: &'a GpuTensorHandle,
    pub axes: ProviderModeAxes,
    pub want_frequency: bool,
    pub want_ties: bool,
}

/// Sorted tied modal values for each output slice.
///
/// `values` stores all tied values for all slices in ascending per-slice order.
/// `offsets` and `counts` are one entry per output slice and describe the
/// slice-local range inside `values`. Runtime callers convert this ragged
/// representation into MATLAB cell output `C`.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderModeTiedSets {
    pub values: HostTensorOwned,
    pub offsets: Vec<usize>,
    pub counts: Vec<usize>,
}

/// Result of a provider-backed `mode` reduction.
///
/// `values` is the MATLAB `M` output and stays resident. `frequencies`, when
/// requested, is the MATLAB `F` output and stays resident. `ties`, when
/// requested, carries the MATLAB `C` tied-set data in a host ragged form because
/// the public runtime value is a cell array.
#[derive(Debug, Clone, PartialEq)]
pub struct ProviderModeResult {
    pub values: GpuTensorHandle,
    pub frequencies: Option<GpuTensorHandle>,
    pub ties: Option<ProviderModeTiedSets>,
}

/// Direction used when computing prefix sums on the device.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderScanDirection {
    Forward,
    Reverse,
}

/// Spacing input for provider-backed trapezoidal integration.
#[derive(Debug, Clone, Copy)]
pub enum ProviderTrapezoidSpacing<'a> {
    Unit,
    Scalar(f64),
    ScalarHandle(&'a GpuTensorHandle),
    Vector(&'a GpuTensorHandle),
    Tensor(&'a GpuTensorHandle),
}

/// Sort direction used by acceleration providers.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum SortOrder {
    Ascend,
    Descend,
}

/// Comparison strategy applied during sorting.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum SortComparison {
    Auto,
    Real,
    Abs,
}

/// Host-resident outputs returned by provider-backed sort operations.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct SortResult {
    pub values: HostTensorOwned,
    pub indices: HostTensorOwned,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct SortRowsColumnSpec {
    pub index: usize,
    pub order: SortOrder,
}

/// Ordering applied by provider-backed `unique` operations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum UniqueOrder {
    Sorted,
    Stable,
}

/// Occurrence selection for provider-backed `unique` operations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum UniqueOccurrence {
    First,
    Last,
}

/// Options controlling provider-backed `unique` operations.
#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct UniqueOptions {
    pub rows: bool,
    pub order: UniqueOrder,
    pub occurrence: UniqueOccurrence,
    pub treat_missing_as_distinct: bool,
    pub explicit_order: bool,
    pub explicit_occurrence: bool,
}

/// Host-resident outputs returned by provider-backed `unique` operations.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct UniqueResult {
    pub values: HostTensorOwned,
    pub ia: HostTensorOwned,
    pub ic: HostTensorOwned,
}

/// Ordering applied by provider-backed `union` operations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum UnionOrder {
    Sorted,
    Stable,
}

/// Options controlling provider-backed `union` operations.
#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct UnionOptions {
    pub rows: bool,
    pub order: UnionOrder,
}

/// Host-resident outputs returned by provider-backed `union` operations.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct UnionResult {
    pub values: HostTensorOwned,
    pub ia: HostTensorOwned,
    pub ib: HostTensorOwned,
}

/// Parameterisation of 2-D filters generated by `fspecial`.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub enum FspecialFilter {
    Average {
        rows: u32,
        cols: u32,
    },
    Disk {
        radius: f64,
        size: u32,
    },
    Gaussian {
        rows: u32,
        cols: u32,
        sigma: f64,
    },
    Laplacian {
        alpha: f64,
    },
    Log {
        rows: u32,
        cols: u32,
        sigma: f64,
    },
    Motion {
        length: u32,
        kernel_size: u32,
        angle_degrees: f64,
        oversample: u32,
    },
    Prewitt,
    Sobel,
    Unsharp {
        alpha: f64,
    },
}

/// Request dispatched to acceleration providers for `fspecial` kernels.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct FspecialRequest {
    pub filter: FspecialFilter,
}

/// Padding strategy used by `imfilter`.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ImfilterPadding {
    Constant,
    Replicate,
    Symmetric,
    Circular,
}

/// Output sizing mode used by `imfilter`.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ImfilterShape {
    Same,
    Full,
    Valid,
}

/// Correlation vs convolution behaviour for `imfilter`.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ImfilterMode {
    Correlation,
    Convolution,
}

/// Options supplied to acceleration providers for `imfilter`.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ImfilterOptions {
    pub padding: ImfilterPadding,
    pub constant_value: f64,
    pub shape: ImfilterShape,
    pub mode: ImfilterMode,
}

impl Default for ImfilterOptions {
    fn default() -> Self {
        Self {
            padding: ImfilterPadding::Constant,
            constant_value: 0.0,
            shape: ImfilterShape::Same,
            mode: ImfilterMode::Correlation,
        }
    }
}

/// Ordering applied by provider-backed `setdiff` operations.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum SetdiffOrder {
    Sorted,
    Stable,
}

/// Options controlling provider-backed `setdiff` operations.
#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct SetdiffOptions {
    pub rows: bool,
    pub order: SetdiffOrder,
}

/// Host-resident outputs returned by provider-backed `setdiff` operations.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct SetdiffResult {
    pub values: HostTensorOwned,
    pub ia: HostTensorOwned,
}

/// Options controlling provider-backed `ismember` operations.
#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct IsMemberOptions {
    pub rows: bool,
}

/// Host-resident logical output returned by providers.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct HostLogicalOwned {
    pub data: Vec<u8>,
    pub shape: Vec<usize>,
}

/// Host-resident outputs returned by provider-backed `ismember` operations.
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct IsMemberResult {
    pub mask: HostLogicalOwned,
    pub loc: HostTensorOwned,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderConvMode {
    Full,
    Same,
    Valid,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ProviderConvOrientation {
    Row,
    Column,
}

#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderConv1dOptions {
    pub mode: ProviderConvMode,
    pub orientation: ProviderConvOrientation,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderIirFilterOptions {
    /// Zero-based dimension along which filtering should be applied.
    pub dim: usize,
    /// Optional initial conditions (state vector) residing on the device.
    pub zi: Option<GpuTensorHandle>,
    /// Caller has already validated that `a` is the scalar denominator `[1]`.
    ///
    /// Providers may use this FIR-only path to avoid reading coefficient buffers
    /// back to the host for normalization.
    pub unit_denominator: bool,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderIirFilterResult {
    /// Filtered output tensor, matching the input signal shape.
    pub output: GpuTensorHandle,
    /// Final conditions for the filter state (same shape as the requested `zi` layout).
    pub final_state: Option<GpuTensorHandle>,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ProviderMoments2 {
    pub mean: GpuTensorHandle,
    pub ex2: GpuTensorHandle,
}

#[derive(Debug, Clone, Copy, Default, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderDispatchStats {
    /// Number of GPU dispatches recorded for this category.
    pub count: u64,
    /// Accumulated wall-clock time of dispatches in nanoseconds (host measured).
    pub total_wall_time_ns: u64,
}

#[derive(Debug, Clone, Default, PartialEq, Eq, Serialize, Deserialize)]
pub struct ProviderFallbackStat {
    pub reason: String,
    pub count: u64,
}

#[derive(Debug, Clone, Default, PartialEq, Serialize, Deserialize)]
pub struct ProviderTelemetry {
    pub fused_elementwise: ProviderDispatchStats,
    pub fused_reduction: ProviderDispatchStats,
    pub matmul: ProviderDispatchStats,
    pub linsolve: ProviderDispatchStats,
    pub mldivide: ProviderDispatchStats,
    pub mrdivide: ProviderDispatchStats,
    pub upload_bytes: u64,
    pub download_bytes: u64,
    pub solve_fallbacks: Vec<ProviderFallbackStat>,
    pub fusion_cache_hits: u64,
    pub fusion_cache_misses: u64,
    pub bind_group_cache_hits: u64,
    pub bind_group_cache_misses: u64,
    /// Optional per-layout bind group cache counters (layout tags and their hit/miss counts)
    pub bind_group_cache_by_layout: Option<Vec<BindGroupLayoutTelemetry>>,
    /// Recent kernel launch metadata (bounded log; newest last)
    pub kernel_launches: Vec<KernelLaunchTelemetry>,
}

#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct BindGroupLayoutTelemetry {
    pub tag: String,
    pub hits: u64,
    pub misses: u64,
}

#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
pub struct KernelAttrTelemetry {
    pub key: String,
    pub value: u64,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct KernelLaunchTelemetry {
    pub kernel: String,
    pub precision: Option<String>,
    pub shape: Vec<KernelAttrTelemetry>,
    pub tuning: Vec<KernelAttrTelemetry>,
}

pub type AccelProviderFuture<'a, T> = Pin<Box<dyn Future<Output = anyhow::Result<T>> + 'a>>;
pub type AccelDownloadFuture<'a> = AccelProviderFuture<'a, crate::HostTensorOwned>;
pub type AccelIntegerDownloadFuture<'a> = AccelProviderFuture<'a, crate::HostIntegerTensorOwned>;
pub type AccelNumericDownloadFuture<'a> = AccelProviderFuture<'a, crate::HostNumericTensorOwned>;

fn unsupported_future<T>(message: &'static str) -> AccelProviderFuture<'static, T> {
    Box::pin(async move { Err(anyhow::anyhow!(message)) })
}

/// Device/provider interface that backends implement and register into the runtime layer
pub trait AccelProvider: Send + Sync {
    fn upload(&self, host: &crate::HostTensorView) -> anyhow::Result<GpuTensorHandle>;
    fn download<'a>(&'a self, h: &'a GpuTensorHandle) -> AccelDownloadFuture<'a>;

    /// Upload one native numeric payload through the shared all-class transfer
    /// contract. Providers should override this method as they adopt native
    /// single and complex integer storage. The default adapter preserves exact
    /// existing double and real-integer routes and rejects every representation
    /// those routes cannot express.
    fn upload_numeric(
        &self,
        host: &crate::HostNumericTensorView,
    ) -> anyhow::Result<GpuTensorHandle> {
        host.validate()?;
        let handle = match (host.data, host.storage) {
            (crate::HostNumericDataView::F64(data), GpuTensorStorage::Real) => {
                self.upload(&crate::HostTensorView {
                    data,
                    shape: host.shape,
                })
            }
            (data, GpuTensorStorage::Real) if data.element_type().integer_type().is_some() => {
                let data = match data {
                    crate::HostNumericDataView::I8(data) => crate::HostIntegerDataView::I8(data),
                    crate::HostNumericDataView::I16(data) => crate::HostIntegerDataView::I16(data),
                    crate::HostNumericDataView::I32(data) => crate::HostIntegerDataView::I32(data),
                    crate::HostNumericDataView::I64(data) => crate::HostIntegerDataView::I64(data),
                    crate::HostNumericDataView::U8(data) => crate::HostIntegerDataView::U8(data),
                    crate::HostNumericDataView::U16(data) => crate::HostIntegerDataView::U16(data),
                    crate::HostNumericDataView::U32(data) => crate::HostIntegerDataView::U32(data),
                    crate::HostNumericDataView::U64(data) => crate::HostIntegerDataView::U64(data),
                    crate::HostNumericDataView::F64(_) | crate::HostNumericDataView::F32(_) => {
                        unreachable!("guarded integer numeric transfer")
                    }
                };
                self.upload_integer(&crate::HostIntegerTensorView {
                    data,
                    shape: host.shape,
                })
            }
            (data, storage) => Err(anyhow!(
                "provider must implement native {:?} {:?} uploads through upload_numeric",
                data.element_type(),
                storage
            )),
        }?;
        let expected_element = host.data.element_type();
        if handle.shape != host.shape
            || handle.descriptor.element_type != Some(expected_element)
            || handle.descriptor.storage != Some(host.storage)
        {
            let _ = self.free(&handle);
            return Err(anyhow!(
                "provider returned an invalid numeric upload descriptor: expected shape {:?}, element {:?}, storage {:?}; got shape {:?}, element {:?}, storage {:?}",
                host.shape,
                expected_element,
                host.storage,
                handle.shape,
                handle.descriptor.element_type,
                handle.descriptor.storage
            ));
        }
        Ok(handle)
    }

    /// Download one native numeric payload through the shared all-class
    /// transfer contract. The default adapter keeps exact real-integer and
    /// native-double downloads available while refusing to present widened
    /// single data as native `f32`.
    fn download_numeric<'a>(&'a self, h: &'a GpuTensorHandle) -> AccelNumericDownloadFuture<'a> {
        Box::pin(async move {
            let element_type = h.descriptor.element_type.ok_or_else(|| {
                anyhow!("numeric handle is missing its durable element descriptor")
            })?;
            let storage = h.descriptor.storage.ok_or_else(|| {
                anyhow!("numeric handle is missing its durable storage descriptor")
            })?;
            if element_type.integer_type().is_some() {
                if storage != GpuTensorStorage::Real {
                    return Err(anyhow!(
                        "provider must implement native complex integer downloads through download_numeric"
                    ));
                }
                let downloaded = self.download_integer(h).await?;
                let numeric: crate::HostNumericTensorOwned = downloaded.into();
                numeric.validate()?;
                return Ok(numeric);
            }
            if element_type == NumericElementType::F32 {
                return Err(anyhow!(
                    "provider must implement native single downloads through download_numeric"
                ));
            }
            let downloaded = self.download(h).await?;
            let numeric = crate::HostNumericTensorOwned {
                data: crate::HostNumericDataOwned::F64(downloaded.data),
                shape: downloaded.shape,
                storage: downloaded.storage,
            };
            numeric.validate()?;
            Ok(numeric)
        })
    }

    /// Upload exact native integer storage without converting through f64.
    /// Providers that do not expose native integer buffers must reject this
    /// operation rather than silently changing the value class or precision.
    fn upload_integer(
        &self,
        _host: &crate::HostIntegerTensorView,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "provider does not support exact native integer buffers"
        ))
    }

    /// Download exact native integer storage without converting through f64.
    fn download_integer<'a>(&'a self, _h: &'a GpuTensorHandle) -> AccelIntegerDownloadFuture<'a> {
        Box::pin(async {
            Err(anyhow::anyhow!(
                "provider does not support exact native integer buffers"
            ))
        })
    }

    fn free(&self, h: &GpuTensorHandle) -> anyhow::Result<()>;
    fn device_info(&self) -> String;
    /// Returns the stable identifier used to route this provider's handles.
    /// Distinct concurrently registered provider instances must return distinct
    /// identifiers; reusing an identifier would make handle ownership ambiguous.
    fn device_id(&self) -> u32 {
        0
    }

    /// Declares provider policy for sharing `GpuTensorHandle` values across
    /// spawned async boundaries.
    ///
    /// Default is conservative rejection. Providers that can safely support
    /// cross-task sharing should override this.
    fn spawn_handle_concurrency(&self) -> SpawnHandleConcurrency {
        SpawnHandleConcurrency::Reject
    }

    /// Returns a versioned, side-effect-free description of the provider
    /// surfaces placement may consider. The default advertises only the
    /// mandatory floating-point transfer boundary.
    fn capability_snapshot(&self) -> ProviderCapabilitySnapshot {
        ProviderCapabilitySnapshot::conservative(self)
    }

    /// Determines whether a representation-specific operation can execute
    /// without attempting allocation, compilation, transfer, or dispatch.
    fn query_feasibility(&self, query: &ProviderFeasibilityQuery) -> ProviderFeasibility {
        query.conservative_transfer_feasibility(self.precision())
    }

    /// Returns a side-effect-free component cost estimate for an already
    /// feasible operation. Providers may return `None` until they have a
    /// trustworthy prior or observation; placement then supplies a bounded
    /// low-confidence policy prior rather than probing by execution.
    fn estimate_cost(&self, _query: &ProviderCostQuery) -> Option<ProviderCostEstimate> {
        None
    }

    /// Returns a side-effect-free resource snapshot for placement admission.
    /// Providers with allocator/queue telemetry should override this default;
    /// the conservative snapshot exposes declared capacity when available,
    /// leaves unknown queue state explicit, and never probes the device.
    fn placement_resources(&self) -> runmat_execution::ProviderResourceSnapshot {
        let info = self.device_info_struct();
        let capacity = info.memory_bytes;
        runmat_execution::ProviderResourceSnapshot {
            device_id: self.device_id(),
            capacity_bytes: capacity,
            live_bytes: 0,
            reclaimable_bytes: 0,
            scratch_available_bytes: capacity,
            queue_depth: None,
            queue_limit: None,
            lost: false,
            epoch: 0,
        }
    }

    /// Export a shared GPU context handle, allowing downstream systems (plotting, visualization)
    /// to reuse the same device/queue without copying tensor data back to the host.
    fn export_context(&self, _kind: AccelContextKind) -> Option<AccelContextHandle> {
        None
    }

    /// Export a provider-owned WGPU buffer for zero-copy integrations.
    #[cfg(feature = "wgpu")]
    fn export_wgpu_buffer(&self, _handle: &GpuTensorHandle) -> Option<WgpuBufferRef> {
        let _ = _handle;
        None
    }

    /// Gather logical elements from `source` at the provided zero-based linear `indices`,
    /// materialising a dense tensor with the specified `output_shape`.
    ///
    /// Indices address MATLAB logical elements, not raw provider storage lanes. Providers that
    /// store complex values as interleaved lanes must copy both lanes for each selected element
    /// and preserve the source storage kind on the output handle.
    fn gather_linear(
        &self,
        _source: &GpuTensorHandle,
        _indices: &[u32],
        _output_shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("gather_linear not supported by provider"))
    }

    /// Scatter logical elements from `values` into `target` at the provided zero-based linear
    /// `indices`.
    ///
    /// Indices address MATLAB logical elements, not raw provider storage lanes. Providers must
    /// ensure `values` has one logical element per index, copy all lanes for the handle storage
    /// kind, and update `target` in place without changing its shape or storage.
    fn scatter_linear(
        &self,
        _target: &GpuTensorHandle,
        _indices: &[u32],
        _values: &GpuTensorHandle,
    ) -> anyhow::Result<()> {
        Err(anyhow::anyhow!("scatter_linear not supported by provider"))
    }

    /// Structured device information (optional to override). Default adapts from `device_info()`.
    fn device_info_struct(&self) -> ApiDeviceInfo {
        ApiDeviceInfo {
            device_id: 0,
            name: self.device_info(),
            vendor: String::new(),
            memory_bytes: None,
            backend: None,
        }
    }

    fn precision(&self) -> ProviderPrecision {
        ProviderPrecision::F64
    }

    /// Read a single scalar at linear index from a device tensor, returning it as f64.
    fn read_scalar(&self, _h: &GpuTensorHandle, _linear_index: usize) -> anyhow::Result<f64> {
        Err(anyhow::anyhow!("read_scalar not supported by provider"))
    }

    /// Allocate a zero-initialised tensor with the provided shape on the device.
    fn zeros(&self, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("zeros not supported by provider"))
    }

    /// Allocate a zero-initialised tensor with the provided shape and storage layout.
    ///
    /// `shape` is the logical MATLAB shape. Providers must allocate enough raw lanes for the
    /// requested storage kind, e.g. two interleaved numeric lanes per logical element for complex
    /// tensors.
    fn zeros_with_storage(
        &self,
        shape: &[usize],
        storage: GpuTensorStorage,
    ) -> anyhow::Result<GpuTensorHandle> {
        if storage == GpuTensorStorage::Real {
            self.zeros(shape)
        } else {
            Err(anyhow::anyhow!(
                "zeros_with_storage not supported by provider for {storage:?}"
            ))
        }
    }

    /// Allocate a one-initialised tensor with the provided shape on the device.
    fn ones(&self, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("ones not supported by provider"))
    }

    /// Allocate a zero-initialised tensor matching the prototype tensor.
    fn zeros_like(&self, prototype: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        self.zeros(&prototype.shape)
    }

    /// Allocate an exact native-integer zero buffer using the element class of
    /// `prototype`. Implementations must reject non-native-integer prototypes
    /// rather than allocating a floating compatibility buffer.
    fn zeros_integer_like(
        &self,
        _prototype: &GpuTensorHandle,
        _shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "zeros_integer_like not supported by provider"
        ))
    }

    /// Allocate a tensor filled with a constant value on the device.
    fn fill(&self, shape: &[usize], value: f64) -> anyhow::Result<GpuTensorHandle> {
        if value == 0.0 {
            return self.zeros(shape);
        }
        if let Ok(base) = self.zeros(shape) {
            match self.scalar_add(&base, value) {
                Ok(out) => {
                    let _ = self.free(&base);
                    return Ok(out);
                }
                Err(_) => {
                    let _ = self.free(&base);
                }
            }
        }
        let len: usize = shape.iter().copied().product();
        let data = vec![value; len];
        let view = HostTensorView { data: &data, shape };
        self.upload(&view)
    }

    /// Allocate a tensor filled with a constant value, matching a prototype's residency.
    fn fill_like(
        &self,
        prototype: &GpuTensorHandle,
        value: f64,
    ) -> anyhow::Result<GpuTensorHandle> {
        if value == 0.0 {
            return self.zeros_like(prototype);
        }
        if let Ok(base) = self.zeros_like(prototype) {
            match self.scalar_add(&base, value) {
                Ok(out) => {
                    let _ = self.free(&base);
                    return Ok(out);
                }
                Err(_) => {
                    let _ = self.free(&base);
                }
            }
        }
        self.fill(&prototype.shape, value)
    }

    /// Allocate a one-initialised tensor matching the prototype tensor.
    fn ones_like(&self, prototype: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        self.ones(&prototype.shape)
    }

    /// Allocate an identity tensor with ones along the leading diagonal of the first two axes.
    fn eye(&self, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("eye not supported by provider"))
    }

    /// Allocate an identity tensor matching the prototype tensor's shape.
    fn eye_like(&self, prototype: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        self.eye(&prototype.shape)
    }

    /// Construct MATLAB-style coordinate grids from axis vectors.
    fn meshgrid(&self, _axes: &[MeshgridAxisView<'_>]) -> anyhow::Result<ProviderMeshgridResult> {
        Err(anyhow::anyhow!("meshgrid not supported by provider"))
    }

    /// Construct MATLAB-style N-D coordinate grids from resident GPU axis vectors.
    fn ndgrid(&self, _request: &ProviderNdgridRequest<'_>) -> anyhow::Result<ProviderNdgridResult> {
        Err(anyhow::anyhow!("ndgrid not supported by provider"))
    }

    /// Compute vectorized Black-Scholes European call and put prices on resident GPU inputs.
    fn black_scholes_price(
        &self,
        _request: &ProviderBlackScholesPriceRequest<'_>,
    ) -> anyhow::Result<ProviderBlackScholesPriceResult> {
        Err(anyhow::anyhow!(
            "black_scholes_price not supported by provider"
        ))
    }

    /// Apply the Adam optimizer update to resident parameter and state tensors.
    fn adam_update(
        &self,
        _request: &ProviderAdamUpdateRequest<'_>,
    ) -> anyhow::Result<ProviderAdamUpdateResult> {
        Err(anyhow::anyhow!("adam_update not supported by provider"))
    }

    /// Compute per-element cross-entropy loss terms for resident prediction and target tensors.
    fn crossentropy_terms(
        &self,
        _request: &ProviderCrossentropyRequest<'_>,
    ) -> anyhow::Result<ProviderCrossentropyResult> {
        Err(anyhow::anyhow!(
            "crossentropy_terms not supported by provider"
        ))
    }

    /// Construct a diagonal matrix from a vector-like tensor. `offset` matches MATLAB semantics.
    fn diag_from_vector(
        &self,
        _vector: &GpuTensorHandle,
        _offset: isize,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "diag_from_vector not supported by provider"
        ))
    }

    /// Construct a diagonal matrix with an explicit shape from a vector-like tensor.
    /// `offset` matches MATLAB semantics; values that do not fit inside the requested
    /// rectangle are ignored.
    fn diag_from_vector_sized(
        &self,
        _vector: &GpuTensorHandle,
        _offset: isize,
        _rows: usize,
        _cols: usize,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "diag_from_vector_sized not supported by provider"
        ))
    }

    /// Extract a diagonal from a matrix-like tensor. The result is always a column vector.
    fn diag_extract(
        &self,
        _matrix: &GpuTensorHandle,
        _offset: isize,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("diag_extract not supported by provider"))
    }

    /// Apply a lower-triangular mask to the first two dimensions of a tensor.
    fn tril<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _offset: isize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow!("tril not supported by provider")) })
    }

    /// Apply an upper-triangular mask to the first two dimensions of a tensor.
    fn triu<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _offset: isize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow!("triu not supported by provider")) })
    }

    /// Evaluate a polynomial expressed by `coefficients` at each element in `points`.
    fn polyval(
        &self,
        _coefficients: &GpuTensorHandle,
        _points: &GpuTensorHandle,
        _options: &ProviderPolyvalOptions,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("polyval not supported by provider"))
    }

    /// Fit a polynomial of degree `degree` to `(x, y)` samples. Optional weights must match `x`.
    fn polyfit<'a>(
        &'a self,
        _x: &'a GpuTensorHandle,
        _y: &'a GpuTensorHandle,
        _degree: usize,
        _weights: Option<&'a GpuTensorHandle>,
    ) -> AccelProviderFuture<'a, ProviderPolyfitResult> {
        Box::pin(async move { Err(anyhow::anyhow!("polyfit not supported by provider")) })
    }

    /// Differentiate a polynomial represented as a vector of coefficients.
    fn polyder_single<'a>(
        &'a self,
        _polynomial: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("polyder_single not supported by provider")) })
    }

    /// Apply the product rule to polynomials `p` and `q`.
    fn polyder_product<'a>(
        &'a self,
        _p: &'a GpuTensorHandle,
        _q: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("polyder_product not supported by provider")) })
    }

    /// Apply the quotient rule to polynomials `u` and `v`.
    fn polyder_quotient<'a>(
        &'a self,
        _u: &'a GpuTensorHandle,
        _v: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, ProviderPolyderQuotient> {
        Box::pin(async move {
            Err(anyhow::anyhow!(
                "polyder_quotient not supported by provider"
            ))
        })
    }

    /// Integrate a polynomial represented as a vector of coefficients and append a constant term.
    fn polyint(
        &self,
        _polynomial: &GpuTensorHandle,
        _constant: f64,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("polyint not supported by provider"))
    }

    /// Allocate a tensor filled with random values drawn from U(0, 1).
    fn random_uniform(&self, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("random_uniform not supported by provider"))
    }

    /// Allocate a tensor filled with random values matching the prototype shape.
    fn random_uniform_like(&self, prototype: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        self.random_uniform(&prototype.shape)
    }

    /// Allocate a tensor filled with standard normal (mean 0, stddev 1) random values.
    fn random_normal(&self, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("random_normal not supported by provider"))
    }

    /// Allocate a tensor of standard normal values matching a prototype's shape.
    fn random_normal_like(&self, prototype: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        self.random_normal(&prototype.shape)
    }

    /// Exponentially-distributed random values with mean `mu`.
    fn random_exponential(&self, _mu: f64, _shape: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "random_exponential not supported by provider"
        ))
    }

    /// Normal random values with mean `mu` and standard deviation `sigma`.
    fn random_normrnd(
        &self,
        _mu: f64,
        _sigma: f64,
        _shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("random_normrnd not supported by provider"))
    }

    /// Uniform random values on the interval `[a, b)`.
    fn random_unifrnd(
        &self,
        _a: f64,
        _b: f64,
        _shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("random_unifrnd not supported by provider"))
    }

    fn stochastic_evolution(
        &self,
        _state: &GpuTensorHandle,
        _drift: f64,
        _scale: f64,
        _steps: u32,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "stochastic_evolution not supported by provider"
        ))
    }

    /// Set the provider RNG state to align with the host RNG.
    fn set_rng_state(&self, _state: u64) -> anyhow::Result<()> {
        Err(anyhow::anyhow!("set_rng_state not supported by provider"))
    }

    /// Generate a 2-D correlation kernel matching MATLAB's `fspecial` builtin.
    fn fspecial(&self, _request: &FspecialRequest) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("fspecial not supported by provider"))
    }

    /// Evaluate the `peaks` test surface on an n×n grid spanning [-3,3]×[-3,3].
    /// Returns the Z matrix (n×n) as a GPU tensor.
    fn peaks(&self, _n: usize) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("peaks not supported by provider"))
    }

    /// Evaluate the `peaks` formula element-wise on caller-supplied GPU coordinate tensors.
    /// X and Y must have the same shape. Returns a Z tensor of the same shape.
    fn peaks_xy(
        &self,
        _x: &GpuTensorHandle,
        _y: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("peaks_xy not supported by provider"))
    }

    fn hann_window(&self, _len: usize, _periodic: bool) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("hann_window not supported by provider"))
    }

    fn hamming_window(&self, _len: usize, _periodic: bool) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("hamming_window not supported by provider"))
    }

    fn blackman_window(&self, _len: usize, _periodic: bool) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("blackman_window not supported by provider"))
    }

    /// Apply an N-D correlation/convolution with padding semantics matching MATLAB's `imfilter`.
    fn imfilter<'a>(
        &'a self,
        _image: &'a GpuTensorHandle,
        _kernel: &'a GpuTensorHandle,
        _options: &'a ImfilterOptions,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("imfilter not supported by provider")
    }

    /// Allocate a random permutation of 1..=n, returning the first k elements.
    fn random_permutation(&self, _n: usize, _k: usize) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow!("random_permutation not supported by provider"))
    }

    /// Allocate a random permutation matching the prototype residency.
    fn random_permutation_like(
        &self,
        _prototype: &GpuTensorHandle,
        n: usize,
        k: usize,
    ) -> anyhow::Result<GpuTensorHandle> {
        self.random_permutation(n, k)
    }

    /// Compute a covariance matrix across the columns of `matrix`.
    fn covariance<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _second: Option<&'a GpuTensorHandle>,
        _weights: Option<&'a GpuTensorHandle>,
        _options: &'a CovarianceOptions,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("covariance not supported by provider")
    }

    /// Compute a correlation coefficient matrix across the columns of `matrix`.
    fn corrcoef<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _options: &'a CorrcoefOptions,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("corrcoef not supported by provider")
    }

    /// Convert a resident covariance matrix into correlation and sigma outputs.
    fn covariance_to_correlation(
        &self,
        _matrix: &GpuTensorHandle,
    ) -> anyhow::Result<ProviderCovarianceToCorrelationResult> {
        Err(anyhow::anyhow!(
            "covariance_to_correlation not supported by provider"
        ))
    }

    // Optional operator hooks (default to unsupported)
    fn linspace(&self, _start: f64, _stop: f64, _count: usize) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("linspace not supported by provider"))
    }
    fn elem_add<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_add not supported by provider")
    }
    fn elem_mul<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_mul not supported by provider")
    }
    fn elem_max<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_max not supported by provider")
    }
    fn elem_min<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_min not supported by provider")
    }
    fn elem_sub<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_sub not supported by provider")
    }
    fn elem_div<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_div not supported by provider")
    }
    /// Compute an exact MATLAB-style remainder for compatible native integer tensors.
    fn elem_rem<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_rem not supported by provider")
    }
    /// Compute an exact MATLAB-style modulus for compatible native integer tensors.
    fn elem_mod<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_mod not supported by provider")
    }
    fn elem_pow<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_pow not supported by provider")
    }

    /// Construct complex-interleaved GPU storage from a real-valued tensor,
    /// using zero for the imaginary lane.
    fn complex_from_real<'a>(
        &'a self,
        _real: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("complex_from_real not supported by provider")
    }

    /// Construct complex-interleaved GPU storage from real and imaginary tensors.
    ///
    /// Implementations should support equal shapes and scalar expansion for either
    /// operand, matching MATLAB's `complex(real, imag)` size rules.
    fn complex_from_real_imag<'a>(
        &'a self,
        _real: &'a GpuTensorHandle,
        _imag: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("complex_from_real_imag not supported by provider")
    }

    /// Map a resident real-valued symbol tensor through a complex constellation
    /// table and return complex-interleaved GPU storage.
    fn modulate_constellation<'a>(
        &'a self,
        _request: ProviderModulationRequest<'a>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("modulate_constellation not supported by provider")
    }

    /// Group a resident real/logical bit tensor into symbols, map through a complex
    /// constellation table, and return complex-interleaved GPU storage.
    fn modulate_bits_constellation<'a>(
        &'a self,
        _request: ProviderBitModulationRequest<'a>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("modulate_bits_constellation not supported by provider")
    }

    fn elem_hypot<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_hypot not supported by provider")
    }
    fn elem_ge<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_ge not supported by provider")
    }
    fn elem_le<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_le not supported by provider")
    }
    fn elem_lt<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_lt not supported by provider")
    }
    fn elem_gt<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_gt not supported by provider")
    }
    fn elem_eq<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_eq not supported by provider")
    }
    fn elem_ne<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_ne not supported by provider")
    }
    fn logical_and(
        &self,
        _a: &GpuTensorHandle,
        _b: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_and not supported by provider"))
    }
    fn logical_or(
        &self,
        _a: &GpuTensorHandle,
        _b: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_or not supported by provider"))
    }
    fn logical_xor(
        &self,
        _a: &GpuTensorHandle,
        _b: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_xor not supported by provider"))
    }
    fn logical_not(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_not not supported by provider"))
    }
    fn logical_islogical(&self, a: &GpuTensorHandle) -> anyhow::Result<bool> {
        Ok(handle_is_logical(a))
    }
    fn logical_isreal(&self, _a: &GpuTensorHandle) -> anyhow::Result<bool> {
        Err(anyhow::anyhow!("logical_isreal not supported by provider"))
    }
    fn logical_isfinite(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "logical_isfinite not supported by provider"
        ))
    }
    fn logical_isnan(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_isnan not supported by provider"))
    }
    fn logical_isinf(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("logical_isinf not supported by provider"))
    }
    fn elem_atan2<'a>(
        &'a self,
        _y: &'a GpuTensorHandle,
        _x: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("elem_atan2 not supported by provider")
    }
    // Unary elementwise operations (optional)
    fn unary_sin<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_sin not supported by provider")
    }
    fn unary_sinc<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_sinc not supported by provider")
    }
    fn unary_gamma<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_gamma not supported by provider")
    }
    fn unary_gammaln<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_gammaln not supported by provider")
    }
    fn unary_erf<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_erf not supported by provider")
    }
    fn unary_erfcinv<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_erfcinv not supported by provider")
    }
    fn unary_factorial<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_factorial not supported by provider")
    }
    fn unary_asinh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_asinh not supported by provider")
    }
    fn unary_sinh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_sinh not supported by provider")
    }
    fn unary_cosh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_cosh not supported by provider")
    }
    fn unary_asin<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_asin not supported by provider")
    }
    fn unary_acos<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_acos not supported by provider")
    }
    fn unary_acosh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_acosh not supported by provider")
    }
    fn unary_tan<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_tan not supported by provider")
    }
    fn unary_tanh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_tanh not supported by provider")
    }
    fn unary_atan<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_atan not supported by provider")
    }
    fn unary_atanh<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_atanh not supported by provider")
    }
    fn unary_ceil<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_ceil not supported by provider")
    }
    fn unary_floor<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_floor not supported by provider")
    }
    fn unary_round<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_round not supported by provider")
    }
    fn round_digits<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _digits: i32,
        _significant: bool,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("round_digits not supported by provider")
    }
    fn unary_fix<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_fix not supported by provider")
    }
    fn unary_cos<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_cos not supported by provider")
    }
    fn unary_angle<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_angle not supported by provider")
    }
    fn unary_imag<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_imag not supported by provider")
    }
    fn unary_real<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_real not supported by provider")
    }
    fn unary_conj<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_conj not supported by provider")
    }
    fn unary_abs<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_abs not supported by provider")
    }
    fn unary_sign<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_sign not supported by provider")
    }
    fn unary_heaviside<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_heaviside not supported by provider")
    }
    fn unary_exp<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_exp not supported by provider")
    }
    /// Apply an exponential linear unit activation with the supplied alpha.
    fn activation_elu<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _alpha: f64,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("activation_elu not supported by provider")
    }
    /// Normalize each row of a real two-dimensional tensor with stable softmax.
    fn activation_softmax_rows<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("activation_softmax_rows not supported by provider")
    }
    fn unary_expm1<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_expm1 not supported by provider")
    }
    fn unary_log<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_log not supported by provider")
    }
    fn unary_log2<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_log2 not supported by provider")
    }
    fn unary_log10<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_log10 not supported by provider")
    }
    fn unary_log1p<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_log1p not supported by provider")
    }
    fn unary_sqrt<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_sqrt not supported by provider")
    }
    fn unary_double<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_double not supported by provider")
    }
    fn unary_single<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_single not supported by provider")
    }
    fn unary_pow2<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_pow2 not supported by provider")
    }
    fn unary_nextpow2<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("unary_nextpow2 not supported by provider")
    }
    fn pow2_scale(
        &self,
        _mantissa: &GpuTensorHandle,
        _exponent: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("pow2_scale not supported by provider"))
    }
    // Left-scalar operations (broadcast with scalar on the left)
    fn scalar_rsub(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_rsub not supported by provider"))
    }
    fn scalar_rdiv(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_rdiv not supported by provider"))
    }
    // Scalar operations: apply op with scalar right-hand side (broadcast over a)
    fn scalar_add(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_add not supported by provider"))
    }
    fn scalar_sub(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_sub not supported by provider"))
    }
    fn scalar_mul(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_mul not supported by provider"))
    }
    fn scalar_max(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_max not supported by provider"))
    }
    fn scalar_min(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_min not supported by provider"))
    }
    fn scalar_div(&self, _a: &GpuTensorHandle, _scalar: f64) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scalar_div not supported by provider"))
    }
    fn sort_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
        _order: SortOrder,
        _comparison: SortComparison,
    ) -> AccelProviderFuture<'a, SortResult> {
        unsupported_future("sort_dim not supported by provider")
    }
    fn sort_rows<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _columns: &'a [SortRowsColumnSpec],
        _comparison: SortComparison,
    ) -> AccelProviderFuture<'a, SortResult> {
        unsupported_future("sort_rows not supported by provider")
    }
    fn matmul<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("matmul not supported by provider")
    }

    fn syrk(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("syrk not supported by provider"))
    }
    fn pagefun(&self, _request: &PagefunRequest) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("pagefun not supported by provider"))
    }

    /// Optional: matrix multiplication with an epilogue applied before store.
    ///
    /// The default implementation falls back to `matmul` when the epilogue is effectively a no-op
    /// (alpha=1, beta=0, no row/col scales), and otherwise returns `Err`.
    fn matmul_epilogue<'a>(
        &'a self,
        a: &'a GpuTensorHandle,
        b: &'a GpuTensorHandle,
        epilogue: &'a MatmulEpilogue,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move {
            if epilogue.is_noop() {
                return self.matmul(a, b).await;
            }
            Err(anyhow::anyhow!("matmul_epilogue not supported by provider"))
        })
    }
    fn image_normalize<'a>(
        &'a self,
        _input: &'a GpuTensorHandle,
        _desc: &'a ImageNormalizeDescriptor,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("image_normalize fusion not supported by provider")
    }
    fn matmul_power_step<'a>(
        &'a self,
        _lhs: &'a GpuTensorHandle,
        _rhs: &'a GpuTensorHandle,
        _epilogue: &'a PowerStepEpilogue,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("matmul_power_step normalization not supported by provider")
    }
    fn linsolve<'a>(
        &'a self,
        _lhs: &'a GpuTensorHandle,
        _rhs: &'a GpuTensorHandle,
        _options: &'a ProviderLinsolveOptions,
    ) -> AccelProviderFuture<'a, ProviderLinsolveResult> {
        unsupported_future("linsolve not supported by provider")
    }
    fn inv<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _options: ProviderInvOptions,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("inv not supported by provider")
    }
    fn pinv<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _options: ProviderPinvOptions,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("pinv not supported by provider")
    }
    fn cond<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _norm: ProviderCondNorm,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("cond not supported by provider")) })
    }
    fn norm<'a>(
        &'a self,
        _tensor: &'a GpuTensorHandle,
        _order: ProviderNormOrder,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("norm not supported by provider")) })
    }
    fn interp1<'a>(
        &'a self,
        _request: &'a ProviderInterp1Request<'a>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("interp1 not supported by provider")
    }
    fn rank<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _tolerance: Option<f64>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("rank not supported by provider")) })
    }
    fn rcond<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("rcond not supported by provider")) })
    }
    fn mldivide<'a>(
        &'a self,
        _lhs: &'a GpuTensorHandle,
        _rhs: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("mldivide not supported by provider")) })
    }
    fn mrdivide<'a>(
        &'a self,
        _lhs: &'a GpuTensorHandle,
        _rhs: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        Box::pin(async move { Err(anyhow::anyhow!("mrdivide not supported by provider")) })
    }
    fn eig<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _compute_left: bool,
    ) -> AccelProviderFuture<'a, ProviderEigResult> {
        Box::pin(async move { Err(anyhow::anyhow!("eig not supported by provider")) })
    }
    fn lu<'a>(&'a self, _a: &'a GpuTensorHandle) -> AccelProviderFuture<'a, ProviderLuResult> {
        Box::pin(async move { Err(anyhow::anyhow!("lu not supported by provider")) })
    }

    fn chol<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _lower: bool,
    ) -> AccelProviderFuture<'a, ProviderCholResult> {
        Box::pin(async move { Err(anyhow::anyhow!("chol not supported by provider")) })
    }
    fn qr<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _options: ProviderQrOptions,
    ) -> AccelProviderFuture<'a, ProviderQrResult> {
        Box::pin(async move { Err(anyhow::anyhow!("qr not supported by provider")) })
    }
    fn take_matmul_sources(
        &self,
        _product: &GpuTensorHandle,
    ) -> Option<(GpuTensorHandle, GpuTensorHandle)> {
        None
    }
    fn qr_power_iter<'a>(
        &'a self,
        product: &'a GpuTensorHandle,
        _product_lhs: Option<&'a GpuTensorHandle>,
        q_handle: &'a GpuTensorHandle,
        options: &'a ProviderQrOptions,
    ) -> AccelProviderFuture<'a, Option<ProviderQrPowerIterResult>> {
        let _ = (product, q_handle, options);
        Box::pin(async move { Ok(None) })
    }
    fn transpose(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("transpose not supported by provider"))
    }
    fn conv1d(
        &self,
        _signal: &GpuTensorHandle,
        _kernel: &GpuTensorHandle,
        _options: ProviderConv1dOptions,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("conv1d not supported by provider"))
    }
    fn conv2d(
        &self,
        _signal: &GpuTensorHandle,
        _kernel: &GpuTensorHandle,
        _mode: ProviderConvMode,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("conv2d not supported by provider"))
    }
    fn iir_filter<'a>(
        &'a self,
        _b: &'a GpuTensorHandle,
        _a: &'a GpuTensorHandle,
        _x: &'a GpuTensorHandle,
        _options: ProviderIirFilterOptions,
    ) -> AccelProviderFuture<'a, ProviderIirFilterResult> {
        Box::pin(async move { Err(anyhow::anyhow!("iir_filter not supported by provider")) })
    }
    fn uniform_spectral_estimate<'a>(
        &'a self,
        _request: &'a ProviderSpectralRequest<'a>,
    ) -> AccelProviderFuture<'a, ProviderSpectralResult> {
        unsupported_future("uniform_spectral_estimate not supported by provider")
    }
    fn signal_envelope<'a>(
        &'a self,
        _request: &'a ProviderEnvelopeRequest<'a>,
    ) -> AccelProviderFuture<'a, ProviderEnvelopeResult> {
        unsupported_future("signal_envelope not supported by provider")
    }
    fn signal_hilbert<'a>(
        &'a self,
        _request: &'a ProviderHilbertRequest<'a>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("signal_hilbert not supported by provider")
    }
    /// Reorder tensor dimensions according to `order`, expressed as zero-based indices.
    fn permute(
        &self,
        _handle: &GpuTensorHandle,
        _order: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("permute not supported by provider"))
    }
    fn flip(&self, _handle: &GpuTensorHandle, _axes: &[usize]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("flip not supported by provider"))
    }
    fn circshift(
        &self,
        _handle: &GpuTensorHandle,
        _shifts: &[isize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("circshift not supported by provider"))
    }
    fn diff_dim(
        &self,
        _handle: &GpuTensorHandle,
        _order: usize,
        _dim: usize,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("diff_dim not supported by provider"))
    }
    fn gradient_dim(
        &self,
        _handle: &GpuTensorHandle,
        _dim: usize,
        _spacing: f64,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("gradient_dim not supported by provider"))
    }
    fn gradient_dim_with_coordinates(
        &self,
        _handle: &GpuTensorHandle,
        _dim: usize,
        _coordinates: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "gradient_dim_with_coordinates not supported by provider"
        ))
    }
    /// Perform an in-place FFT along a zero-based dimension, optionally padding/truncating to `len`.
    fn fft_dim<'a>(
        &'a self,
        _handle: &'a GpuTensorHandle,
        _len: Option<usize>,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("fft_dim not supported by provider")
    }
    fn ifft_dim<'a>(
        &'a self,
        _handle: &'a GpuTensorHandle,
        _len: Option<usize>,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("ifft_dim not supported by provider")
    }
    fn fft_extract_real<'a>(
        &'a self,
        _handle: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("fft_extract_real not supported by provider")
    }
    fn unique<'a>(
        &'a self,
        _handle: &'a GpuTensorHandle,
        _options: &'a UniqueOptions,
    ) -> AccelProviderFuture<'a, UniqueResult> {
        Box::pin(async move { Err(anyhow::anyhow!("unique not supported by provider")) })
    }
    fn union<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
        _options: &'a UnionOptions,
    ) -> AccelProviderFuture<'a, UnionResult> {
        Box::pin(async move { Err(anyhow::anyhow!("union not supported by provider")) })
    }
    fn setdiff<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
        _options: &'a SetdiffOptions,
    ) -> AccelProviderFuture<'a, SetdiffResult> {
        Box::pin(async move { Err(anyhow::anyhow!("setdiff not supported by provider")) })
    }
    fn ismember<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _b: &'a GpuTensorHandle,
        _options: &'a IsMemberOptions,
    ) -> AccelProviderFuture<'a, IsMemberResult> {
        Box::pin(async move { Err(anyhow::anyhow!("ismember not supported by provider")) })
    }
    fn reshape(
        &self,
        handle: &GpuTensorHandle,
        new_shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        let mut updated = handle.clone();
        updated.shape = new_shape.to_vec();
        Ok(updated)
    }
    /// Concatenate the provided tensors along the 1-based dimension `dim`.
    fn cat(&self, _dim: usize, _inputs: &[GpuTensorHandle]) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("cat not supported by provider"))
    }
    fn repmat(
        &self,
        _handle: &GpuTensorHandle,
        _reps: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("repmat not supported by provider"))
    }
    /// Compute the Kronecker product of two tensors, matching MATLAB semantics.
    fn kron(&self, _a: &GpuTensorHandle, _b: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("kron not supported by provider"))
    }
    /// Compute the cross product of 3-element vectors along a matching dimension.
    fn cross(
        &self,
        _lhs: &GpuTensorHandle,
        _rhs: &GpuTensorHandle,
        _dim: Option<usize>,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("cross not supported by provider"))
    }
    fn reduce_sum<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_sum not supported by provider")
    }
    fn reduce_sum_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_sum_dim not supported by provider")
    }
    /// Reduce a native integer gpuArray while preserving its exact element
    /// class. This is intentionally separate from `reduce_sum`, whose MATLAB
    /// default output semantics are floating point for integer inputs.
    fn reduce_integer_sum_native<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_sum_native not supported by provider")
    }
    /// Dimension form of [`Self::reduce_integer_sum_native`].
    fn reduce_integer_sum_native_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_sum_native_dim not supported by provider")
    }
    fn dot<'a>(
        &'a self,
        _lhs: &'a GpuTensorHandle,
        _rhs: &'a GpuTensorHandle,
        _dim: Option<usize>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("dot not supported by provider")
    }
    fn reduce_nnz<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_nnz not supported by provider")
    }
    fn reduce_nnz_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_nnz_dim not supported by provider")
    }
    fn reduce_prod<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_prod not supported by provider")
    }
    fn reduce_prod_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_prod_dim not supported by provider")
    }
    /// Reduce a native integer gpuArray by product while preserving its exact
    /// element class, distinct from MATLAB's default floating-point output.
    fn reduce_integer_prod_native<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_prod_native not supported by provider")
    }
    /// Dimension form of [`Self::reduce_integer_prod_native`].
    fn reduce_integer_prod_native_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_prod_native_dim not supported by provider")
    }
    /// Compute MATLAB's explicit native-output mean for a native integer
    /// gpuArray. Providers must avoid floating-point accumulation so int64 and
    /// uint64 values remain exact before the final class-preserving rounding.
    fn reduce_integer_mean_native<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_mean_native not supported by provider")
    }
    /// Dimension form of [`Self::reduce_integer_mean_native`].
    fn reduce_integer_mean_native_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_mean_native_dim not supported by provider")
    }
    /// Multi-dimension form of [`Self::reduce_integer_mean_native`]. All
    /// requested zero-based dimensions must be reduced in one logical pass so
    /// native integer output rounds only once.
    fn reduce_integer_mean_native_dims<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dims_zero_based: &'a [usize],
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_integer_mean_native_dims not supported by provider")
    }
    /// Convert a real gpuArray to exact native integer storage while preserving
    /// device residency. Floating-point inputs use MATLAB-compatible saturating
    /// rounding; native integer inputs use exact class-to-class clamping.
    fn cast_to_integer<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _target: IntegerElementType,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("cast_to_integer not supported by provider")
    }
    fn reduce_mean<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_mean not supported by provider")
    }
    /// Reduce mean across multiple zero-based dimensions in one device pass.
    fn reduce_mean_nd<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dims_zero_based: &'a [usize],
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_mean_nd not supported by provider")
    }
    /// Reduce moments across multiple zero-based dimensions in one device pass.
    /// Returns mean (E[x]) and mean of squares (E[x^2]).
    fn reduce_moments_nd<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dims_zero_based: &'a [usize],
    ) -> AccelProviderFuture<'a, ProviderMoments2> {
        unsupported_future("reduce_moments_nd not supported by provider")
    }
    fn reduce_mean_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_mean_dim not supported by provider")
    }
    fn reduce_std<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _normalization: ProviderStdNormalization,
        _nan_mode: ProviderNanMode,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_std not supported by provider")
    }
    fn reduce_std_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
        _normalization: ProviderStdNormalization,
        _nan_mode: ProviderNanMode,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_std_dim not supported by provider")
    }
    fn reduce_any<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _omit_nan: bool,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_any not supported by provider")
    }
    fn reduce_any_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
        _omit_nan: bool,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_any_dim not supported by provider")
    }
    fn reduce_all<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _omit_nan: bool,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_all not supported by provider")
    }
    fn reduce_all_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
        _omit_nan: bool,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_all_dim not supported by provider")
    }
    fn reduce_median<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_median not supported by provider")
    }
    fn reduce_median_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_median_dim not supported by provider")
    }
    fn mode_values<'a>(
        &'a self,
        _request: &'a ProviderModeRequest<'a>,
    ) -> AccelProviderFuture<'a, ProviderModeResult> {
        unsupported_future("mode_values not supported by provider")
    }
    fn moving_window<'a>(
        &'a self,
        _request: &'a ProviderMovingWindowRequest<'a>,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("moving_window not supported by provider")
    }
    fn reduce_min<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_min not supported by provider")
    }
    fn reduce_min_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, ReduceDimResult> {
        unsupported_future("reduce_min_dim not supported by provider")
    }
    fn reduce_max<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
    ) -> AccelProviderFuture<'a, GpuTensorHandle> {
        unsupported_future("reduce_max not supported by provider")
    }
    fn reduce_max_dim<'a>(
        &'a self,
        _a: &'a GpuTensorHandle,
        _dim: usize,
    ) -> AccelProviderFuture<'a, ReduceDimResult> {
        unsupported_future("reduce_max_dim not supported by provider")
    }
    fn cumsum_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
        _nan_mode: ProviderNanMode,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("cumsum_scan not supported by provider"))
    }
    fn integer_cumsum_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "integer_cumsum_scan not supported by provider"
        ))
    }
    fn trapz_dim(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _spacing: ProviderTrapezoidSpacing<'_>,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("trapz_dim not supported by provider"))
    }
    fn cumtrapz_dim(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _spacing: ProviderTrapezoidSpacing<'_>,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("cumtrapz_dim not supported by provider"))
    }
    fn cumprod_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
        _nan_mode: ProviderNanMode,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("cumprod_scan not supported by provider"))
    }
    fn integer_cumprod_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "integer_cumprod_scan not supported by provider"
        ))
    }
    fn cummin_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
        _nan_mode: ProviderNanMode,
    ) -> anyhow::Result<ProviderCumminResult> {
        Err(anyhow::anyhow!("cummin_scan not supported by provider"))
    }
    fn integer_cummin_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
    ) -> anyhow::Result<ProviderCumminResult> {
        Err(anyhow::anyhow!(
            "integer_cummin_scan not supported by provider"
        ))
    }
    fn cummax_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
        _nan_mode: ProviderNanMode,
    ) -> anyhow::Result<ProviderCummaxResult> {
        Err(anyhow::anyhow!("cummax_scan not supported by provider"))
    }
    fn integer_cummax_scan(
        &self,
        _input: &GpuTensorHandle,
        _dim: usize,
        _direction: ProviderScanDirection,
    ) -> anyhow::Result<ProviderCummaxResult> {
        Err(anyhow::anyhow!(
            "integer_cummax_scan not supported by provider"
        ))
    }

    fn find(
        &self,
        _a: &GpuTensorHandle,
        _limit: Option<usize>,
        _direction: FindDirection,
    ) -> anyhow::Result<ProviderFindResult> {
        Err(anyhow::anyhow!("find not supported by provider"))
    }

    fn fused_elementwise(
        &self,
        _shader: &str,
        _inputs: &[GpuTensorHandle],
        _output_shape: &[usize],
        _len: usize,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!(
            "fused_elementwise not supported by provider"
        ))
    }

    /// Execute a single fused elementwise kernel that writes `num_outputs` output buffers in one
    /// dispatch. The shader is expected to declare `output0`, `output1`, … `output{N-1}` storage
    /// bindings (at binding indices `inputs.len()` through `inputs.len() + num_outputs - 1`) and a
    /// uniform `params` binding at `inputs.len() + num_outputs`.
    ///
    /// Providers that do not override this method fall back to calling `fused_elementwise` once
    /// per output, which preserves correctness at the cost of the O(N²) dispatch overhead this
    /// method is designed to eliminate.
    fn fused_elementwise_multi(
        &self,
        _shader: &str,
        _inputs: &[GpuTensorHandle],
        _output_shape: &[usize],
        _len: usize,
        _num_outputs: usize,
    ) -> anyhow::Result<Vec<GpuTensorHandle>> {
        Err(anyhow::anyhow!(
            "fused_elementwise_multi not supported by provider"
        ))
    }

    /// Build a numeric tensor where NaNs in `a` are replaced with 0.0 (device side).
    fn map_nan_to_zero(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("map_nan_to_zero not supported by provider"))
    }

    /// Build a numeric mask tensor with 1.0 where value is not NaN and 0.0 where value is NaN.
    fn not_nan_mask(&self, _a: &GpuTensorHandle) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("not_nan_mask not supported by provider"))
    }

    /// Generic fused reduction entrypoint.
    ///
    /// The shader is expected to implement a column-major reduction across `reduce_len` with
    /// `num_slices` independent slices (e.g., columns). Providers should create a uniform buffer
    /// compatible with the expected `Params/MParams` struct in the shader and dispatch
    /// `num_slices` workgroups with `workgroup_size` threads, or an equivalent strategy.
    #[allow(clippy::too_many_arguments)]
    fn fused_reduction(
        &self,
        _shader: &str,
        _inputs: &[GpuTensorHandle],
        _output_shape: &[usize],
        _reduce_len: usize,
        _num_slices: usize,
        _workgroup_size: u32,
        _flavor: ReductionFlavor,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("fused_reduction not supported by provider"))
    }

    /// Optionally pre-compile commonly used pipelines to amortize first-dispatch costs.
    fn warmup(&self) {}

    /// Returns (cache_hits, cache_misses) for fused pipeline cache, if supported.
    fn fused_cache_counters(&self) -> (u64, u64) {
        (0, 0)
    }

    /// Returns the duration of the last provider warmup in milliseconds, if known.
    fn last_warmup_millis(&self) -> Option<u64> {
        None
    }

    /// Returns a snapshot of provider telemetry counters if supported.
    fn telemetry_snapshot(&self) -> ProviderTelemetry {
        let (hits, misses) = self.fused_cache_counters();
        ProviderTelemetry {
            fused_elementwise: ProviderDispatchStats::default(),
            fused_reduction: ProviderDispatchStats::default(),
            matmul: ProviderDispatchStats::default(),
            linsolve: ProviderDispatchStats::default(),
            mldivide: ProviderDispatchStats::default(),
            mrdivide: ProviderDispatchStats::default(),
            upload_bytes: 0,
            download_bytes: 0,
            solve_fallbacks: Vec::new(),
            fusion_cache_hits: hits,
            fusion_cache_misses: misses,
            bind_group_cache_hits: 0,
            bind_group_cache_misses: 0,
            bind_group_cache_by_layout: None,
            kernel_launches: Vec::new(),
        }
    }

    /// Reset all telemetry counters maintained by the provider, if supported.
    fn reset_telemetry(&self) {}

    /// Default reduction workgroup size the provider prefers.
    fn default_reduction_workgroup_size(&self) -> u32 {
        256
    }

    /// Threshold above which provider will prefer two-pass reduction.
    fn two_pass_threshold(&self) -> usize {
        1024
    }

    /// Current two-pass mode preference (auto/forced on/off).
    fn reduction_two_pass_mode(&self) -> ReductionTwoPassMode {
        ReductionTwoPassMode::Auto
    }

    /// Fast-path: write a GPU column in a matrix from a GPU vector, returning a new handle.
    /// Expected: `values.shape == [rows, 1]` (or `[rows]`) and `col_index < cols`.
    fn scatter_column(
        &self,
        _matrix: &GpuTensorHandle,
        _col_index: usize,
        _values: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scatter_column not supported by provider"))
    }

    /// Fast-path: write a GPU row in a matrix from a GPU vector, returning a new handle.
    /// Expected: `values.shape == [1, cols]` (or `[cols]`) and `row_index < rows`.
    fn scatter_row(
        &self,
        _matrix: &GpuTensorHandle,
        _row_index: usize,
        _values: &GpuTensorHandle,
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("scatter_row not supported by provider"))
    }

    fn sub2ind(
        &self,
        _dims: &[usize],
        _strides: &[usize],
        _inputs: &[&GpuTensorHandle],
        _scalar_mask: &[bool],
        _len: usize,
        _output_shape: &[usize],
    ) -> anyhow::Result<GpuTensorHandle> {
        Err(anyhow::anyhow!("sub2ind not supported by provider"))
    }

    /// Returns true if the provider offers a device-side `ind2sub` implementation.
    fn supports_ind2sub(&self) -> bool {
        false
    }

    /// Convert linear indices into per-dimension subscripts on the device.
    fn ind2sub(
        &self,
        _dims: &[usize],
        _strides: &[usize],
        _indices: &GpuTensorHandle,
        _total: usize,
        _len: usize,
        _output_shape: &[usize],
    ) -> anyhow::Result<Vec<GpuTensorHandle>> {
        Err(anyhow::anyhow!("ind2sub not supported by provider"))
    }

    /// Determine if a matrix is symmetric (or skew-symmetric) without gathering it to the host.
    fn issymmetric(
        &self,
        _matrix: &GpuTensorHandle,
        _kind: ProviderSymmetryKind,
        _tolerance: f64,
    ) -> anyhow::Result<bool> {
        Err(anyhow::anyhow!(
            "issymmetric predicate not supported by provider"
        ))
    }

    /// Determine if a matrix is Hermitian (or skew-Hermitian) without gathering it to the host.
    fn ishermitian<'a>(
        &'a self,
        _matrix: &'a GpuTensorHandle,
        _kind: ProviderHermitianKind,
        _tolerance: f64,
    ) -> AccelProviderFuture<'a, bool> {
        Box::pin(async move {
            Err(anyhow::anyhow!(
                "ishermitian predicate not supported by provider"
            ))
        })
    }

    /// Inspect the bandwidth of a matrix without gathering it back to the host.
    fn bandwidth(&self, _matrix: &GpuTensorHandle) -> anyhow::Result<ProviderBandwidth> {
        Err(anyhow::anyhow!("bandwidth not supported by provider"))
    }

    /// Compute the symmetric reverse Cuthill-McKee permutation for the matrix.
    ///
    /// Implementations may execute on the device or gather to the host. The permutation should be
    /// returned as zero-based indices.
    fn sym_rcm<'a>(&'a self, _matrix: &'a GpuTensorHandle) -> AccelProviderFuture<'a, Vec<usize>> {
        Box::pin(async move { Err(anyhow::anyhow!("sym_rcm not supported by provider")) })
    }
}

static GLOBAL_PROVIDER: Lazy<RwLock<Option<&'static dyn AccelProvider>>> =
    Lazy::new(|| RwLock::new(None));
static PROVIDER_REGISTRY: Lazy<RwLock<HashMap<u32, &'static dyn AccelProvider>>> =
    Lazy::new(|| RwLock::new(HashMap::new()));
static DEVICE_ID_COUNTER: AtomicU32 = AtomicU32::new(1);

#[cfg(not(target_arch = "wasm32"))]
thread_local! {
    static THREAD_PROVIDER: Cell<Option<&'static dyn AccelProvider>> = Cell::new(None);
}

#[cfg(target_arch = "wasm32")]
static WASM_THREAD_PROVIDER: Lazy<Mutex<Option<&'static dyn AccelProvider>>> =
    Lazy::new(|| Mutex::new(None));

#[cfg(not(target_arch = "wasm32"))]
fn replace_thread_provider(
    provider: Option<&'static dyn AccelProvider>,
) -> Option<&'static dyn AccelProvider> {
    THREAD_PROVIDER.with(|cell| {
        let prev = cell.get();
        cell.set(provider);
        prev
    })
}

#[cfg(target_arch = "wasm32")]
fn replace_thread_provider(
    provider: Option<&'static dyn AccelProvider>,
) -> Option<&'static dyn AccelProvider> {
    let mut slot = WASM_THREAD_PROVIDER
        .lock()
        .expect("wasm provider mutex poisoned");
    let prev = *slot;
    *slot = provider;
    prev
}

#[cfg(not(target_arch = "wasm32"))]
fn current_thread_provider() -> Option<&'static dyn AccelProvider> {
    THREAD_PROVIDER.with(|cell| cell.get())
}

#[cfg(target_arch = "wasm32")]
fn current_thread_provider() -> Option<&'static dyn AccelProvider> {
    WASM_THREAD_PROVIDER
        .lock()
        .expect("wasm provider mutex poisoned")
        .as_ref()
        .copied()
}

/// Register a global acceleration provider.
///
/// # Safety
/// - The caller must guarantee that `p` is valid for the entire program lifetime
///   (e.g., a `'static` singleton), as the runtime stores a raw reference globally.
/// - Concurrent callers must ensure registration happens once or is properly
///   synchronized; this function does not enforce thread-safety for re-registration.
/// - Distinct live provider instances that can own handles must not share a
///   `device_id`; the handle registry uses that identifier as its ownership key.
pub unsafe fn register_provider(p: &'static dyn AccelProvider) {
    replace_thread_provider(Some(p));
    if let Ok(mut guard) = GLOBAL_PROVIDER.write() {
        *guard = Some(p);
    }
    register_provider_for_device(p.device_id(), p);
}

unsafe fn register_provider_for_device(device_id: u32, provider: &'static dyn AccelProvider) {
    if let Ok(mut guard) = PROVIDER_REGISTRY.write() {
        guard.insert(device_id, provider);
    }
}

pub fn provider() -> Option<&'static dyn AccelProvider> {
    if let Some(p) = current_thread_provider() {
        return Some(p);
    }
    GLOBAL_PROVIDER
        .read()
        .ok()
        .and_then(|guard| guard.as_ref().copied())
}

/// Clear the active provider selection without invalidating providers that own live handles.
///
/// Registered device owners are retained because a [`GpuTensorHandle`] may outlive its provider's
/// tenure as the global default. Removing its owner here would make that otherwise-valid handle
/// impossible to gather or operate on. Tests that need a different default can safely call this
/// function and register another provider; device ids keep the handle namespaces distinct.
pub fn clear_provider() {
    replace_thread_provider(None);
    if let Ok(mut guard) = GLOBAL_PROVIDER.write() {
        *guard = None;
    }
}

pub fn provider_for_device(device_id: u32) -> Option<&'static dyn AccelProvider> {
    if let Some(registered) = PROVIDER_REGISTRY
        .read()
        .ok()
        .and_then(|guard| guard.get(&device_id).copied())
    {
        return Some(registered);
    }
    if let Some(thread_provider) = current_thread_provider() {
        if thread_provider.device_id() == device_id {
            return Some(thread_provider);
        }
    }
    // Preserve legacy behavior: when no explicit per-device registration exists,
    // fall back to the globally active provider regardless of handle device id.
    GLOBAL_PROVIDER
        .read()
        .ok()
        .and_then(|guard| guard.as_ref().copied())
}

pub fn provider_for_handle(handle: &GpuTensorHandle) -> Option<&'static dyn AccelProvider> {
    provider_for_device(handle.device_id)
}

pub fn spawn_handle_concurrency_for(handle: &GpuTensorHandle) -> Option<SpawnHandleConcurrency> {
    provider_for_handle(handle).map(AccelProvider::spawn_handle_concurrency)
}

pub fn next_device_id() -> u32 {
    DEVICE_ID_COUNTER.fetch_add(1, Ordering::Relaxed)
}

pub struct ThreadProviderGuard {
    prev: Option<&'static dyn AccelProvider>,
}

impl ThreadProviderGuard {
    pub fn set(provider: Option<&'static dyn AccelProvider>) -> Self {
        let prev = replace_thread_provider(provider);
        ThreadProviderGuard { prev }
    }
}

impl Drop for ThreadProviderGuard {
    fn drop(&mut self) {
        let prev = self.prev.take();
        replace_thread_provider(prev);
    }
}

pub fn set_thread_provider(provider: Option<&'static dyn AccelProvider>) {
    replace_thread_provider(provider);
}

/// Convenience: perform elementwise add via provider if possible; otherwise return None
pub async fn try_elem_add(a: &GpuTensorHandle, b: &GpuTensorHandle) -> Option<GpuTensorHandle> {
    if a.device_id == b.device_id {
        let p = provider_for_handle(a)?;
        if let Ok(h) = p.elem_add(a, b).await {
            return Some(h);
        }
    }
    None
}

/// Convenience: perform elementwise hypot via provider if possible; otherwise return None
pub async fn try_elem_hypot(a: &GpuTensorHandle, b: &GpuTensorHandle) -> Option<GpuTensorHandle> {
    if a.device_id == b.device_id {
        let p = provider_for_handle(a)?;
        if let Ok(h) = p.elem_hypot(a, b).await {
            return Some(h);
        }
    }
    None
}

/// Convenience: perform elementwise max via provider if possible; otherwise return None
pub async fn try_elem_max(a: &GpuTensorHandle, b: &GpuTensorHandle) -> Option<GpuTensorHandle> {
    if a.device_id == b.device_id {
        let p = provider_for_handle(a)?;
        if let Ok(h) = p.elem_max(a, b).await {
            return Some(h);
        }
    }
    None
}

/// Convenience: perform elementwise min via provider if possible; otherwise return None
pub async fn try_elem_min(a: &GpuTensorHandle, b: &GpuTensorHandle) -> Option<GpuTensorHandle> {
    if a.device_id == b.device_id {
        let p = provider_for_handle(a)?;
        if let Ok(h) = p.elem_min(a, b).await {
            return Some(h);
        }
    }
    None
}

/// Convenience: perform elementwise atan2 via provider if possible; otherwise return None
pub async fn try_elem_atan2(y: &GpuTensorHandle, x: &GpuTensorHandle) -> Option<GpuTensorHandle> {
    if y.device_id == x.device_id {
        let p = provider_for_handle(y)?;
        if let Ok(h) = p.elem_atan2(y, x).await {
            return Some(h);
        }
    }
    None
}

/// Physical element type carried by a numeric provider transfer.
///
/// This type is exhaustive over RunMat's real numeric classes and lives in the
/// acceleration API so providers do not depend on runtime or builtin storage
/// types.
#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, Serialize, Deserialize)]
pub enum NumericElementType {
    F64,
    F32,
    I8,
    I16,
    I32,
    I64,
    U8,
    U16,
    U32,
    U64,
}

impl NumericElementType {
    pub const fn element_size(self) -> usize {
        match self {
            Self::F64 | Self::I64 | Self::U64 => 8,
            Self::F32 | Self::I32 | Self::U32 => 4,
            Self::I16 | Self::U16 => 2,
            Self::I8 | Self::U8 => 1,
        }
    }

    pub const fn class_name(self) -> &'static str {
        match self {
            Self::F64 => "double",
            Self::F32 => "single",
            Self::I8 => "int8",
            Self::I16 => "int16",
            Self::I32 => "int32",
            Self::I64 => "int64",
            Self::U8 => "uint8",
            Self::U16 => "uint16",
            Self::U32 => "uint32",
            Self::U64 => "uint64",
        }
    }

    pub const fn precision(self) -> Option<ProviderPrecision> {
        match self {
            Self::F64 => Some(ProviderPrecision::F64),
            Self::F32 => Some(ProviderPrecision::F32),
            Self::I8
            | Self::I16
            | Self::I32
            | Self::I64
            | Self::U8
            | Self::U16
            | Self::U32
            | Self::U64 => None,
        }
    }

    pub const fn integer_type(self) -> Option<IntegerElementType> {
        match self {
            Self::F64 | Self::F32 => None,
            Self::I8 => Some(IntegerElementType::I8),
            Self::I16 => Some(IntegerElementType::I16),
            Self::I32 => Some(IntegerElementType::I32),
            Self::I64 => Some(IntegerElementType::I64),
            Self::U8 => Some(IntegerElementType::U8),
            Self::U16 => Some(IntegerElementType::U16),
            Self::U32 => Some(IntegerElementType::U32),
            Self::U64 => Some(IntegerElementType::U64),
        }
    }
}

impl From<IntegerElementType> for NumericElementType {
    fn from(value: IntegerElementType) -> Self {
        match value {
            IntegerElementType::I8 => Self::I8,
            IntegerElementType::I16 => Self::I16,
            IntegerElementType::I32 => Self::I32,
            IntegerElementType::I64 => Self::I64,
            IntegerElementType::U8 => Self::U8,
            IntegerElementType::U16 => Self::U16,
            IntegerElementType::U32 => Self::U32,
            IntegerElementType::U64 => Self::U64,
        }
    }
}

/// Borrowed native numeric buffer used for provider transfers.
#[derive(Debug, Clone, Copy)]
pub enum HostNumericDataView<'a> {
    F64(&'a [f64]),
    F32(&'a [f32]),
    I8(&'a [i8]),
    I16(&'a [i16]),
    I32(&'a [i32]),
    I64(&'a [i64]),
    U8(&'a [u8]),
    U16(&'a [u16]),
    U32(&'a [u32]),
    U64(&'a [u64]),
}

impl HostNumericDataView<'_> {
    pub const fn element_type(self) -> NumericElementType {
        match self {
            Self::F64(_) => NumericElementType::F64,
            Self::F32(_) => NumericElementType::F32,
            Self::I8(_) => NumericElementType::I8,
            Self::I16(_) => NumericElementType::I16,
            Self::I32(_) => NumericElementType::I32,
            Self::I64(_) => NumericElementType::I64,
            Self::U8(_) => NumericElementType::U8,
            Self::U16(_) => NumericElementType::U16,
            Self::U32(_) => NumericElementType::U32,
            Self::U64(_) => NumericElementType::U64,
        }
    }

    pub const fn len(self) -> usize {
        match self {
            Self::F64(data) => data.len(),
            Self::F32(data) => data.len(),
            Self::I8(data) => data.len(),
            Self::I16(data) => data.len(),
            Self::I32(data) => data.len(),
            Self::I64(data) => data.len(),
            Self::U8(data) => data.len(),
            Self::U16(data) => data.len(),
            Self::U32(data) => data.len(),
            Self::U64(data) => data.len(),
        }
    }

    pub const fn is_empty(self) -> bool {
        self.len() == 0
    }
}

/// Borrowed tensor transfer with one native numeric payload and explicit
/// real/complex layout.
#[derive(Debug, Clone, Copy)]
pub struct HostNumericTensorView<'a> {
    pub data: HostNumericDataView<'a>,
    pub shape: &'a [usize],
    pub storage: GpuTensorStorage,
}

impl HostNumericTensorView<'_> {
    pub fn validate(self) -> anyhow::Result<()> {
        let logical_len = self
            .shape
            .iter()
            .try_fold(1usize, |len, &dim| len.checked_mul(dim))
            .ok_or_else(|| anyhow!("numeric transfer shape overflows usize"))?;
        let expected = match self.storage {
            GpuTensorStorage::Real => logical_len,
            GpuTensorStorage::ComplexInterleaved => logical_len
                .checked_mul(2)
                .ok_or_else(|| anyhow!("complex numeric transfer length overflows usize"))?,
        };
        if self.data.len() != expected {
            return Err(anyhow!(
                "{} transfer has {} elements for shape {:?} and {:?} storage; expected {}",
                self.data.element_type().class_name(),
                self.data.len(),
                self.shape,
                self.storage,
                expected
            ));
        }
        Ok(())
    }

    pub const fn element_type(self) -> NumericElementType {
        self.data.element_type()
    }
}

/// Owned native numeric buffer returned by a provider download.
#[derive(Debug, Clone, PartialEq)]
pub enum HostNumericDataOwned {
    F64(Vec<f64>),
    F32(Vec<f32>),
    I8(Vec<i8>),
    I16(Vec<i16>),
    I32(Vec<i32>),
    I64(Vec<i64>),
    U8(Vec<u8>),
    U16(Vec<u16>),
    U32(Vec<u32>),
    U64(Vec<u64>),
}

impl HostNumericDataOwned {
    pub const fn element_type(&self) -> NumericElementType {
        match self {
            Self::F64(_) => NumericElementType::F64,
            Self::F32(_) => NumericElementType::F32,
            Self::I8(_) => NumericElementType::I8,
            Self::I16(_) => NumericElementType::I16,
            Self::I32(_) => NumericElementType::I32,
            Self::I64(_) => NumericElementType::I64,
            Self::U8(_) => NumericElementType::U8,
            Self::U16(_) => NumericElementType::U16,
            Self::U32(_) => NumericElementType::U32,
            Self::U64(_) => NumericElementType::U64,
        }
    }

    pub fn as_view(&self) -> HostNumericDataView<'_> {
        match self {
            Self::F64(data) => HostNumericDataView::F64(data),
            Self::F32(data) => HostNumericDataView::F32(data),
            Self::I8(data) => HostNumericDataView::I8(data),
            Self::I16(data) => HostNumericDataView::I16(data),
            Self::I32(data) => HostNumericDataView::I32(data),
            Self::I64(data) => HostNumericDataView::I64(data),
            Self::U8(data) => HostNumericDataView::U8(data),
            Self::U16(data) => HostNumericDataView::U16(data),
            Self::U32(data) => HostNumericDataView::U32(data),
            Self::U64(data) => HostNumericDataView::U64(data),
        }
    }

    pub fn len(&self) -> usize {
        self.as_view().len()
    }

    pub fn is_empty(&self) -> bool {
        self.len() == 0
    }
}

/// Owned tensor transfer with one native numeric payload and explicit
/// real/complex layout.
#[derive(Debug, Clone, PartialEq)]
pub struct HostNumericTensorOwned {
    pub data: HostNumericDataOwned,
    pub shape: Vec<usize>,
    pub storage: GpuTensorStorage,
}

impl HostNumericTensorOwned {
    pub fn as_view(&self) -> HostNumericTensorView<'_> {
        HostNumericTensorView {
            data: self.data.as_view(),
            shape: &self.shape,
            storage: self.storage,
        }
    }

    pub fn validate(&self) -> anyhow::Result<()> {
        self.as_view().validate()
    }
}

// Minimal host tensor views to avoid depending on runmat-builtins and cycles
#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct HostTensorOwned {
    pub data: Vec<f64>,
    pub shape: Vec<usize>,
    pub storage: GpuTensorStorage,
}

#[derive(Debug)]
pub struct HostTensorView<'a> {
    pub data: &'a [f64],
    pub shape: &'a [usize],
}

/// Borrowed exact native-integer buffer used for provider transfers.
#[derive(Debug, Clone, Copy)]
pub enum HostIntegerDataView<'a> {
    I8(&'a [i8]),
    I16(&'a [i16]),
    I32(&'a [i32]),
    I64(&'a [i64]),
    U8(&'a [u8]),
    U16(&'a [u16]),
    U32(&'a [u32]),
    U64(&'a [u64]),
}

impl HostIntegerDataView<'_> {
    pub fn element_type(self) -> IntegerElementType {
        match self {
            Self::I8(_) => IntegerElementType::I8,
            Self::I16(_) => IntegerElementType::I16,
            Self::I32(_) => IntegerElementType::I32,
            Self::I64(_) => IntegerElementType::I64,
            Self::U8(_) => IntegerElementType::U8,
            Self::U16(_) => IntegerElementType::U16,
            Self::U32(_) => IntegerElementType::U32,
            Self::U64(_) => IntegerElementType::U64,
        }
    }

    pub fn len(self) -> usize {
        match self {
            Self::I8(data) => data.len(),
            Self::I16(data) => data.len(),
            Self::I32(data) => data.len(),
            Self::I64(data) => data.len(),
            Self::U8(data) => data.len(),
            Self::U16(data) => data.len(),
            Self::U32(data) => data.len(),
            Self::U64(data) => data.len(),
        }
    }

    pub fn is_empty(self) -> bool {
        self.len() == 0
    }
}

/// Borrowed exact integer tensor transfer payload.
#[derive(Debug, Clone, Copy)]
pub struct HostIntegerTensorView<'a> {
    pub data: HostIntegerDataView<'a>,
    pub shape: &'a [usize],
}

/// Owned exact native-integer buffer returned by a provider download.
#[derive(Debug, Clone, PartialEq, Eq)]
pub enum HostIntegerDataOwned {
    I8(Vec<i8>),
    I16(Vec<i16>),
    I32(Vec<i32>),
    I64(Vec<i64>),
    U8(Vec<u8>),
    U16(Vec<u16>),
    U32(Vec<u32>),
    U64(Vec<u64>),
}

impl HostIntegerDataOwned {
    pub fn element_type(&self) -> IntegerElementType {
        match self {
            Self::I8(_) => IntegerElementType::I8,
            Self::I16(_) => IntegerElementType::I16,
            Self::I32(_) => IntegerElementType::I32,
            Self::I64(_) => IntegerElementType::I64,
            Self::U8(_) => IntegerElementType::U8,
            Self::U16(_) => IntegerElementType::U16,
            Self::U32(_) => IntegerElementType::U32,
            Self::U64(_) => IntegerElementType::U64,
        }
    }
}

/// Exact integer tensor returned by a provider download.
#[derive(Debug, Clone, PartialEq, Eq)]
pub struct HostIntegerTensorOwned {
    pub data: HostIntegerDataOwned,
    pub shape: Vec<usize>,
}

impl<'a> From<HostTensorView<'a>> for HostNumericTensorView<'a> {
    fn from(value: HostTensorView<'a>) -> Self {
        Self {
            data: HostNumericDataView::F64(value.data),
            shape: value.shape,
            storage: GpuTensorStorage::Real,
        }
    }
}

impl<'a> From<&'a HostTensorOwned> for HostNumericTensorView<'a> {
    fn from(value: &'a HostTensorOwned) -> Self {
        Self {
            data: HostNumericDataView::F64(&value.data),
            shape: &value.shape,
            storage: value.storage,
        }
    }
}

impl<'a> From<HostIntegerDataView<'a>> for HostNumericDataView<'a> {
    fn from(value: HostIntegerDataView<'a>) -> Self {
        match value {
            HostIntegerDataView::I8(data) => Self::I8(data),
            HostIntegerDataView::I16(data) => Self::I16(data),
            HostIntegerDataView::I32(data) => Self::I32(data),
            HostIntegerDataView::I64(data) => Self::I64(data),
            HostIntegerDataView::U8(data) => Self::U8(data),
            HostIntegerDataView::U16(data) => Self::U16(data),
            HostIntegerDataView::U32(data) => Self::U32(data),
            HostIntegerDataView::U64(data) => Self::U64(data),
        }
    }
}

impl<'a> From<HostIntegerTensorView<'a>> for HostNumericTensorView<'a> {
    fn from(value: HostIntegerTensorView<'a>) -> Self {
        Self {
            data: value.data.into(),
            shape: value.shape,
            storage: GpuTensorStorage::Real,
        }
    }
}

impl From<HostIntegerDataOwned> for HostNumericDataOwned {
    fn from(value: HostIntegerDataOwned) -> Self {
        match value {
            HostIntegerDataOwned::I8(data) => Self::I8(data),
            HostIntegerDataOwned::I16(data) => Self::I16(data),
            HostIntegerDataOwned::I32(data) => Self::I32(data),
            HostIntegerDataOwned::I64(data) => Self::I64(data),
            HostIntegerDataOwned::U8(data) => Self::U8(data),
            HostIntegerDataOwned::U16(data) => Self::U16(data),
            HostIntegerDataOwned::U32(data) => Self::U32(data),
            HostIntegerDataOwned::U64(data) => Self::U64(data),
        }
    }
}

impl From<HostIntegerTensorOwned> for HostNumericTensorOwned {
    fn from(value: HostIntegerTensorOwned) -> Self {
        Self {
            data: value.data.into(),
            shape: value.shape,
            storage: GpuTensorStorage::Real,
        }
    }
}

/// Lightweight 1-D axis view used by provider meshgrid hooks.
#[derive(Debug)]
pub struct MeshgridAxisView<'a> {
    pub data: &'a [f64],
}

/// Provider-side meshgrid result containing coordinate tensor handles.
#[derive(Debug, Clone)]
pub struct ProviderMeshgridResult {
    pub outputs: Vec<GpuTensorHandle>,
}

/// Single resident GPU axis supplied to provider-side `ndgrid`.
#[derive(Debug, Clone, Copy)]
pub struct ProviderNdgridAxis<'a> {
    pub handle: &'a GpuTensorHandle,
}

/// Provider-side `ndgrid` request. `output_shape` is already MATLAB-normalized
/// by the runtime, including the single-axis `[n, 1]` case.
#[derive(Debug, Clone, Copy)]
pub struct ProviderNdgridRequest<'a> {
    pub axes: &'a [ProviderNdgridAxis<'a>],
    pub output_shape: &'a [usize],
    pub output_count: usize,
}

/// Provider-side ndgrid result containing coordinate tensor handles.
#[derive(Debug, Clone)]
pub struct ProviderNdgridResult {
    pub outputs: Vec<GpuTensorHandle>,
}

/// Single broadcastable resident GPU input supplied to provider-side Black-Scholes pricing.
#[derive(Debug, Clone, Copy)]
pub struct ProviderBlackScholesPriceInput<'a> {
    pub handle: &'a GpuTensorHandle,
    /// Input shape aligned to `output_shape` rank using MATLAB implicit-expansion rules.
    pub shape: &'a [usize],
    /// Column-major strides for `shape`.
    pub strides: &'a [usize],
}

/// Provider-side Black-Scholes price request. Inputs are ordered as
/// Price, Strike, Rate, Time, Volatility, Yield.
#[derive(Debug, Clone, Copy)]
pub struct ProviderBlackScholesPriceRequest<'a> {
    pub inputs: &'a [ProviderBlackScholesPriceInput<'a>],
    pub output_shape: &'a [usize],
    pub len: usize,
}

/// Provider-side Black-Scholes price result containing call and put handles.
#[derive(Debug, Clone)]
pub struct ProviderBlackScholesPriceResult {
    pub call: GpuTensorHandle,
    pub put: GpuTensorHandle,
}

/// Provider-side covariance-to-correlation result containing the correlation
/// matrix and column-vector standard deviations.
#[derive(Debug, Clone)]
pub struct ProviderCovarianceToCorrelationResult {
    pub correlation: GpuTensorHandle,
    pub sigma: GpuTensorHandle,
}

/// Provider-side Adam optimizer update request.
#[derive(Debug, Clone, Copy)]
pub struct ProviderAdamUpdateRequest<'a> {
    pub parameters: &'a GpuTensorHandle,
    pub gradient: &'a GpuTensorHandle,
    pub average_grad: Option<&'a GpuTensorHandle>,
    pub average_sq_grad: Option<&'a GpuTensorHandle>,
    pub iteration: usize,
    pub learn_rate: f64,
    pub gradient_decay_factor: f64,
    pub squared_gradient_decay_factor: f64,
    pub epsilon: f64,
}

/// Provider-side Adam optimizer update result.
#[derive(Debug, Clone)]
pub struct ProviderAdamUpdateResult {
    pub parameters: GpuTensorHandle,
    pub average_grad: GpuTensorHandle,
    pub average_sq_grad: GpuTensorHandle,
}

/// Provider-side cross-entropy mode.
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum ProviderCrossentropyMode {
    SingleLabel,
    MultiLabel,
}

/// Provider-side cross-entropy request.
#[derive(Debug, Clone, Copy)]
pub struct ProviderCrossentropyRequest<'a> {
    pub predictions: &'a GpuTensorHandle,
    pub targets: &'a GpuTensorHandle,
    pub weights: Option<&'a GpuTensorHandle>,
    pub mask: Option<&'a GpuTensorHandle>,
    pub mode: ProviderCrossentropyMode,
}

/// Provider-side cross-entropy result containing per-element loss terms.
#[derive(Debug, Clone)]
pub struct ProviderCrossentropyResult {
    pub losses: GpuTensorHandle,
}

/// Descriptor for GEMM epilogues applied to `C = A * B` before storing to `C`.
///
/// Supported operations:
/// - Scale by `alpha` and add scalar `beta`.
/// - Multiply output by per-row and/or per-column scale vectors (broadcasted).
#[derive(Debug, Clone, Copy, PartialEq, Eq, Serialize, Deserialize)]
pub enum ScaleOp {
    Multiply,
    Divide,
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct MatmulEpilogue {
    /// Scalar multiply applied to each output element.
    pub alpha: f64,
    /// Scalar add applied to each output element after scaling.
    pub beta: f64,
    /// Optional per-row scale (length m). When present, output[row, col] *= row_scale[row].
    pub row_scale: Option<GpuTensorHandle>,
    /// Optional per-column scale (length n). When present, output[row, col] *= col_scale[col].
    pub col_scale: Option<GpuTensorHandle>,
    /// Row scale operation (multiply or divide). Ignored when `row_scale` is None.
    pub row_op: ScaleOp,
    /// Column scale operation (multiply or divide). Ignored when `col_scale` is None.
    pub col_op: ScaleOp,
    /// Optional lower clamp bound applied after scale/bias.
    #[serde(default)]
    pub clamp_min: Option<f64>,
    /// Optional upper clamp bound applied after scale/bias.
    #[serde(default)]
    pub clamp_max: Option<f64>,
    /// Optional power exponent applied after clamp (final operation in the epilogue).
    #[serde(default)]
    pub pow_exponent: Option<f64>,
    /// Optional output buffer for the diagonal of the result (length min(m, n)).
    #[serde(default)]
    pub diag_output: Option<GpuTensorHandle>,
}

impl MatmulEpilogue {
    pub fn noop() -> Self {
        Self {
            alpha: 1.0,
            beta: 0.0,
            row_scale: None,
            col_scale: None,
            row_op: ScaleOp::Multiply,
            col_op: ScaleOp::Multiply,
            clamp_min: None,
            clamp_max: None,
            pow_exponent: None,
            diag_output: None,
        }
    }
    pub fn is_noop(&self) -> bool {
        self.alpha == 1.0
            && self.beta == 0.0
            && self.row_scale.is_none()
            && self.col_scale.is_none()
            && self.clamp_min.is_none()
            && self.clamp_max.is_none()
            && self.pow_exponent.is_none()
            && self.diag_output.is_none()
    }
}

#[derive(Debug, Clone, Copy, PartialEq, Serialize, Deserialize)]
pub struct PowerStepEpilogue {
    pub epsilon: f64,
}

impl Default for PowerStepEpilogue {
    fn default() -> Self {
        Self { epsilon: 0.0 }
    }
}

#[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
pub struct ImageNormalizeDescriptor {
    pub batch: usize,
    pub height: usize,
    pub width: usize,
    pub epsilon: f64,
    #[serde(default)]
    pub gain: Option<f64>,
    #[serde(default)]
    pub bias: Option<f64>,
    #[serde(default)]
    pub gamma: Option<f64>,
    #[serde(default = "default_image_normalize_clamp_zero")]
    pub clamp_zero: bool,
}

fn default_image_normalize_clamp_zero() -> bool {
    true
}

#[cfg(test)]
mod tests {
    use super::*;

    struct TestProvider {
        device_id: u32,
        name: &'static str,
        spawn_concurrency: SpawnHandleConcurrency,
    }

    struct DefaultNumericAdapterProvider;

    impl AccelProvider for DefaultNumericAdapterProvider {
        fn upload(&self, host: &HostTensorView) -> anyhow::Result<GpuTensorHandle> {
            assert_eq!(host.data, &[1.0, 2.0]);
            Ok(GpuTensorHandle {
                shape: host.shape.to_vec(),
                device_id: 404,
                buffer_id: 1,
                descriptor: GpuTensorDescriptor::numeric(
                    NumericElementType::F64,
                    GpuTensorStorage::Real,
                ),
            })
        }

        fn download<'a>(&'a self, h: &'a GpuTensorHandle) -> AccelDownloadFuture<'a> {
            Box::pin(async move {
                Ok(HostTensorOwned {
                    data: vec![1.0, 2.0],
                    shape: h.shape.clone(),
                    storage: GpuTensorStorage::Real,
                })
            })
        }

        fn upload_integer(&self, host: &HostIntegerTensorView) -> anyhow::Result<GpuTensorHandle> {
            assert!(matches!(host.data, HostIntegerDataView::U64([1, u64::MAX])));
            Ok(GpuTensorHandle {
                shape: host.shape.to_vec(),
                device_id: 404,
                buffer_id: 2,
                descriptor: GpuTensorDescriptor::numeric(
                    NumericElementType::U64,
                    GpuTensorStorage::Real,
                ),
            })
        }

        fn download_integer<'a>(
            &'a self,
            h: &'a GpuTensorHandle,
        ) -> AccelIntegerDownloadFuture<'a> {
            Box::pin(async move {
                Ok(HostIntegerTensorOwned {
                    data: HostIntegerDataOwned::U64(vec![1, u64::MAX]),
                    shape: h.shape.clone(),
                })
            })
        }

        fn free(&self, _h: &GpuTensorHandle) -> anyhow::Result<()> {
            Ok(())
        }

        fn device_info(&self) -> String {
            "legacy-numeric-provider".to_string()
        }
    }

    impl AccelProvider for TestProvider {
        fn upload(&self, _host: &HostTensorView) -> anyhow::Result<GpuTensorHandle> {
            Err(anyhow!("test provider upload should not be called"))
        }

        fn download<'a>(&'a self, _h: &'a GpuTensorHandle) -> AccelDownloadFuture<'a> {
            unsupported_future("test provider download should not be called")
        }

        fn free(&self, _h: &GpuTensorHandle) -> anyhow::Result<()> {
            Err(anyhow!("test provider free should not be called"))
        }

        fn device_info(&self) -> String {
            self.name.to_string()
        }

        fn device_id(&self) -> u32 {
            self.device_id
        }

        fn spawn_handle_concurrency(&self) -> SpawnHandleConcurrency {
            self.spawn_concurrency
        }
    }

    static PROVIDER_TEST_LOCK: Lazy<std::sync::Mutex<()>> = Lazy::new(|| std::sync::Mutex::new(()));
    static PROVIDER_A: TestProvider = TestProvider {
        device_id: 101,
        name: "provider-a",
        spawn_concurrency: SpawnHandleConcurrency::ImmutableShare,
    };
    static PROVIDER_B: TestProvider = TestProvider {
        device_id: 202,
        name: "provider-b",
        spawn_concurrency: SpawnHandleConcurrency::Reject,
    };
    static PROVIDER_C: TestProvider = TestProvider {
        device_id: 303,
        name: "provider-c",
        spawn_concurrency: SpawnHandleConcurrency::CopyOnWrite,
    };

    fn resolve_ready<F: Future>(future: F) -> F::Output {
        struct NoopWake;
        impl std::task::Wake for NoopWake {
            fn wake(self: std::sync::Arc<Self>) {}
        }
        let waker = std::task::Waker::from(std::sync::Arc::new(NoopWake));
        let mut context = std::task::Context::from_waker(&waker);
        let mut future = std::pin::pin!(future);
        match future.as_mut().poll(&mut context) {
            std::task::Poll::Ready(output) => output,
            std::task::Poll::Pending => panic!("test adapter future unexpectedly pending"),
        }
    }

    fn register_test_providers() {
        clear_provider();
        unsafe {
            register_provider(&PROVIDER_A);
            register_provider(&PROVIDER_B);
        }
    }

    fn test_handle(device_id: u32) -> GpuTensorHandle {
        GpuTensorHandle::new(vec![1], device_id, 42)
    }

    #[test]
    fn numeric_transfer_contract_is_exhaustive_and_native() {
        let f64_values = [1.0_f64, 2.0];
        let f32_values = [1.0_f32, 2.0];
        let i8_values = [1_i8, 2];
        let i16_values = [1_i16, 2];
        let i32_values = [1_i32, 2];
        let i64_values = [1_i64, 9_007_199_254_740_993];
        let u8_values = [1_u8, 2];
        let u16_values = [1_u16, 2];
        let u32_values = [1_u32, 2];
        let u64_values = [1_u64, 9_007_199_254_740_993];
        let cases = [
            (
                HostNumericDataView::F64(&f64_values),
                NumericElementType::F64,
            ),
            (
                HostNumericDataView::F32(&f32_values),
                NumericElementType::F32,
            ),
            (HostNumericDataView::I8(&i8_values), NumericElementType::I8),
            (
                HostNumericDataView::I16(&i16_values),
                NumericElementType::I16,
            ),
            (
                HostNumericDataView::I32(&i32_values),
                NumericElementType::I32,
            ),
            (
                HostNumericDataView::I64(&i64_values),
                NumericElementType::I64,
            ),
            (HostNumericDataView::U8(&u8_values), NumericElementType::U8),
            (
                HostNumericDataView::U16(&u16_values),
                NumericElementType::U16,
            ),
            (
                HostNumericDataView::U32(&u32_values),
                NumericElementType::U32,
            ),
            (
                HostNumericDataView::U64(&u64_values),
                NumericElementType::U64,
            ),
        ];

        for (data, element_type) in cases {
            let transfer = HostNumericTensorView {
                data,
                shape: &[1, 2],
                storage: GpuTensorStorage::Real,
            };
            transfer.validate().expect("matching native transfer");
            assert_eq!(transfer.element_type(), element_type);
            assert_eq!(element_type.class_name(), data.element_type().class_name());
            assert_eq!(
                element_type.element_size(),
                data.element_type().element_size()
            );
        }

        assert_eq!(
            NumericElementType::F64.precision(),
            Some(ProviderPrecision::F64)
        );
        assert_eq!(
            NumericElementType::F32.precision(),
            Some(ProviderPrecision::F32)
        );
        assert_eq!(
            NumericElementType::U64.integer_type(),
            Some(IntegerElementType::U64)
        );
        assert_eq!(
            NumericElementType::I64.integer_type(),
            Some(IntegerElementType::I64)
        );
        assert_eq!(NumericElementType::F64.integer_type(), None);
    }

    #[test]
    fn numeric_transfer_contract_validates_real_and_complex_lengths() {
        let real = HostNumericTensorOwned {
            data: HostNumericDataOwned::F32(vec![1.0, 2.0]),
            shape: vec![2, 1],
            storage: GpuTensorStorage::Real,
        };
        real.validate().expect("native single real transfer");

        let complex = HostNumericTensorOwned {
            data: HostNumericDataOwned::U64(vec![1, u64::MAX, 2, u64::MAX - 1]),
            shape: vec![1, 2],
            storage: GpuTensorStorage::ComplexInterleaved,
        };
        complex.validate().expect("native complex integer transfer");

        let mismatch = HostNumericTensorView {
            data: HostNumericDataView::I16(&[1, 2, 3]),
            shape: &[2, 2],
            storage: GpuTensorStorage::Real,
        };
        let error = mismatch.validate().expect_err("shape mismatch must reject");
        assert!(error.to_string().contains("expected 4"));

        let overflow = HostNumericTensorView {
            data: HostNumericDataView::F64(&[]),
            shape: &[usize::MAX, 2],
            storage: GpuTensorStorage::Real,
        };
        assert!(overflow.validate().is_err());
    }

    #[test]
    fn gpu_handle_descriptor_survives_clone_serialization_and_semantic_metadata_cleanup() {
        let handle = GpuTensorHandle::new(vec![2, 3], 17, 29)
            .with_numeric_descriptor(
                NumericElementType::U64,
                GpuTensorStorage::ComplexInterleaved,
            )
            .with_provenance(GpuHandleProvenance::Explicit);
        let cloned = handle.clone();
        clear_handle_metadata(&handle);
        assert_eq!(handle_integer_type(&cloned), Some(IntegerElementType::U64));
        assert_eq!(handle_precision(&cloned), None);
        assert_eq!(
            handle_storage(&cloned),
            GpuTensorStorage::ComplexInterleaved
        );
        assert_eq!(
            handle_provenance(&cloned),
            Some(GpuHandleProvenance::Explicit)
        );
        assert_eq!(handle_class_name(&cloned).as_deref(), Some("uint64"));

        let encoded = serde_json::to_string(&cloned).expect("serialize durable handle");
        let decoded: GpuTensorHandle =
            serde_json::from_str(&encoded).expect("deserialize durable handle");
        assert_eq!(decoded, cloned);

        let legacy: GpuTensorHandle =
            serde_json::from_str(r#"{"shape":[1,1],"device_id":5,"buffer_id":8}"#)
                .expect("deserialize legacy handle");
        assert_eq!(legacy.descriptor, GpuTensorDescriptor::default());
        assert_eq!(handle_storage(&legacy), GpuTensorStorage::Real);
    }

    #[test]
    fn durable_physical_descriptor_is_the_only_numeric_authority() {
        let floating = GpuTensorHandle::new(vec![1, 2], 23, 31).with_numeric_descriptor(
            NumericElementType::F32,
            GpuTensorStorage::ComplexInterleaved,
        );

        assert_eq!(handle_precision(&floating), Some(ProviderPrecision::F32));
        assert_eq!(handle_integer_type(&floating), None);
        assert_eq!(
            handle_storage(&floating),
            GpuTensorStorage::ComplexInterleaved
        );
        assert_eq!(handle_class_name(&floating).as_deref(), Some("single"));
        clear_handle_metadata(&floating);

        let integer = GpuTensorHandle::new(vec![1, 2], 37, 41)
            .with_numeric_descriptor(NumericElementType::U64, GpuTensorStorage::Real);

        assert_eq!(handle_precision(&integer), None);
        assert_eq!(handle_integer_type(&integer), Some(IntegerElementType::U64));
        assert_eq!(handle_class_name(&integer).as_deref(), Some("uint64"));
        clear_handle_metadata(&integer);
    }

    #[test]
    fn removed_physical_side_metadata_surface_stays_absent() {
        let source = include_str!("lib.rs");
        let forbidden = [
            ["set_handle_", "precision"].concat(),
            ["set_handle_", "integer_type"].concat(),
            ["set_handle_", "storage"].concat(),
            ["clear_handle_", "precision"].concat(),
            ["clear_handle_", "integer_type"].concat(),
            ["clear_handle_", "storage"].concat(),
            ["HANDLE_", "PRECISIONS"].concat(),
            ["HANDLE_", "INTEGER_TYPES"].concat(),
            ["HANDLE_", "STORAGES"].concat(),
        ];
        for symbol in forbidden {
            assert!(
                !source.contains(&symbol),
                "removed symbol returned: {symbol}"
            );
        }
    }

    #[test]
    fn numeric_transfer_default_adapter_preserves_double_and_integer_storage() {
        let provider = DefaultNumericAdapterProvider;
        let double = HostNumericTensorView {
            data: HostNumericDataView::F64(&[1.0, 2.0]),
            shape: &[1, 2],
            storage: GpuTensorStorage::Real,
        };
        let double_handle = provider
            .upload_numeric(&double)
            .expect("double upload adapter");
        let double_download = resolve_ready(provider.download_numeric(&double_handle))
            .expect("double download adapter");
        assert_eq!(
            double_download.data,
            HostNumericDataOwned::F64(vec![1.0, 2.0])
        );

        let integer = HostNumericTensorView {
            data: HostNumericDataView::U64(&[1, u64::MAX]),
            shape: &[1, 2],
            storage: GpuTensorStorage::Real,
        };
        let integer_handle = provider
            .upload_numeric(&integer)
            .expect("integer upload adapter");
        let integer_download = resolve_ready(provider.download_numeric(&integer_handle))
            .expect("integer download adapter");
        assert_eq!(
            integer_download.data,
            HostNumericDataOwned::U64(vec![1, u64::MAX])
        );
        clear_handle_metadata(&integer_handle);
    }

    #[test]
    fn numeric_transfer_default_adapter_rejects_widened_single_and_complex_integer() {
        let provider = DefaultNumericAdapterProvider;
        let single = HostNumericTensorView {
            data: HostNumericDataView::F32(&[1.0, 2.0]),
            shape: &[1, 2],
            storage: GpuTensorStorage::Real,
        };
        assert!(provider.upload_numeric(&single).is_err());

        let complex_integer = HostNumericTensorView {
            data: HostNumericDataView::I16(&[1, 2, 3, 4]),
            shape: &[1, 2],
            storage: GpuTensorStorage::ComplexInterleaved,
        };
        assert!(provider.upload_numeric(&complex_integer).is_err());

        let single_handle = GpuTensorHandle {
            shape: vec![1, 2],
            device_id: 404,
            buffer_id: 3,
            descriptor: GpuTensorDescriptor::numeric(
                NumericElementType::F32,
                GpuTensorStorage::Real,
            ),
        };
        let error = resolve_ready(provider.download_numeric(&single_handle))
            .expect_err("widened single download must reject");
        assert!(error.to_string().contains("native single"));
        clear_handle_metadata(&single_handle);

        let complex_integer_handle = GpuTensorHandle {
            shape: vec![1, 2],
            device_id: 404,
            buffer_id: 4,
            descriptor: GpuTensorDescriptor::numeric(
                NumericElementType::I16,
                GpuTensorStorage::ComplexInterleaved,
            ),
        };
        let error = resolve_ready(provider.download_numeric(&complex_integer_handle))
            .expect_err("complex integer default download must reject");
        assert!(error.to_string().contains("complex integer"));
        clear_handle_metadata(&complex_integer_handle);

        let descriptorless = GpuTensorHandle::new(vec![1, 2], 404, 5);
        let error = resolve_ready(provider.download_numeric(&descriptorless))
            .expect_err("descriptorless numeric download must reject");
        assert!(error.to_string().contains("durable element descriptor"));
    }

    #[test]
    fn conservative_feasibility_rejects_operations_without_execution() {
        let query = ProviderFeasibilityQuery {
            operation: ProviderOperationIdentity::new("legacy.elementwise"),
            family: ProviderOperationFamily::Elementwise,
            inputs: vec![ProviderRepresentation {
                element_type: ProviderElementType::F64,
                storage: ProviderStorage::DenseReal,
                layout: ProviderLayout::ColumnMajorContiguous,
                shape: vec![4, 4],
                residency: ProviderResidency::Host,
            }],
            outputs: Vec::new(),
            workload: ProviderWorkload {
                elements: Some(16),
                ..ProviderWorkload::default()
            },
        };

        assert!(matches!(
            PROVIDER_A.query_feasibility(&query),
            ProviderFeasibility::Rejected {
                rejection: ProviderRejection {
                    code: ProviderRejectionCode::UnsupportedOperation,
                    ..
                }
            }
        ));
        assert!(!PROVIDER_A
            .capability_snapshot()
            .supports(ProviderOperationFamily::Elementwise));

        let snapshot = PROVIDER_A.capability_snapshot();
        let encoded = serde_json::to_string(&snapshot).expect("serialize capability snapshot");
        let decoded: ProviderCapabilitySnapshot =
            serde_json::from_str(&encoded).expect("deserialize capability snapshot");
        assert_eq!(decoded, snapshot);
    }

    #[test]
    fn semantic_handle_metadata_is_namespaced_by_device_and_buffer() {
        let mut first = test_handle(PROVIDER_A.device_id())
            .with_numeric_descriptor(NumericElementType::F32, GpuTensorStorage::Real);
        let mut second = test_handle(PROVIDER_B.device_id()).with_numeric_descriptor(
            NumericElementType::U64,
            GpuTensorStorage::ComplexInterleaved,
        );
        clear_handle_metadata(&first);
        clear_handle_metadata(&second);

        set_handle_class_name(&first, "single");
        set_handle_class_name(&second, "uint64");
        set_handle_logical(&first, true);
        record_handle_transpose(&first, 2, 3);
        mark_handle_automatic(&mut first);
        mark_handle_explicit(&mut second);

        assert_eq!(handle_precision(&first), Some(ProviderPrecision::F32));
        assert_eq!(handle_precision(&second), None);
        assert_eq!(handle_class_name(&first).as_deref(), Some("logical"));
        assert_eq!(handle_class_name(&second).as_deref(), Some("uint64"));
        assert_eq!(handle_integer_type(&first), None);
        assert_eq!(handle_integer_type(&second), Some(IntegerElementType::U64));
        assert!(handle_is_logical(&first));
        assert!(!handle_is_logical(&second));
        assert_eq!(
            handle_transpose_info(&first).map(|info| (info.base_rows, info.base_cols)),
            Some((2, 3))
        );
        assert_eq!(handle_transpose_info(&second), None);
        assert_eq!(handle_storage(&first), GpuTensorStorage::Real);
        assert_eq!(
            handle_storage(&second),
            GpuTensorStorage::ComplexInterleaved
        );
        assert_eq!(
            handle_provenance(&first),
            Some(GpuHandleProvenance::Automatic)
        );
        assert!(handle_is_explicit(&second));

        clear_handle_metadata(&first);
        assert_eq!(handle_precision(&first), Some(ProviderPrecision::F32));
        assert_eq!(handle_class_name(&first).as_deref(), Some("single"));
        assert!(!handle_is_logical(&first));
        assert_eq!(handle_transpose_info(&first), None);
        assert_eq!(
            handle_provenance(&first),
            Some(GpuHandleProvenance::Automatic)
        );
        assert_eq!(handle_precision(&second), None);
        assert_eq!(handle_integer_type(&second), Some(IntegerElementType::U64));
        assert!(handle_is_explicit(&second));
        clear_handle_metadata(&second);
    }

    fn spectral_request<'a>(
        input: &'a GpuTensorHandle,
        frame_mode: ProviderSpectralFrameMode,
    ) -> ProviderSpectralRequest<'a> {
        static WINDOW: [f64; 4] = [1.0, 1.0, 1.0, 1.0];
        ProviderSpectralRequest {
            input,
            input_len: 16,
            input_complex: false,
            window: &WINDOW,
            nfft: 8,
            frame_count: 3,
            frame_mode,
            range: ProviderSpectralRange::Onesided,
            denominator: 1.0,
        }
    }

    #[test]
    fn provider_envelope_shape_guard_rejects_equal_len_layout_spoofing() {
        assert!(provider_envelope_input_shape_matches(&[2, 3], 2, 3));
        assert!(provider_envelope_input_shape_matches(&[6, 1], 6, 1));
        assert!(provider_envelope_input_shape_matches(&[1, 6], 6, 1));
        assert!(provider_envelope_input_shape_matches(&[6], 6, 1));

        assert!(!provider_envelope_input_shape_matches(&[3, 2], 2, 3));
        assert!(!provider_envelope_input_shape_matches(&[6], 2, 3));
        assert!(!provider_envelope_input_shape_matches(&[2, 1, 3], 2, 3));
    }

    #[test]
    fn provider_for_device_prefers_registered_device_over_thread_provider() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        register_test_providers();
        let _thread_provider = ThreadProviderGuard::set(Some(&PROVIDER_B));

        let provider = provider_for_device(PROVIDER_A.device_id()).expect("provider for device");

        assert_eq!(provider.device_info(), PROVIDER_A.name);
        clear_provider();
    }

    #[test]
    fn provider_for_handle_uses_handle_device_owner() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        register_test_providers();
        let _thread_provider = ThreadProviderGuard::set(Some(&PROVIDER_B));

        let provider =
            provider_for_handle(&test_handle(PROVIDER_A.device_id())).expect("provider for handle");

        assert_eq!(provider.device_info(), PROVIDER_A.name);
        clear_provider();
    }

    #[test]
    fn clearing_active_provider_preserves_live_handle_owner() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        register_test_providers();
        let handle = test_handle(PROVIDER_A.device_id());

        clear_provider();

        assert!(provider().is_none());
        let owner = provider_for_handle(&handle).expect("registered handle owner survives clear");
        assert_eq!(owner.device_info(), PROVIDER_A.name);
    }

    #[test]
    fn concurrent_registration_keeps_each_threads_active_provider() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        clear_provider();

        let first = std::thread::spawn(|| {
            unsafe { register_provider(&PROVIDER_A) };
            std::thread::yield_now();
            provider().map(AccelProvider::device_info)
        });
        let second = std::thread::spawn(|| {
            unsafe { register_provider(&PROVIDER_B) };
            std::thread::yield_now();
            provider().map(AccelProvider::device_info)
        });

        assert_eq!(
            first.join().expect("first registration thread"),
            Some(PROVIDER_A.name.to_string())
        );
        assert_eq!(
            second.join().expect("second registration thread"),
            Some(PROVIDER_B.name.to_string())
        );
        clear_provider();
    }

    #[test]
    fn spawn_handle_concurrency_for_uses_registered_owner() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        register_test_providers();
        let _thread_provider = ThreadProviderGuard::set(Some(&PROVIDER_B));

        let concurrency = spawn_handle_concurrency_for(&test_handle(PROVIDER_A.device_id()))
            .expect("spawn concurrency");

        assert_eq!(concurrency, PROVIDER_A.spawn_concurrency);
        clear_provider();
    }

    #[test]
    fn provider_keeps_thread_local_active_provider_semantics() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        register_test_providers();
        let _thread_provider = ThreadProviderGuard::set(Some(&PROVIDER_A));

        let active = provider().expect("active provider");

        assert_eq!(active.device_info(), PROVIDER_A.name);
        clear_provider();
    }

    #[test]
    fn unregistered_thread_provider_only_matches_own_device_before_global_fallback() {
        let _lock = PROVIDER_TEST_LOCK
            .lock()
            .expect("provider test lock poisoned");
        clear_provider();
        unsafe {
            register_provider(&PROVIDER_A);
        }
        let _thread_provider = ThreadProviderGuard::set(Some(&PROVIDER_C));

        let own_device = provider_for_device(PROVIDER_C.device_id()).expect("own provider");
        let fallback = provider_for_device(404).expect("global fallback provider");

        assert_eq!(own_device.device_info(), PROVIDER_C.name);
        assert_eq!(fallback.device_info(), PROVIDER_A.name);
        clear_provider();
    }

    #[test]
    fn uniform_spectral_request_validates_sliding_input_coverage() {
        let input = test_handle(PROVIDER_A.device_id());
        let mut request = spectral_request(&input, ProviderSpectralFrameMode::Sliding { hop: 6 });
        assert!(validate_uniform_spectral_request(&request).is_ok());

        request.input_len = 15;
        assert!(validate_uniform_spectral_request(&request).is_err());
    }

    #[test]
    fn uniform_spectral_request_rejects_sliding_coverage_overflow() {
        let input = test_handle(PROVIDER_A.device_id());
        let mut request = spectral_request(&input, ProviderSpectralFrameMode::Sliding { hop: 2 });
        request.frame_count = usize::MAX;

        assert!(validate_uniform_spectral_request(&request).is_err());
    }

    #[test]
    fn uniform_spectral_request_validates_folded_input_coverage() {
        let input = test_handle(PROVIDER_A.device_id());
        let mut request = spectral_request(
            &input,
            ProviderSpectralFrameMode::FoldedColumns { input_rows: 5 },
        );
        assert!(validate_uniform_spectral_request(&request).is_ok());

        request.frame_mode = ProviderSpectralFrameMode::FoldedColumns { input_rows: 0 };
        assert!(validate_uniform_spectral_request(&request).is_err());

        request.frame_mode = ProviderSpectralFrameMode::FoldedColumns { input_rows: 6 };
        assert!(validate_uniform_spectral_request(&request).is_err());
    }

    #[test]
    fn uniform_spectral_request_rejects_folded_coverage_overflow() {
        let input = test_handle(PROVIDER_A.device_id());
        let request = spectral_request(
            &input,
            ProviderSpectralFrameMode::FoldedColumns {
                input_rows: usize::MAX,
            },
        );

        assert!(validate_uniform_spectral_request(&request).is_err());
    }

    #[test]
    fn image_normalize_descriptor_omitted_clamp_zero_defaults_true() {
        let payload = r#"{
            "batch": 2,
            "height": 4,
            "width": 5,
            "epsilon": 0.000001
        }"#;

        let desc: ImageNormalizeDescriptor =
            serde_json::from_str(payload).expect("deserialize descriptor");

        assert!(
            desc.clamp_zero,
            "legacy serialized descriptors should default to clamped image normalize"
        );
    }

    #[test]
    fn image_normalize_descriptor_explicit_false_preserves_unclamped() {
        let payload = r#"{
            "batch": 2,
            "height": 4,
            "width": 5,
            "epsilon": 0.000001,
            "clamp_zero": false
        }"#;

        let desc: ImageNormalizeDescriptor =
            serde_json::from_str(payload).expect("deserialize descriptor");

        assert!(
            !desc.clamp_zero,
            "explicit clamp_zero=false should preserve unclamped semantics"
        );
    }
}