runmat_runtime/builtins/math/rounding/

mod.rs

//! MATLAB-compatible `mod` builtin plus rounding helpers for RunMat.

pub(crate) mod ceil;
pub(crate) mod fix;
pub(crate) mod floor;
pub(crate) mod rem;
pub(crate) mod round;

use runmat_accelerate_api::GpuTensorHandle;
use runmat_builtins::{
    BuiltinCompletionPolicy, BuiltinDescriptor, BuiltinErrorDescriptor, BuiltinIntegerBackendRule,
    BuiltinIntegerCapabilityDescriptor, BuiltinIntegerComputationDomain,
    BuiltinIntegerInputAvailability, BuiltinIntegerInputCapability, BuiltinIntegerOutputClassRule,
    BuiltinIntegerOverflowRule, BuiltinIntegerOverloadKind, BuiltinIntegerScalarDoubleRule,
    BuiltinOutputMode, BuiltinParamArity, BuiltinParamDescriptor, BuiltinParamType,
    BuiltinSignatureDescriptor,
};
use runmat_macros::runtime_builtin;
use runmat_value::{ComplexTensor, NumericDType, Tensor, Value};

use crate::builtins::common::broadcast::BroadcastPlan;
use crate::builtins::common::spec::{
    BroadcastSemantics, BuiltinFusionSpec, BuiltinGpuSpec, ConstantStrategy, FusionError,
    FusionExprContext, FusionKernelTemplate, GpuOpKind, ProviderHook, ReductionNaN,
    ResidencyPolicy, ScalarType, ShapeRequirements,
};
use crate::builtins::common::{gpu_helpers, tensor};
use crate::builtins::math::elementwise::integer_arithmetic::{
    reject_integer_logical_operands, try_integer_remainder, IntegerRemainderOp,
};
use crate::builtins::math::type_resolvers::numeric_binary_type;
use crate::{build_runtime_error, BuiltinResult, RuntimeError};

#[runmat_macros::register_gpu_spec(builtin_path = "crate::builtins::math::rounding")]
pub const GPU_SPEC: BuiltinGpuSpec = BuiltinGpuSpec {
    name: "mod",
    op_kind: GpuOpKind::Elementwise,
    supported_precisions: &[ScalarType::F32, ScalarType::F64],
    broadcast: BroadcastSemantics::Matlab,
    provider_hooks: &[
        ProviderHook::Binary {
            name: "elem_div",
            commutative: false,
        },
        ProviderHook::Unary { name: "unary_floor" },
        ProviderHook::Binary {
            name: "elem_mul",
            commutative: false,
        },
        ProviderHook::Binary {
            name: "elem_sub",
            commutative: false,
        },
    ],
    constant_strategy: ConstantStrategy::InlineLiteral,
    residency: ResidencyPolicy::NewHandle,
    nan_mode: ReductionNaN::Include,
    two_pass_threshold: None,
    workgroup_size: None,
    accepts_nan_mode: false,
    notes:
        "Native integer providers may execute exact mod directly; floating fallback gathers and reuploads when a dedicated semantically complete provider hook is unavailable.",
};

#[runmat_macros::register_fusion_spec(builtin_path = "crate::builtins::math::rounding")]
pub const FUSION_SPEC: BuiltinFusionSpec = BuiltinFusionSpec {
    name: "mod",
    shape: ShapeRequirements::BroadcastCompatible,
    constant_strategy: ConstantStrategy::InlineLiteral,
    elementwise: Some(FusionKernelTemplate {
        scalar_precisions: &[ScalarType::F32, ScalarType::F64],
        wgsl_body: |ctx: &FusionExprContext| {
            let a = ctx
                .inputs
                .first()
                .ok_or(FusionError::MissingInput(0))?;
            let b = ctx.inputs.get(1).ok_or(FusionError::MissingInput(1))?;
            Ok(format!(
                "select({a} - {b} * floor({a} / {b}), {a}, {b} == 0.0)"
            ))
        },
    }),
    reduction: None,
    emits_nan: true,
    notes: "Fusion applies a - b * floor(a / b), including the documented mod(a, 0) = a convention; providers may substitute specialised kernels when available.",
};

const BUILTIN_NAME: &str = "mod";

const MOD_OUTPUT: [BuiltinParamDescriptor; 1] = [BuiltinParamDescriptor {
    name: "R",
    ty: BuiltinParamType::NumericArray,
    arity: BuiltinParamArity::Required,
    default: None,
    description: "Element-wise modulus result.",
}];
const MOD_INPUTS: [BuiltinParamDescriptor; 2] = [
    BuiltinParamDescriptor {
        name: "A",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Real dividend input (numeric/logical/char).",
    },
    BuiltinParamDescriptor {
        name: "B",
        ty: BuiltinParamType::Any,
        arity: BuiltinParamArity::Required,
        default: None,
        description: "Real divisor input (numeric/logical/char).",
    },
];
const MOD_SIGNATURES: [BuiltinSignatureDescriptor; 1] = [BuiltinSignatureDescriptor {
    label: "R = mod(A, B)",
    inputs: &MOD_INPUTS,
    outputs: &MOD_OUTPUT,
}];
const MOD_ERROR_INVALID_INPUT: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.MOD.INVALID_INPUT",
    identifier: Some("RunMat:mod:InvalidInput"),
    when: "Inputs are complex or cannot be interpreted as real numeric, logical, or char operands.",
    message: "mod: invalid input",
};
const MOD_ERROR_SIZE_MISMATCH: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.MOD.SIZE_MISMATCH",
    identifier: Some("RunMat:mod:SizeMismatch"),
    when: "Operands are not broadcast-compatible.",
    message: "mod: array sizes are not compatible for broadcasting",
};
const MOD_ERROR_INTERNAL: BuiltinErrorDescriptor = BuiltinErrorDescriptor {
    code: "RM.MOD.INTERNAL",
    identifier: Some("RunMat:mod:Internal"),
    when: "Internal tensor conversion, allocation, or provider composition failed.",
    message: "mod: internal error",
};
const MOD_ERRORS: [BuiltinErrorDescriptor; 3] = [
    MOD_ERROR_INVALID_INPUT,
    MOD_ERROR_SIZE_MISMATCH,
    MOD_ERROR_INTERNAL,
];
pub const MOD_DESCRIPTOR: BuiltinDescriptor = BuiltinDescriptor {
    signatures: &MOD_SIGNATURES,
    output_mode: BuiltinOutputMode::Fixed,
    completion_policy: BuiltinCompletionPolicy::Public,
    errors: &MOD_ERRORS,
};

const MOD_INTEGER_INPUTS: [BuiltinIntegerInputCapability; 2] = [
    BuiltinIntegerInputCapability {
        name: "A",
        classes: &crate::builtins::common::integer_capability::ALL_INTEGER_CLASSES,
        availability: BuiltinIntegerInputAvailability::Documented,
        scalar_double: BuiltinIntegerScalarDoubleRule::AllowedExceptWith64BitInteger,
        notes: "A is an integer array or a compatible real scalar double.",
    },
    BuiltinIntegerInputCapability {
        name: "B",
        classes: &crate::builtins::common::integer_capability::ALL_INTEGER_CLASSES,
        availability: BuiltinIntegerInputAvailability::Documented,
        scalar_double: BuiltinIntegerScalarDoubleRule::AllowedExceptWith64BitInteger,
        notes: "B is an integer array or a compatible real scalar double.",
    },
];

pub const INTEGER_CAPABILITIES: [BuiltinIntegerCapabilityDescriptor; 1] =
    [BuiltinIntegerCapabilityDescriptor {
        form: "R = mod(A, B)",
        inputs: &MOD_INTEGER_INPUTS,
        computation_domain: BuiltinIntegerComputationDomain::ExactInteger,
        output_class: BuiltinIntegerOutputClassRule::PreserveNondoubleInput,
        overflow: BuiltinIntegerOverflowRule::NotApplicable,
        backend: BuiltinIntegerBackendRule::HostAndGpu,
        overload: BuiltinIntegerOverloadKind::BroadcastCompatible,
        notes: "Native integer storage uses exact floor-remainder semantics; providers may execute exact integer mod directly or gather for a semantically complete fallback.",
    }];

fn mod_error_with_detail(
    error: &'static BuiltinErrorDescriptor,
    detail: impl AsRef<str>,
) -> RuntimeError {
    mod_error_with_message(format!("{}: {}", error.message, detail.as_ref()), error)
}

fn mod_error_with_message(
    message: impl Into<String>,
    error: &'static BuiltinErrorDescriptor,
) -> RuntimeError {
    let mut builder = build_runtime_error(message).with_builtin(BUILTIN_NAME);
    if let Some(identifier) = error.identifier {
        builder = builder.with_identifier(identifier);
    }
    builder.build()
}

#[runtime_builtin(
    name = "mod",
    category = "math/rounding",
    summary = "MATLAB-compatible real modulus a - b .* floor(a./b) with broadcasting.",
    keywords = "mod,modulus,remainder,gpu",
    accel = "binary",
    type_resolver(numeric_binary_type),
    descriptor(crate::builtins::math::rounding::MOD_DESCRIPTOR),
    integer_capabilities(crate::builtins::math::rounding::INTEGER_CAPABILITIES),
    builtin_path = "crate::builtins::math::rounding"
)]
async fn mod_builtin(lhs: Value, rhs: Value) -> BuiltinResult<Value> {
    if matches!(&lhs, Value::Complex(_, _) | Value::ComplexTensor(_))
        || matches!(&rhs, Value::Complex(_, _) | Value::ComplexTensor(_))
    {
        return Err(mod_error_with_detail(
            &MOD_ERROR_INVALID_INPUT,
            "inputs must be real",
        ));
    }
    crate::builtins::common::validation::reject_typed_complex_integer(&lhs, BUILTIN_NAME)?;
    crate::builtins::common::validation::reject_typed_complex_integer(&rhs, BUILTIN_NAME)?;
    reject_integer_logical_operands(&lhs, &rhs, BUILTIN_NAME)
        .map_err(|error| mod_error_with_detail(&MOD_ERROR_INVALID_INPUT, error))?;
    match (lhs, rhs) {
        (Value::GpuTensor(a), Value::GpuTensor(b)) => mod_gpu_pair(a, b).await,
        (Value::GpuTensor(a), other) => {
            let gathered = gpu_helpers::gather_tensor_async(&a).await?;
            mod_host(Value::Tensor(gathered), other)
        }
        (other, Value::GpuTensor(b)) => {
            let gathered = gpu_helpers::gather_tensor_async(&b).await?;
            mod_host(other, Value::Tensor(gathered))
        }
        (left, right) => mod_host(left, right),
    }
}

async fn mod_gpu_pair(a: GpuTensorHandle, b: GpuTensorHandle) -> BuiltinResult<Value> {
    if runmat_accelerate_api::handle_integer_type(&a).is_some()
        && runmat_accelerate_api::handle_integer_type(&b).is_some()
        && a.device_id == b.device_id
    {
        if let Some(provider) = runmat_accelerate_api::provider_for_handle(&a) {
            if let Ok(out) = provider.elem_mod(&a, &b).await {
                return Ok(gpu_helpers::resident_gpu_value(out));
            }
        }
    }
    let left = gpu_helpers::gather_tensor_async(&a).await?;
    let right = gpu_helpers::gather_tensor_async(&b).await?;
    let result = mod_host(Value::Tensor(left), Value::Tensor(right))?;
    if a.device_id == b.device_id {
        if let Some(provider) = runmat_accelerate_api::provider_for_handle(&a) {
            return upload_mod_result(provider, result);
        }
    }
    Ok(result)
}

fn upload_mod_result(
    provider: &dyn runmat_accelerate_api::AccelProvider,
    value: Value,
) -> BuiltinResult<Value> {
    let tensor = match value {
        Value::Tensor(tensor) => tensor,
        Value::Num(value) => Tensor::new(vec![value], vec![1, 1])
            .map_err(|err| mod_error_with_detail(&MOD_ERROR_INTERNAL, err))?,
        other => return Ok(other),
    };
    let handle = gpu_helpers::upload_tensor(provider, &tensor)
        .map_err(|err| mod_error_with_detail(&MOD_ERROR_INTERNAL, err))?;
    Ok(gpu_helpers::resident_gpu_value(handle))
}

fn mod_host(lhs: Value, rhs: Value) -> BuiltinResult<Value> {
    if let Some(result) = try_integer_remainder(&lhs, &rhs, IntegerRemainderOp::Mod, BUILTIN_NAME)
        .map_err(|error| mod_error_with_detail(&MOD_ERROR_INVALID_INPUT, error))?
    {
        return Ok(result);
    }
    if let Some(result) = scalar_mod_value(&lhs, &rhs) {
        return Ok(result);
    }
    let left = value_into_numeric_array(lhs)?;
    let right = value_into_numeric_array(rhs)?;
    match align_numeric_arrays(left, right)? {
        NumericPair::Real(a, b) => compute_mod_real(&a, &b),
        NumericPair::Complex(a, b) => compute_mod_complex(&a, &b),
    }
}

fn compute_mod_real(a: &Tensor, b: &Tensor) -> BuiltinResult<Value> {
    let plan = BroadcastPlan::new(&a.shape, &b.shape)
        .map_err(|err| mod_error_with_detail(&MOD_ERROR_SIZE_MISMATCH, err))?;
    let dtype = if a.numeric_dtype() == NumericDType::F32 && b.numeric_dtype() == NumericDType::F32
    {
        NumericDType::F32
    } else {
        NumericDType::F64
    };
    if plan.is_empty() {
        let tensor = Tensor::new_with_dtype(Vec::new(), plan.output_shape().to_vec(), dtype)
            .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
        return Ok(tensor::tensor_into_value(tensor));
    }
    let mut result = vec![0.0f64; plan.len()];
    for (out_idx, idx_a, idx_b) in plan.iter() {
        let aval = tensor::tensor_value_f64(a, idx_a);
        let bval = tensor::tensor_value_f64(b, idx_b);
        result[out_idx] = mod_real_scalar(aval, bval);
    }
    let tensor = Tensor::new_with_dtype(result, plan.output_shape().to_vec(), dtype)
        .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
    Ok(tensor::tensor_into_value(tensor))
}

fn compute_mod_complex(a: &ComplexTensor, b: &ComplexTensor) -> BuiltinResult<Value> {
    let plan = BroadcastPlan::new(&a.shape, &b.shape)
        .map_err(|err| mod_error_with_detail(&MOD_ERROR_SIZE_MISMATCH, err))?;
    if plan.is_empty() {
        let tensor = ComplexTensor::new(Vec::new(), plan.output_shape().to_vec())
            .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
        return Ok(complex_tensor_into_value(tensor));
    }
    let mut result = vec![(0.0f64, 0.0f64); plan.len()];
    for (out_idx, idx_a, idx_b) in plan.iter() {
        let (ar, ai) = a.materialize_f64()[idx_a];
        let (br, bi) = b.materialize_f64()[idx_b];
        result[out_idx] = mod_complex_scalar(ar, ai, br, bi);
    }
    let tensor = ComplexTensor::new(result, plan.output_shape().to_vec())
        .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
    Ok(complex_tensor_into_value(tensor))
}

fn mod_real_scalar(a: f64, b: f64) -> f64 {
    if a.is_nan() || b.is_nan() {
        return f64::NAN;
    }
    if b == 0.0 {
        return a;
    }
    if !a.is_finite() && b.is_finite() {
        return f64::NAN;
    }
    let quotient = (a / b).floor();
    let mut remainder = a - b * quotient;
    if remainder == 0.0 {
        remainder = 0.0;
    }
    if b.is_infinite() && a.is_finite() {
        // MATLAB sign-correction: mod(a, ±Inf) returns a when signs match, ±Inf otherwise.
        if a == 0.0 {
            return 0.0;
        }
        return if a.signum() == b.signum() { a } else { b };
    }
    if !remainder.is_finite() && !a.is_finite() {
        return f64::NAN;
    }
    let same_sign = remainder == 0.0 || remainder.signum() == b.signum();
    if !same_sign {
        remainder += b;
    }
    if remainder == -0.0 {
        remainder = 0.0;
    }
    remainder
}

fn mod_complex_scalar(ar: f64, ai: f64, br: f64, bi: f64) -> (f64, f64) {
    if (ar.is_nan() || ai.is_nan()) || (br.is_nan() || bi.is_nan()) {
        return (f64::NAN, f64::NAN);
    }
    if br == 0.0 && bi == 0.0 {
        return (f64::NAN, f64::NAN);
    }
    if !ar.is_finite() || !ai.is_finite() {
        return (f64::NAN, f64::NAN);
    }
    let (qr, qi) = complex_div(ar, ai, br, bi);
    if !qr.is_finite() && !qi.is_finite() && br.is_finite() && bi.is_finite() {
        return (f64::NAN, f64::NAN);
    }
    let (fr, fi) = (qr.floor(), qi.floor());
    let (mulr, muli) = complex_mul(br, bi, fr, fi);
    let (rr, ri) = (ar - mulr, ai - muli);
    (normalize_zero(rr), normalize_zero(ri))
}

fn scalar_real_value(value: &Value) -> Option<f64> {
    match value {
        Value::Num(n) => Some(*n),
        Value::Int(i) => Some(i.to_f64()),
        Value::Bool(b) => Some(if *b { 1.0 } else { 0.0 }),
        Value::Tensor(t) if tensor::is_scalar_tensor(t) => Some(tensor::tensor_value_f64(t, 0)),
        Value::LogicalArray(l) if l.data.len() == 1 => Some(if l.data[0] != 0 { 1.0 } else { 0.0 }),
        Value::CharArray(ca) if ca.rows * ca.cols == 1 => {
            Some(ca.data.first().map(|&ch| ch as u32 as f64).unwrap_or(0.0))
        }
        _ => None,
    }
}

fn scalar_complex_value(value: &Value) -> Option<(f64, f64)> {
    match value {
        Value::Complex(re, im) => Some((*re, *im)),
        Value::ComplexTensor(ct) if tensor::complex_tensor_element_len(ct) == 1 => {
            let value = tensor::complex_tensor_value_complex64(ct, 0);
            Some((value.re, value.im))
        }
        _ => None,
    }
}

fn scalar_mod_value(lhs: &Value, rhs: &Value) -> Option<Value> {
    let left = scalar_complex_value(lhs).or_else(|| scalar_real_value(lhs).map(|v| (v, 0.0)))?;
    let right = scalar_complex_value(rhs).or_else(|| scalar_real_value(rhs).map(|v| (v, 0.0)))?;
    let (ar, ai) = left;
    let (br, bi) = right;
    if ai != 0.0 || bi != 0.0 {
        let (re, im) = mod_complex_scalar(ar, ai, br, bi);
        return Some(Value::Complex(re, im));
    }
    Some(Value::Num(mod_real_scalar(ar, br)))
}

fn normalize_zero(value: f64) -> f64 {
    if value == -0.0 {
        0.0
    } else {
        value
    }
}

fn complex_mul(ar: f64, ai: f64, br: f64, bi: f64) -> (f64, f64) {
    (ar * br - ai * bi, ar * bi + ai * br)
}

fn complex_div(ar: f64, ai: f64, br: f64, bi: f64) -> (f64, f64) {
    let denom = br * br + bi * bi;
    if denom == 0.0 {
        return (f64::NAN, f64::NAN);
    }
    ((ar * br + ai * bi) / denom, (ai * br - ar * bi) / denom)
}

fn complex_tensor_into_value(tensor: ComplexTensor) -> Value {
    if tensor::complex_tensor_element_len(&tensor) == 1 {
        let value = tensor::complex_tensor_value_complex64(&tensor, 0);
        Value::Complex(value.re, value.im)
    } else {
        Value::ComplexTensor(tensor)
    }
}

fn value_into_numeric_array(value: Value) -> BuiltinResult<NumericArray> {
    match value {
        Value::Complex(re, im) => {
            let tensor = ComplexTensor::new(vec![(re, im)], vec![1, 1])
                .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
            Ok(NumericArray::Complex(tensor))
        }
        Value::ComplexTensor(ct) => Ok(NumericArray::Complex(ct)),
        Value::CharArray(ca) => {
            let data: Vec<f64> = ca.data.iter().map(|&ch| ch as u32 as f64).collect();
            let tensor = Tensor::new(data, vec![ca.rows, ca.cols])
                .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?;
            Ok(NumericArray::Real(tensor))
        }
        Value::String(_) | Value::StringArray(_) => Err(mod_error_with_detail(
            &MOD_ERROR_INVALID_INPUT,
            "expected numeric input, got string",
        )),
        Value::GpuTensor(_) => Err(mod_error_with_detail(
            &MOD_ERROR_INTERNAL,
            "internal error converting GPU tensor",
        )),
        other => {
            let tensor = tensor::value_into_tensor_for(BUILTIN_NAME, other)
                .map_err(|err| mod_error_with_detail(&MOD_ERROR_INVALID_INPUT, err))?;
            Ok(NumericArray::Real(tensor))
        }
    }
}

enum NumericArray {
    Real(Tensor),
    Complex(ComplexTensor),
}

enum NumericPair {
    Real(Tensor, Tensor),
    Complex(ComplexTensor, ComplexTensor),
}

fn align_numeric_arrays(lhs: NumericArray, rhs: NumericArray) -> BuiltinResult<NumericPair> {
    match (lhs, rhs) {
        (NumericArray::Real(a), NumericArray::Real(b)) => Ok(NumericPair::Real(a, b)),
        (left, right) => {
            let lc = into_complex(left)?;
            let rc = into_complex(right)?;
            Ok(NumericPair::Complex(lc, rc))
        }
    }
}

fn into_complex(input: NumericArray) -> BuiltinResult<ComplexTensor> {
    match input {
        NumericArray::Real(t) => {
            let shape = t.shape.clone();
            let complex = t
                .into_numeric_storage()
                .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))?
                .materialize_f64()
                .into_iter()
                .map(|real| (real, 0.0))
                .collect();
            ComplexTensor::new(complex, shape)
                .map_err(|e| mod_error_with_detail(&MOD_ERROR_INTERNAL, e))
        }
        NumericArray::Complex(ct) => Ok(ct),
    }
}

#[cfg(test)]
pub(crate) mod tests {
    use super::*;
    use crate::builtins::common::test_support;
    use crate::RuntimeError;
    use futures::executor::block_on;
    use runmat_builtins::{ResolveContext, Type};
    use runmat_value::{CharArray, ComplexTensor, IntValue, IntegerStorage, LogicalArray, Tensor};

    fn mod_builtin(lhs: Value, rhs: Value) -> BuiltinResult<Value> {
        block_on(super::mod_builtin(lhs, rhs))
    }

    #[test]
    fn mod_real_arrays_preserve_native_single_storage_including_empty() {
        let lhs = Tensor::from_f32(vec![5.5, -5.5], vec![1, 2]).unwrap();
        let rhs = Tensor::from_f32(vec![2.0, 2.0], vec![1, 2]).unwrap();
        let output = compute_mod_real(&lhs, &rhs).unwrap();
        let Value::Tensor(output) = output else {
            panic!("expected native-single tensor")
        };
        assert_eq!(
            output.into_numeric_storage().unwrap(),
            runmat_value::NumericStorage::F32(vec![1.5, 0.5])
        );

        let lhs = Tensor::from_f32(Vec::new(), vec![0, 2]).unwrap();
        let rhs = Tensor::from_f32(Vec::new(), vec![0, 2]).unwrap();
        let Value::Tensor(output) = compute_mod_real(&lhs, &rhs).unwrap() else {
            panic!("expected empty native-single tensor")
        };
        assert_eq!(
            output.into_numeric_storage().unwrap(),
            runmat_value::NumericStorage::F32(Vec::new())
        );
    }

    fn assert_error_contains(error: RuntimeError, needle: &str) {
        assert!(
            error.message().contains(needle),
            "unexpected error: {}",
            error.message()
        );
    }

    #[test]
    fn mod_descriptor_signatures_cover_core_forms() {
        let labels: Vec<&str> = MOD_DESCRIPTOR
            .signatures
            .iter()
            .map(|sig| sig.label)
            .collect();
        assert!(labels.contains(&"R = mod(A, B)"));
    }

    #[test]
    fn mod_type_preserves_tensor_shape() {
        let out = numeric_binary_type(
            &[
                Type::Tensor {
                    shape: Some(vec![Some(2), Some(3)]),
                },
                Type::Tensor {
                    shape: Some(vec![Some(2), Some(3)]),
                },
            ],
            &ResolveContext::new(Vec::new()),
        );
        assert_eq!(
            out,
            Type::Tensor {
                shape: Some(vec![Some(2), Some(3)])
            }
        );
    }

    #[test]
    fn mod_type_scalar_and_tensor_returns_tensor() {
        let out = numeric_binary_type(
            &[
                Type::Num,
                Type::Tensor {
                    shape: Some(vec![Some(4), Some(1)]),
                },
            ],
            &ResolveContext::new(Vec::new()),
        );
        assert_eq!(
            out,
            Type::Tensor {
                shape: Some(vec![Some(4), Some(1)])
            }
        );
    }

    #[test]
    fn mod_type_scalar_returns_num() {
        let out = numeric_binary_type(&[Type::Num, Type::Int], &ResolveContext::new(Vec::new()));
        assert_eq!(out, Type::Num);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_positive_values() {
        let result = mod_builtin(Value::Num(17.0), Value::Num(5.0)).expect("mod");
        match result {
            Value::Num(v) => assert!((v - 2.0).abs() < 1e-12),
            other => panic!("expected scalar result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_negative_divisor_keeps_sign() {
        let tensor = Tensor::new(vec![-7.0, -3.0, 4.0, 9.0], vec![4, 1]).unwrap();
        let divisor = Tensor::new(vec![-4.0], vec![1, 1]).unwrap();
        let result =
            mod_builtin(Value::Tensor(tensor), Value::Tensor(divisor)).expect("mod broadcast");
        match result {
            Value::Tensor(out) => {
                assert_eq!(out.materialize_f64(), vec![-3.0, -3.0, 0.0, -3.0]);
            }
            other => panic!("expected tensor result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_negative_numerator_positive_divisor() {
        let result = mod_builtin(Value::Num(-3.0), Value::Num(2.0)).expect("mod");
        match result {
            Value::Num(v) => assert!((v - 1.0).abs() < 1e-12),
            other => panic!("expected scalar result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_zero_divisor_returns_dividend() {
        let result = mod_builtin(Value::Num(3.0), Value::Num(0.0)).expect("mod");
        match result {
            Value::Num(v) => assert_eq!(v, 3.0),
            other => panic!("expected dividend, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_matrix_scalar_broadcast() {
        let matrix = Tensor::new(vec![4.5, 7.1, -2.3, 0.4], vec![2, 2]).unwrap();
        let result = mod_builtin(Value::Tensor(matrix), Value::Num(2.0)).expect("broadcast");
        match result {
            Value::Tensor(t) => {
                assert_eq!(t.shape, vec![2, 2]);
                let expected = [0.5, 1.1, 1.7, 0.4];
                for (a, b) in t.materialize_f64().iter().zip(expected.iter()) {
                    assert!((a - b).abs() < 1e-12);
                }
            }
            other => panic!("expected tensor result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_rejects_complex_operands() {
        let complex =
            ComplexTensor::new(vec![(3.0, 4.0), (-2.0, 5.0)], vec![1, 2]).expect("complex tensor");
        let divisor = ComplexTensor::new(vec![(2.0, 1.0)], vec![1, 1]).expect("divisor");
        assert!(mod_builtin(Value::ComplexTensor(complex), Value::ComplexTensor(divisor)).is_err());
        assert!(mod_builtin(Value::Complex(2.0, 0.0), Value::Num(1.0)).is_err());
    }

    #[test]
    fn mod_complex_scalar_helpers_read_typed_integer_complex_storage_without_mirror() {
        let complex = ComplexTensor::new_integer(
            runmat_value::IntegerComplexStorage::new(
                IntegerStorage::I16(vec![9]),
                IntegerStorage::I16(vec![-2]),
            )
            .expect("integer complex storage"),
            vec![1, 1],
        )
        .expect("integer complex tensor");

        assert_eq!(
            scalar_complex_value(&Value::ComplexTensor(complex.clone())),
            Some((9.0, -2.0))
        );
        assert_eq!(
            complex_tensor_into_value(complex),
            Value::Complex(9.0, -2.0)
        );
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_char_array_support() {
        let chars = CharArray::new("ABC".chars().collect(), 1, 3).unwrap();
        let result = mod_builtin(Value::CharArray(chars), Value::Num(5.0)).expect("mod");
        match result {
            Value::Tensor(t) => assert_eq!(t.materialize_f64(), vec![0.0, 1.0, 2.0]),
            other => panic!("expected tensor result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_string_input_errors() {
        let err = mod_builtin(Value::from("abc"), Value::Num(3.0))
            .expect_err("string inputs should error");
        let identifier = err.identifier().map(str::to_string);
        assert_error_contains(err, "expected numeric input");
        assert_eq!(identifier.as_deref(), MOD_ERROR_INVALID_INPUT.identifier);
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_logical_array_support() {
        let logical = LogicalArray::new(vec![1, 0, 1, 0], vec![2, 2]).unwrap();
        let value =
            mod_builtin(Value::LogicalArray(logical), Value::Num(2.0)).expect("logical mod");
        match value {
            Value::Tensor(t) => assert_eq!(t.materialize_f64(), vec![1.0, 0.0, 1.0, 0.0]),
            other => panic!("expected tensor result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_vector_broadcasting() {
        let lhs = Tensor::new(vec![1.0, 2.0], vec![2, 1]).unwrap();
        let rhs = Tensor::new(vec![3.0, 4.0, 5.0], vec![1, 3]).unwrap();
        let result = mod_builtin(Value::Tensor(lhs), Value::Tensor(rhs)).expect("vector broadcast");
        match result {
            Value::Tensor(t) => {
                assert_eq!(t.shape, vec![2, 3]);
                assert_eq!(t.materialize_f64(), vec![1.0, 2.0, 1.0, 2.0, 1.0, 2.0]);
            }
            other => panic!("expected tensor result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_nan_inputs_propagate() {
        let result = mod_builtin(Value::Num(f64::NAN), Value::Num(3.0)).expect("mod");
        match result {
            Value::Num(v) => assert!(v.is_nan()),
            other => panic!("expected NaN result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_gpu_pair_roundtrip() {
        test_support::with_test_provider(|provider| {
            let tensor = Tensor::new(vec![-5.0, -3.0, 0.0, 1.0, 6.0, 9.0], vec![3, 2]).unwrap();
            let divisor = Tensor::new(vec![4.0, 4.0, 4.0, 4.0, 4.0, 4.0], vec![3, 2]).unwrap();
            let a_view = runmat_accelerate_api::HostTensorView {
                data: &tensor.materialize_f64(),
                shape: &tensor.shape,
            };
            let b_view = runmat_accelerate_api::HostTensorView {
                data: &divisor.materialize_f64(),
                shape: &divisor.shape,
            };
            let a_handle = provider.upload(&a_view).expect("upload a");
            let b_handle = provider.upload(&b_view).expect("upload b");
            let result =
                mod_builtin(Value::GpuTensor(a_handle), Value::GpuTensor(b_handle)).expect("mod");
            let gathered = test_support::gather(result).expect("gather result");
            assert_eq!(gathered.shape, vec![3, 2]);
            assert_eq!(
                gathered.materialize_f64(),
                vec![3.0, 1.0, 0.0, 1.0, 2.0, 1.0]
            );
        });
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_int_scalar_preserves_exact_class() {
        let result =
            mod_builtin(Value::Int(IntValue::I32(-7)), Value::Int(IntValue::I32(4))).expect("mod");
        match result {
            Value::Int(IntValue::I32(v)) => assert_eq!(v, 1),
            other => panic!("expected int32 scalar result, got {other:?}"),
        }
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    fn mod_scalar_fast_path_reads_typed_integer_storage_exactly() {
        let lhs =
            Tensor::new_integer(IntegerStorage::I16(vec![7]), vec![1, 1]).expect("lhs tensor");

        assert_eq!(scalar_real_value(&Value::Tensor(lhs.clone())), Some(7.0));

        let result = mod_builtin(Value::Tensor(lhs), Value::Num(4.0)).expect("mod");
        match result {
            Value::Int(IntValue::I16(v)) => assert_eq!(v, 3),
            other => panic!("expected int16 scalar result, got {other:?}"),
        }
    }

    #[test]
    fn mod_dense_integer_arrays_preserve_exact_storage_without_mirror() {
        let lhs = Tensor::new_integer(IntegerStorage::I64(vec![-7, 7]), vec![2, 1]).expect("lhs");
        let rhs =
            Tensor::new_integer(IntegerStorage::I64(vec![4, -4, 0]), vec![1, 3]).expect("rhs");

        let result = mod_builtin(Value::Tensor(lhs), Value::Tensor(rhs)).expect("mod");
        let Value::Tensor(result) = result else {
            panic!("expected integer tensor");
        };
        assert_eq!(result.shape, vec![2, 3]);
        assert_eq!(
            result.integer_storage(),
            Some(&IntegerStorage::I64(vec![1, 3, -3, -1, -7, 7]))
        );

        let lhs =
            Tensor::new_integer(IntegerStorage::U64(vec![u64::MAX]), vec![1, 1]).expect("lhs");
        assert_eq!(
            mod_builtin(Value::Tensor(lhs), Value::Num(3.0)).expect("mod"),
            Value::Int(IntValue::U64(0))
        );
    }

    #[cfg_attr(target_arch = "wasm32", wasm_bindgen_test::wasm_bindgen_test)]
    #[test]
    #[cfg(feature = "wgpu")]
    fn mod_wgpu_matches_cpu() {
        let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
            runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
        );
        let numer = Tensor::new(vec![-5.0, -3.25, 0.0, 1.75, 6.5, 9.0], vec![3, 2]).unwrap();
        let denom = Tensor::new(vec![4.0, -2.5, 3.0, 3.0, 2.0, -5.0], vec![3, 2]).unwrap();
        let cpu_value =
            mod_host(Value::Tensor(numer.clone()), Value::Tensor(denom.clone())).expect("cpu mod");

        let provider = runmat_accelerate_api::provider().expect("wgpu provider registered");
        let numer_handle = provider
            .upload(&runmat_accelerate_api::HostTensorView {
                data: &numer.materialize_f64(),
                shape: &numer.shape,
            })
            .expect("upload numer");
        let denom_handle = provider
            .upload(&runmat_accelerate_api::HostTensorView {
                data: &denom.materialize_f64(),
                shape: &denom.shape,
            })
            .expect("upload denom");

        let gpu_value = block_on(mod_gpu_pair(numer_handle, denom_handle)).expect("gpu mod");
        let gpu_tensor = test_support::gather(gpu_value).expect("gather gpu result");

        let cpu_tensor = match cpu_value {
            Value::Tensor(t) => t,
            Value::Num(n) => Tensor::new(vec![n], vec![1, 1]).expect("scalar tensor"),
            other => panic!("unexpected CPU result {other:?}"),
        };

        assert_eq!(gpu_tensor.shape, cpu_tensor.shape);
        let tol = match provider.precision() {
            runmat_accelerate_api::ProviderPrecision::F64 => 1e-12,
            runmat_accelerate_api::ProviderPrecision::F32 => 1e-5,
        };
        for (gpu, cpu) in gpu_tensor
            .materialize_f64()
            .iter()
            .zip(cpu_tensor.materialize_f64().iter())
        {
            assert!(
                (gpu - cpu).abs() <= tol,
                "|{gpu} - {cpu}| exceeded tolerance {tol}"
            );
        }
    }
}