vector_ta/indicators/

mab.rs

#[cfg(feature = "python")]
use numpy::{IntoPyArray, PyArray1};
#[cfg(feature = "python")]
use pyo3::exceptions::PyValueError;
#[cfg(feature = "python")]
use pyo3::prelude::*;
#[cfg(feature = "python")]
use pyo3::types::PyDict;

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
use serde::{Deserialize, Serialize};
#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
use wasm_bindgen::prelude::*;

use crate::utilities::data_loader::{source_type, Candles};
use crate::utilities::enums::Kernel;
use crate::utilities::helpers::{
    alloc_with_nan_prefix, detect_best_batch_kernel, detect_best_kernel, init_matrix_prefixes,
    make_uninit_matrix,
};
#[cfg(feature = "python")]
use crate::utilities::kernel_validation::validate_kernel;
use aligned_vec::{AVec, CACHELINE_ALIGN};
#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
use core::arch::x86_64::*;
#[cfg(not(target_arch = "wasm32"))]
use rayon::prelude::*;
use std::convert::AsRef;
use std::error::Error;
use std::mem::MaybeUninit;
use thiserror::Error;

impl<'a> AsRef<[f64]> for MabInput<'a> {
    #[inline(always)]
    fn as_ref(&self) -> &[f64] {
        match &self.data {
            MabData::Slice(slice) => slice,
            MabData::Candles { candles, source } => source_type(candles, source),
        }
    }
}

#[derive(Debug, Clone)]
pub enum MabData<'a> {
    Candles {
        candles: &'a Candles,
        source: &'a str,
    },
    Slice(&'a [f64]),
}

#[derive(Debug, Clone)]
pub struct MabOutput {
    pub upperband: Vec<f64>,
    pub middleband: Vec<f64>,
    pub lowerband: Vec<f64>,
}

#[derive(Debug, Clone)]
#[cfg_attr(
    all(target_arch = "wasm32", feature = "wasm"),
    derive(Serialize, Deserialize)
)]
pub struct MabParams {
    pub fast_period: Option<usize>,
    pub slow_period: Option<usize>,
    pub devup: Option<f64>,
    pub devdn: Option<f64>,
    pub fast_ma_type: Option<String>,
    pub slow_ma_type: Option<String>,
}

impl Default for MabParams {
    fn default() -> Self {
        Self {
            fast_period: Some(10),
            slow_period: Some(50),
            devup: Some(1.0),
            devdn: Some(1.0),
            fast_ma_type: Some("sma".to_string()),
            slow_ma_type: Some("sma".to_string()),
        }
    }
}

#[derive(Debug, Clone)]
pub struct MabInput<'a> {
    pub data: MabData<'a>,
    pub params: MabParams,
}

impl<'a> MabInput<'a> {
    pub fn from_candles(candles: &'a Candles, source: &'a str, params: MabParams) -> Self {
        Self {
            data: MabData::Candles { candles, source },
            params,
        }
    }

    pub fn from_slice(slice: &'a [f64], params: MabParams) -> Self {
        Self {
            data: MabData::Slice(slice),
            params,
        }
    }

    pub fn with_default_params(data: MabData<'a>) -> Self {
        Self {
            data,
            params: MabParams::default(),
        }
    }

    pub fn with_default_candles(candles: &'a Candles) -> Self {
        Self::from_candles(candles, "close", MabParams::default())
    }

    pub fn get_fast_period(&self) -> usize {
        self.params.fast_period.unwrap_or(10)
    }

    pub fn get_slow_period(&self) -> usize {
        self.params.slow_period.unwrap_or(50)
    }

    pub fn get_devup(&self) -> f64 {
        self.params.devup.unwrap_or(1.0)
    }

    pub fn get_devdn(&self) -> f64 {
        self.params.devdn.unwrap_or(1.0)
    }

    pub fn get_fast_ma_type(&self) -> &str {
        self.params
            .fast_ma_type
            .as_ref()
            .map(|s| s.as_str())
            .unwrap_or("sma")
    }

    pub fn get_slow_ma_type(&self) -> &str {
        self.params
            .slow_ma_type
            .as_ref()
            .map(|s| s.as_str())
            .unwrap_or("sma")
    }
}

#[derive(Error, Debug)]
pub enum MabError {
    #[error("mab: Input data slice is empty.")]
    EmptyInputData,
    #[error("mab: All values are NaN.")]
    AllValuesNaN,
    #[error("mab: Invalid period: fast={fast} slow={slow} len={data_len}")]
    InvalidPeriod {
        fast: usize,
        slow: usize,
        data_len: usize,
    },
    #[error("mab: Not enough valid data: need={needed} valid={valid}")]
    NotEnoughValidData { needed: usize, valid: usize },
    #[error(
        "mab: Output length mismatch: expected={expected} upper={upper_len} middle={middle_len} lower={lower_len}"
    )]
    OutputLengthMismatch {
        upper_len: usize,
        middle_len: usize,
        lower_len: usize,
        expected: usize,
    },
    #[error("mab: Invalid range (start={start}, end={end}, step={step})")]
    InvalidRange {
        start: usize,
        end: usize,
        step: usize,
    },
    #[error("mab: Invalid range (f64) (start={start}, end={end}, step={step})")]
    InvalidRangeF64 { start: f64, end: f64, step: f64 },
    #[error("mab: non-batch kernel passed to batch path: {0:?}")]
    InvalidKernelForBatch(Kernel),
}

#[inline(always)]
fn mab_validate(input: &MabInput) -> Result<usize, MabError> {
    let data = input.as_ref();
    if data.is_empty() {
        return Err(MabError::EmptyInputData);
    }
    let fast = input.get_fast_period();
    let slow = input.get_slow_period();
    if fast == 0 || slow == 0 || fast > data.len() || slow > data.len() {
        return Err(MabError::InvalidPeriod {
            fast,
            slow,
            data_len: data.len(),
        });
    }

    let first_valid = data
        .iter()
        .position(|&x| !x.is_nan())
        .ok_or(MabError::AllValuesNaN)?;
    let max_period = fast.max(slow);
    if data.len() - first_valid < max_period {
        return Err(MabError::NotEnoughValidData {
            needed: max_period,
            valid: data.len() - first_valid,
        });
    }
    Ok(first_valid)
}

#[inline(always)]
fn mab_prepare<'a>(
    input: &'a MabInput,
    kernel: Kernel,
) -> Result<(&'a [f64], usize, Kernel, usize, usize, f64, f64), MabError> {
    let first_valid = mab_validate(input)?;
    let data = input.as_ref();
    let fast = input.get_fast_period();
    let slow = input.get_slow_period();
    let devup = input.get_devup();
    let devdn = input.get_devdn();
    let chosen = match kernel {
        Kernel::Auto => Kernel::Scalar,
        _ => kernel,
    };
    let warmup = first_valid + fast.max(slow) - 1;
    Ok((data, warmup, chosen, fast, slow, devup, devdn))
}

#[inline(always)]
fn mab_prepare2<'a>(
    input: &'a MabInput,
    kernel: Kernel,
) -> Result<(&'a [f64], Kernel, usize, usize, usize, f64, f64), MabError> {
    let data = input.as_ref();
    if data.is_empty() {
        return Err(MabError::EmptyInputData);
    }
    let fast = input.get_fast_period();
    let slow = input.get_slow_period();
    if fast == 0 || slow == 0 || fast > data.len() || slow > data.len() {
        return Err(MabError::InvalidPeriod {
            fast,
            slow,
            data_len: data.len(),
        });
    }
    let first = data
        .iter()
        .position(|&x| !x.is_nan())
        .ok_or(MabError::AllValuesNaN)?;
    let need = fast.max(slow);

    let need_total = need + fast - 1;
    if data.len() - first < need_total {
        return Err(MabError::NotEnoughValidData {
            needed: need_total,
            valid: data.len() - first,
        });
    }
    let chosen = match kernel {
        Kernel::Auto => Kernel::Scalar,
        k => k,
    };
    let warmup = first + need_total - 1;
    Ok((
        data,
        chosen,
        first,
        warmup,
        fast,
        input.get_devup(),
        input.get_devdn(),
    ))
}

pub fn mab_into_slice(
    upper_dst: &mut [f64],
    middle_dst: &mut [f64],
    lower_dst: &mut [f64],
    input: &MabInput,
    kern: Kernel,
) -> Result<(), MabError> {
    use crate::indicators::ema::{ema, EmaInput, EmaParams};
    use crate::indicators::sma::{sma, SmaInput, SmaParams};

    let (data, chosen, first, warmup, fast_period, devup, devdn) = mab_prepare2(input, kern)?;
    let slow_period = input.get_slow_period();

    let n = data.len();
    if upper_dst.len() != n || middle_dst.len() != n || lower_dst.len() != n {
        return Err(MabError::OutputLengthMismatch {
            upper_len: upper_dst.len(),
            middle_len: middle_dst.len(),
            lower_len: lower_dst.len(),
            expected: n,
        });
    }

    let fast_ma_type = input.get_fast_ma_type();
    let slow_ma_type = input.get_slow_ma_type();

    let fast_ma = match fast_ma_type {
        "ema" => {
            let params = EmaParams {
                period: Some(fast_period),
            };
            ema(&EmaInput::from_slice(data, params))
                .map_err(|_| MabError::NotEnoughValidData {
                    needed: fast_period,
                    valid: n - first,
                })?
                .values
        }
        _ => {
            let params = SmaParams {
                period: Some(fast_period),
            };
            sma(&SmaInput::from_slice(data, params))
                .map_err(|_| MabError::NotEnoughValidData {
                    needed: fast_period,
                    valid: n - first,
                })?
                .values
        }
    };

    let slow_ma = match slow_ma_type {
        "ema" => {
            let params = EmaParams {
                period: Some(slow_period),
            };
            ema(&EmaInput::from_slice(data, params))
                .map_err(|_| MabError::NotEnoughValidData {
                    needed: slow_period,
                    valid: n - first,
                })?
                .values
        }
        _ => {
            let params = SmaParams {
                period: Some(slow_period),
            };
            sma(&SmaInput::from_slice(data, params))
                .map_err(|_| MabError::NotEnoughValidData {
                    needed: slow_period,
                    valid: n - first,
                })?
                .values
        }
    };

    let warmup_end = (warmup + 1).min(n);
    for dst in [&mut *upper_dst, &mut *middle_dst, &mut *lower_dst] {
        for v in &mut dst[..warmup_end] {
            *v = f64::NAN;
        }
    }

    mab_compute_into(
        &fast_ma,
        &slow_ma,
        fast_period,
        devup,
        devdn,
        warmup + 1,
        chosen,
        upper_dst,
        middle_dst,
        lower_dst,
    );

    Ok(())
}

#[cfg(not(all(target_arch = "wasm32", feature = "wasm")))]
pub fn mab_into(
    input: &MabInput,
    upper_dst: &mut [f64],
    middle_dst: &mut [f64],
    lower_dst: &mut [f64],
) -> Result<(), MabError> {
    mab_into_slice(upper_dst, middle_dst, lower_dst, input, Kernel::Auto)
}

pub fn mab(input: &MabInput) -> Result<MabOutput, MabError> {
    let (_data, _warmup, kernel, _fast, _slow, _devup, _devdn) = mab_prepare(input, Kernel::Auto)?;
    mab_with_kernel(input, kernel)
}

pub fn mab_with_kernel(input: &MabInput, kernel: Kernel) -> Result<MabOutput, MabError> {
    let data = input.as_ref();
    let (_, _, _, warmup, _, _, _) = mab_prepare2(input, kernel)?;

    let mut upperband = alloc_with_nan_prefix(data.len(), warmup);
    let mut middleband = alloc_with_nan_prefix(data.len(), warmup);
    let mut lowerband = alloc_with_nan_prefix(data.len(), warmup);

    mab_into_slice(
        &mut upperband,
        &mut middleband,
        &mut lowerband,
        input,
        kernel,
    )?;

    Ok(MabOutput {
        upperband,
        middleband,
        lowerband,
    })
}

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

    fn eq_or_both_nan(a: f64, b: f64) -> bool {
        (a.is_nan() && b.is_nan()) || (a == b)
    }

    #[test]
    fn test_mab_into_matches_api() {
        let n = 256usize;
        let mut data = vec![f64::NAN; n];
        for i in 5..n {
            data[i] = (i as f64).sin() * 0.5 + (i as f64).cos() * 0.25;
        }

        let input = MabInput::from_slice(&data, MabParams::default());

        let base = mab(&input).expect("mab baseline should succeed");

        let mut up = vec![0.0; n];
        let mut mid = vec![0.0; n];
        let mut lo = vec![0.0; n];

        #[cfg(not(all(target_arch = "wasm32", feature = "wasm")))]
        {
            mab_into(&input, &mut up, &mut mid, &mut lo).expect("mab_into should succeed");
        }

        assert_eq!(base.upperband.len(), n);
        assert_eq!(base.middleband.len(), n);
        assert_eq!(base.lowerband.len(), n);
        assert_eq!(up.len(), n);
        assert_eq!(mid.len(), n);
        assert_eq!(lo.len(), n);

        for i in 0..n {
            assert!(
                eq_or_both_nan(base.upperband[i], up[i]),
                "upper mismatch at {}: base={:?} into={:?}",
                i,
                base.upperband[i],
                up[i]
            );
            assert!(
                eq_or_both_nan(base.middleband[i], mid[i]),
                "middle mismatch at {}: base={:?} into={:?}",
                i,
                base.middleband[i],
                mid[i]
            );
            assert!(
                eq_or_both_nan(base.lowerband[i], lo[i]),
                "lower mismatch at {}: base={:?} into={:?}",
                i,
                base.lowerband[i],
                lo[i]
            );
        }
    }
}

#[inline]
fn mab_compute_into(
    fast_ma: &[f64],
    slow_ma: &[f64],
    fast_period: usize,
    devup: f64,
    devdn: f64,
    first_output: usize,
    kernel: Kernel,
    upper: &mut [f64],
    mid: &mut [f64],
    lower: &mut [f64],
) {
    unsafe {
        match kernel {
            Kernel::Scalar | Kernel::ScalarBatch => mab_scalar(
                fast_ma,
                slow_ma,
                fast_period,
                devup,
                devdn,
                first_output,
                upper,
                mid,
                lower,
            ),
            #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
            Kernel::Avx2 | Kernel::Avx2Batch => mab_avx2(
                fast_ma,
                slow_ma,
                fast_period,
                devup,
                devdn,
                first_output,
                upper,
                mid,
                lower,
            ),
            #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
            Kernel::Avx512 | Kernel::Avx512Batch => mab_avx512(
                fast_ma,
                slow_ma,
                fast_period,
                devup,
                devdn,
                first_output,
                upper,
                mid,
                lower,
            ),
            _ => unreachable!(),
        }
    }
}

#[inline(always)]
pub unsafe fn mab_scalar(
    fast_ma: &[f64],
    slow_ma: &[f64],
    fast_period: usize,
    devup: f64,
    devdn: f64,
    first_output: usize,
    upper: &mut [f64],
    mid: &mut [f64],
    lower: &mut [f64],
) {
    let start_idx = if first_output >= fast_period {
        first_output - fast_period + 1
    } else {
        0
    };

    let mut sum_sq = 0.0;
    for i in start_idx..(start_idx + fast_period).min(fast_ma.len()) {
        let diff = fast_ma[i] - slow_ma[i];
        sum_sq += diff * diff;
    }

    if first_output < fast_ma.len() {
        let dev = (sum_sq / fast_period as f64).sqrt();
        mid[first_output] = fast_ma[first_output];
        upper[first_output] = slow_ma[first_output] + devup * dev;
        lower[first_output] = slow_ma[first_output] - devdn * dev;
    }

    for i in (first_output + 1)..fast_ma.len() {
        let old_idx = i - fast_period;
        let old = fast_ma[old_idx] - slow_ma[old_idx];
        let new = fast_ma[i] - slow_ma[i];
        sum_sq += new * new - old * old;
        let dev = (sum_sq / fast_period as f64).sqrt();

        mid[i] = fast_ma[i];
        upper[i] = slow_ma[i] + devup * dev;
        lower[i] = slow_ma[i] - devdn * dev;
    }
}

#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
#[target_feature(enable = "avx2,fma")]
pub unsafe fn mab_avx2(
    fast_ma: &[f64],
    slow_ma: &[f64],
    fast_period: usize,
    devup: f64,
    devdn: f64,
    first_output: usize,
    upper: &mut [f64],
    mid: &mut [f64],
    lower: &mut [f64],
) {
    use core::arch::x86_64::*;
    let n = fast_ma.len();
    if first_output >= n {
        return;
    }
    debug_assert!(fast_period > 0);
    debug_assert!(first_output + 1 >= fast_period);

    let start = first_output + 1 - fast_period;
    let m = n - start;

    let mut diffsq: AVec<f64> = AVec::with_capacity(CACHELINE_ALIGN, m);
    diffsq.set_len(m);
    let mut prefix: AVec<f64> = AVec::with_capacity(CACHELINE_ALIGN, m + 1);
    prefix.set_len(m + 1);

    let f0 = fast_ma.as_ptr().add(start);
    let s0 = slow_ma.as_ptr().add(start);
    let dptr = diffsq.as_mut_ptr();

    let mut k = 0usize;
    while k + 4 <= m {
        let vf = _mm256_loadu_pd(f0.add(k));
        let vs = _mm256_loadu_pd(s0.add(k));
        let vd = _mm256_sub_pd(vf, vs);
        let vd2 = _mm256_mul_pd(vd, vd);
        _mm256_storeu_pd(dptr.add(k), vd2);
        k += 4;
    }
    while k < m {
        let d = *f0.add(k) - *s0.add(k);
        *dptr.add(k) = d * d;
        k += 1;
    }

    let pptr = prefix.as_mut_ptr();
    *pptr = 0.0;
    let mut acc = 0.0f64;
    k = 0;
    while k < m {
        acc += *dptr.add(k);
        *pptr.add(k + 1) = acc;
        k += 1;
    }

    let invf = 1.0 / (fast_period as f64);
    let vinvf = _mm256_set1_pd(invf);
    let vup = _mm256_set1_pd(devup);
    let vdn = _mm256_set1_pd(devdn);

    let mut i = first_output;

    while i < n && (i & 3) != 0 {
        let base = i - start;
        let sum = *pptr.add(base + 1) - *pptr.add(base + 1 - fast_period);
        let dev = (sum * invf).sqrt();
        let sm = *slow_ma.as_ptr().add(i);
        *mid.as_mut_ptr().add(i) = *fast_ma.as_ptr().add(i);
        *upper.as_mut_ptr().add(i) = sm + devup * dev;
        *lower.as_mut_ptr().add(i) = sm - devdn * dev;
        i += 1;
    }

    while i + 3 < n {
        let base = i - start;
        let pend = _mm256_loadu_pd(pptr.add(base + 1));
        let psta = _mm256_loadu_pd(pptr.add(base + 1 - fast_period));
        let vsum = _mm256_sub_pd(pend, psta);
        let vdev = _mm256_sqrt_pd(_mm256_mul_pd(vsum, vinvf));

        let vfast = _mm256_loadu_pd(fast_ma.as_ptr().add(i));
        let vslow = _mm256_loadu_pd(slow_ma.as_ptr().add(i));

        _mm256_storeu_pd(mid.as_mut_ptr().add(i), vfast);
        let vupper = _mm256_fmadd_pd(vup, vdev, vslow);
        let vlower = _mm256_fnmadd_pd(vdn, vdev, vslow);
        _mm256_storeu_pd(upper.as_mut_ptr().add(i), vupper);
        _mm256_storeu_pd(lower.as_mut_ptr().add(i), vlower);

        i += 4;
    }

    while i < n {
        let base = i - start;
        let sum = *pptr.add(base + 1) - *pptr.add(base + 1 - fast_period);
        let dev = (sum * invf).sqrt();
        let sm = *slow_ma.as_ptr().add(i);
        *mid.as_mut_ptr().add(i) = *fast_ma.as_ptr().add(i);
        *upper.as_mut_ptr().add(i) = sm + devup * dev;
        *lower.as_mut_ptr().add(i) = sm - devdn * dev;
        i += 1;
    }
}

#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
#[target_feature(enable = "avx512f")]
pub unsafe fn mab_avx512(
    fast_ma: &[f64],
    slow_ma: &[f64],
    fast_period: usize,
    devup: f64,
    devdn: f64,
    first_output: usize,
    upper: &mut [f64],
    mid: &mut [f64],
    lower: &mut [f64],
) {
    use core::arch::x86_64::*;
    let n = fast_ma.len();
    if first_output >= n {
        return;
    }
    debug_assert!(fast_period > 0);
    debug_assert!(first_output + 1 >= fast_period);

    let start = first_output + 1 - fast_period;
    let m = n - start;

    let mut diffsq: AVec<f64> = AVec::with_capacity(CACHELINE_ALIGN, m);
    diffsq.set_len(m);
    let mut prefix: AVec<f64> = AVec::with_capacity(CACHELINE_ALIGN, m + 1);
    prefix.set_len(m + 1);

    let f0 = fast_ma.as_ptr().add(start);
    let s0 = slow_ma.as_ptr().add(start);
    let dptr = diffsq.as_mut_ptr();

    let mut k = 0usize;
    while k + 8 <= m {
        let vf = _mm512_loadu_pd(f0.add(k));
        let vs = _mm512_loadu_pd(s0.add(k));
        let vd = _mm512_sub_pd(vf, vs);
        let vd2 = _mm512_mul_pd(vd, vd);
        _mm512_storeu_pd(dptr.add(k), vd2);
        k += 8;
    }
    while k < m {
        let d = *f0.add(k) - *s0.add(k);
        *dptr.add(k) = d * d;
        k += 1;
    }

    let pptr = prefix.as_mut_ptr();
    *pptr = 0.0;
    let mut acc = 0.0f64;
    k = 0;
    while k < m {
        acc += *dptr.add(k);
        *pptr.add(k + 1) = acc;
        k += 1;
    }

    let invf = 1.0 / (fast_period as f64);
    let vinvf = _mm512_set1_pd(invf);
    let vup = _mm512_set1_pd(devup);
    let vdn = _mm512_set1_pd(devdn);

    let mut i = first_output;

    while i < n && (i & 7) != 0 {
        let base = i - start;
        let sum = *pptr.add(base + 1) - *pptr.add(base + 1 - fast_period);
        let dev = (sum * invf).sqrt();
        let sm = *slow_ma.as_ptr().add(i);
        *mid.as_mut_ptr().add(i) = *fast_ma.as_ptr().add(i);
        *upper.as_mut_ptr().add(i) = sm + devup * dev;
        *lower.as_mut_ptr().add(i) = sm - devdn * dev;
        i += 1;
    }

    while i + 7 < n {
        let base = i - start;
        let pend = _mm512_loadu_pd(pptr.add(base + 1));
        let psta = _mm512_loadu_pd(pptr.add(base + 1 - fast_period));
        let vsum = _mm512_sub_pd(pend, psta);
        let vdev = _mm512_sqrt_pd(_mm512_mul_pd(vsum, vinvf));

        let vfast = _mm512_loadu_pd(fast_ma.as_ptr().add(i));
        let vslow = _mm512_loadu_pd(slow_ma.as_ptr().add(i));

        _mm512_storeu_pd(mid.as_mut_ptr().add(i), vfast);
        let vupper = _mm512_fmadd_pd(vup, vdev, vslow);
        let vlower = _mm512_fnmadd_pd(vdn, vdev, vslow);
        _mm512_storeu_pd(upper.as_mut_ptr().add(i), vupper);
        _mm512_storeu_pd(lower.as_mut_ptr().add(i), vlower);

        i += 8;
    }

    while i < n {
        let base = i - start;
        let sum = *pptr.add(base + 1) - *pptr.add(base + 1 - fast_period);
        let dev = (sum * invf).sqrt();
        let sm = *slow_ma.as_ptr().add(i);
        *mid.as_mut_ptr().add(i) = *fast_ma.as_ptr().add(i);
        *upper.as_mut_ptr().add(i) = sm + devup * dev;
        *lower.as_mut_ptr().add(i) = sm - devdn * dev;
        i += 1;
    }
}

pub struct MabStream {
    fast_buffer: Vec<f64>,
    slow_buffer: Vec<f64>,
    diffs_buffer: Vec<f64>,
    fast_index: usize,
    slow_index: usize,
    diff_index: usize,
    count: usize,
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: String,
    slow_ma_type: String,
    fast_sum: f64,
    slow_sum: f64,
    fast_ma: f64,
    slow_ma: f64,
    ema_fast: f64,
    ema_slow: f64,
    kernel: Kernel,

    sumsq_diff: f64,
    diffs_filled: usize,
    inv_fast_len: f64,
    ready_threshold: usize,
    k_fast: f64,
    k_slow: f64,
    max_period: usize,
}

impl MabStream {
    pub fn try_new(params: MabParams) -> Result<Self, String> {
        let fast_period = params.fast_period.unwrap_or(10);
        let slow_period = params.slow_period.unwrap_or(50);
        let devup = params.devup.unwrap_or(1.0);
        let devdn = params.devdn.unwrap_or(1.0);
        let fast_ma_type = params.fast_ma_type.unwrap_or_else(|| "sma".to_string());
        let slow_ma_type = params.slow_ma_type.unwrap_or_else(|| "sma".to_string());

        if fast_period == 0 || slow_period == 0 {
            return Err("Period cannot be zero".to_string());
        }

        let max_period = fast_period.max(slow_period);
        let ready_threshold = max_period + fast_period - 1;

        Ok(Self {
            fast_buffer: vec![0.0; fast_period],
            slow_buffer: vec![0.0; slow_period],
            diffs_buffer: vec![0.0; fast_period],
            fast_index: 0,
            slow_index: 0,
            diff_index: 0,
            count: 0,
            fast_period,
            slow_period,
            devup,
            devdn,
            fast_ma_type,
            slow_ma_type,
            fast_sum: 0.0,
            slow_sum: 0.0,
            fast_ma: 0.0,
            slow_ma: 0.0,
            ema_fast: 0.0,
            ema_slow: 0.0,
            kernel: detect_best_kernel(),

            sumsq_diff: 0.0,
            diffs_filled: 0,
            inv_fast_len: 1.0 / fast_period as f64,
            ready_threshold,
            k_fast: 2.0 / (fast_period as f64 + 1.0),
            k_slow: 2.0 / (slow_period as f64 + 1.0),
            max_period,
        })
    }

    pub fn update(&mut self, value: f64) -> Option<(f64, f64, f64)> {
        if !value.is_finite() {
            return None;
        }

        self.count += 1;

        match self.fast_ma_type.as_str() {
            "ema" => {
                if self.count == 1 {
                    self.ema_fast = value;
                } else {
                    self.ema_fast = (1.0 - self.k_fast).mul_add(self.ema_fast, self.k_fast * value);
                }
                self.fast_ma = self.ema_fast;
            }
            _ => {
                if self.count <= self.fast_period {
                    let idx = self.fast_index;
                    self.fast_sum += value;
                    self.fast_buffer[idx] = value;
                    if self.count == self.fast_period {
                        self.fast_ma = self.fast_sum * self.inv_fast_len;
                    }
                    self.fast_index += 1;
                    if self.fast_index == self.fast_period {
                        self.fast_index = 0;
                    }
                } else {
                    let idx = self.fast_index;
                    let old = self.fast_buffer[idx];
                    self.fast_buffer[idx] = value;
                    self.fast_sum += value - old;
                    self.fast_ma = self.fast_sum * self.inv_fast_len;
                    self.fast_index += 1;
                    if self.fast_index == self.fast_period {
                        self.fast_index = 0;
                    }
                }
            }
        }

        match self.slow_ma_type.as_str() {
            "ema" => {
                if self.count == 1 {
                    self.ema_slow = value;
                } else {
                    self.ema_slow = (1.0 - self.k_slow).mul_add(self.ema_slow, self.k_slow * value);
                }
                self.slow_ma = self.ema_slow;
            }
            _ => {
                if self.count <= self.slow_period {
                    let idx = self.slow_index;
                    self.slow_sum += value;
                    self.slow_buffer[idx] = value;
                    if self.count == self.slow_period {
                        self.slow_ma = self.slow_sum / self.slow_period as f64;
                    }
                    self.slow_index += 1;
                    if self.slow_index == self.slow_period {
                        self.slow_index = 0;
                    }
                } else {
                    let idx = self.slow_index;
                    let old = self.slow_buffer[idx];
                    self.slow_buffer[idx] = value;
                    self.slow_sum += value - old;
                    self.slow_ma = self.slow_sum / self.slow_period as f64;
                    self.slow_index += 1;
                    if self.slow_index == self.slow_period {
                        self.slow_index = 0;
                    }
                }
            }
        }

        if self.count < self.max_period {
            return None;
        }

        let diff = self.fast_ma - self.slow_ma;
        let diff2 = diff * diff;

        if self.diffs_filled < self.fast_period {
            self.sumsq_diff += diff2;
            self.diffs_buffer[self.diff_index] = diff2;
            self.diff_index += 1;
            if self.diff_index == self.fast_period {
                self.diff_index = 0;
            }
            self.diffs_filled += 1;
        } else {
            let old2 = self.diffs_buffer[self.diff_index];
            self.sumsq_diff += diff2 - old2;
            self.diffs_buffer[self.diff_index] = diff2;
            self.diff_index += 1;
            if self.diff_index == self.fast_period {
                self.diff_index = 0;
            }
        }

        if self.count < self.ready_threshold || self.diffs_filled < self.fast_period {
            return None;
        }

        let dev = (self.sumsq_diff * self.inv_fast_len).sqrt();

        let upper = dev.mul_add(self.devup, self.slow_ma);
        let middle = self.fast_ma;
        let lower = (-self.devdn * dev).mul_add(1.0, self.slow_ma);

        Some((upper, middle, lower))
    }
}

#[derive(Clone, Debug)]
pub struct MabBatchRange {
    pub fast_period: (usize, usize, usize),
    pub slow_period: (usize, usize, usize),
    pub devup: (f64, f64, f64),
    pub devdn: (f64, f64, f64),
    pub fast_ma_type: (String, String, String),
    pub slow_ma_type: (String, String, String),
}

impl Default for MabBatchRange {
    fn default() -> Self {
        Self {
            fast_period: (10, 10, 0),
            slow_period: (50, 299, 1),
            devup: (1.0, 1.0, 0.0),
            devdn: (1.0, 1.0, 0.0),
            fast_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
            slow_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
        }
    }
}

#[derive(Clone, Debug)]
pub struct MabBatchOutput {
    pub upperbands: Vec<f64>,
    pub middlebands: Vec<f64>,
    pub lowerbands: Vec<f64>,
    pub combos: Vec<MabParams>,
    pub rows: usize,
    pub cols: usize,
}
impl MabBatchOutput {
    pub fn row_for_params(&self, p: &MabParams) -> Option<usize> {
        self.combos.iter().position(|c| {
            c.fast_period == p.fast_period
                && c.slow_period == p.slow_period
                && c.devup == p.devup
                && c.devdn == p.devdn
                && c.fast_ma_type == p.fast_ma_type
                && c.slow_ma_type == p.slow_ma_type
        })
    }

    pub fn upper_slice(&self, row: usize) -> Option<&[f64]> {
        if row < self.rows {
            let start = row * self.cols;
            let end = start + self.cols;
            Some(&self.upperbands[start..end])
        } else {
            None
        }
    }

    pub fn middle_slice(&self, row: usize) -> Option<&[f64]> {
        if row < self.rows {
            let start = row * self.cols;
            let end = start + self.cols;
            Some(&self.middlebands[start..end])
        } else {
            None
        }
    }

    pub fn lower_slice(&self, row: usize) -> Option<&[f64]> {
        if row < self.rows {
            let start = row * self.cols;
            let end = start + self.cols;
            Some(&self.lowerbands[start..end])
        } else {
            None
        }
    }
}

pub(crate) fn expand_grid(p: &MabBatchRange) -> Result<Vec<MabParams>, MabError> {
    fn axis_usize(axis: (usize, usize, usize)) -> Result<Vec<usize>, MabError> {
        let (start, end, step) = axis;
        if step == 0 || start == end {
            return Ok(vec![start]);
        }
        let (lo, hi) = if start <= end {
            (start, end)
        } else {
            (end, start)
        };
        let v: Vec<usize> = (lo..=hi).step_by(step).collect();
        if v.is_empty() {
            return Err(MabError::InvalidRange { start, end, step });
        }
        Ok(v)
    }

    fn axis_f64(axis: (f64, f64, f64)) -> Result<Vec<f64>, MabError> {
        let (start, end, step) = axis;
        const EPS: f64 = 1e-12;
        if step.abs() < EPS || (start - end).abs() < EPS {
            return Ok(vec![start]);
        }
        let step_eff = if start <= end {
            step.abs()
        } else {
            -step.abs()
        };
        let mut v = Vec::new();
        let mut x = start;
        if step_eff > 0.0 {
            while x <= end + EPS {
                v.push(x);
                x += step_eff;
            }
        } else {
            while x >= end - EPS {
                v.push(x);
                x += step_eff;
            }
        }
        if v.is_empty() {
            return Err(MabError::InvalidRangeF64 { start, end, step });
        }
        Ok(v)
    }

    let fast_periods = axis_usize(p.fast_period)?;
    let slow_periods = axis_usize(p.slow_period)?;
    let devups = axis_f64(p.devup)?;
    let devdns = axis_f64(p.devdn)?;

    let mut combos =
        Vec::with_capacity(fast_periods.len() * slow_periods.len() * devups.len() * devdns.len());

    for &fast in &fast_periods {
        for &slow in &slow_periods {
            for &devup in &devups {
                for &devdn in &devdns {
                    combos.push(MabParams {
                        fast_period: Some(fast),
                        slow_period: Some(slow),
                        devup: Some(devup),
                        devdn: Some(devdn),
                        fast_ma_type: Some(p.fast_ma_type.0.clone()),
                        slow_ma_type: Some(p.slow_ma_type.0.clone()),
                    });
                }
            }
        }
    }

    Ok(combos)
}

pub fn mab_batch(input: &[f64], sweep: &MabBatchRange) -> Result<MabBatchOutput, MabError> {
    mab_batch_inner(input, sweep, Kernel::Auto, false)
}

fn mab_batch_inner(
    input: &[f64],
    sweep: &MabBatchRange,
    kernel: Kernel,
    parallel: bool,
) -> Result<MabBatchOutput, MabError> {
    let kernel = match kernel {
        Kernel::Auto => detect_best_batch_kernel(),
        k => k,
    };
    let simd = match kernel {
        Kernel::Avx512Batch => Kernel::Avx512,
        Kernel::Avx2Batch => Kernel::Avx2,
        Kernel::ScalarBatch => Kernel::Scalar,
        other => {
            return Err(MabError::InvalidKernelForBatch(other));
        }
    };

    let combos = expand_grid(sweep)?;
    let rows = combos.len();
    let cols = input.len();
    if cols == 0 {
        return Err(MabError::EmptyInputData);
    }
    rows.checked_mul(cols).ok_or(MabError::InvalidRange {
        start: sweep.fast_period.0,
        end: sweep.fast_period.1,
        step: sweep.fast_period.2,
    })?;

    let first_valid = input
        .iter()
        .position(|x| !x.is_nan())
        .ok_or(MabError::AllValuesNaN)?;
    let valid = cols - first_valid;
    let warmup_prefixes: Vec<usize> = combos
        .iter()
        .map(|p| {
            let fast = p.fast_period.unwrap();
            let slow = p.slow_period.unwrap();
            if fast == 0 || slow == 0 || fast > cols || slow > cols {
                return Err(MabError::InvalidPeriod {
                    fast,
                    slow,
                    data_len: cols,
                });
            }
            let need_total = fast.max(slow) + fast - 1;
            if valid < need_total {
                return Err(MabError::NotEnoughValidData {
                    needed: need_total,
                    valid,
                });
            }
            Ok(first_valid + need_total)
        })
        .collect::<Result<Vec<_>, MabError>>()?;

    let mut upper_buf = make_uninit_matrix(rows, cols);
    let mut middle_buf = make_uninit_matrix(rows, cols);
    let mut lower_buf = make_uninit_matrix(rows, cols);

    init_matrix_prefixes(&mut upper_buf, cols, &warmup_prefixes);
    init_matrix_prefixes(&mut middle_buf, cols, &warmup_prefixes);
    init_matrix_prefixes(&mut lower_buf, cols, &warmup_prefixes);

    let upper_slice =
        unsafe { std::slice::from_raw_parts_mut(upper_buf.as_mut_ptr() as *mut f64, rows * cols) };
    let middle_slice =
        unsafe { std::slice::from_raw_parts_mut(middle_buf.as_mut_ptr() as *mut f64, rows * cols) };
    let lower_slice =
        unsafe { std::slice::from_raw_parts_mut(lower_buf.as_mut_ptr() as *mut f64, rows * cols) };

    let combos = mab_batch_inner_into(
        input,
        sweep,
        simd,
        parallel,
        upper_slice,
        middle_slice,
        lower_slice,
    )?;

    let mut upper_guard = core::mem::ManuallyDrop::new(upper_buf);
    let mut middle_guard = core::mem::ManuallyDrop::new(middle_buf);
    let mut lower_guard = core::mem::ManuallyDrop::new(lower_buf);

    let upperbands = unsafe {
        Vec::from_raw_parts(
            upper_guard.as_mut_ptr() as *mut f64,
            upper_guard.len(),
            upper_guard.capacity(),
        )
    };
    let middlebands = unsafe {
        Vec::from_raw_parts(
            middle_guard.as_mut_ptr() as *mut f64,
            middle_guard.len(),
            middle_guard.capacity(),
        )
    };
    let lowerbands = unsafe {
        Vec::from_raw_parts(
            lower_guard.as_mut_ptr() as *mut f64,
            lower_guard.len(),
            lower_guard.capacity(),
        )
    };

    Ok(MabBatchOutput {
        upperbands,
        middlebands,
        lowerbands,
        combos,
        rows,
        cols,
    })
}

fn mab_batch_inner_into(
    input: &[f64],
    sweep: &MabBatchRange,
    kernel: Kernel,
    parallel: bool,
    upper_out: &mut [f64],
    middle_out: &mut [f64],
    lower_out: &mut [f64],
) -> Result<Vec<MabParams>, MabError> {
    let combos = expand_grid(sweep)?;
    let rows = combos.len();
    let cols = input.len();
    let expected = rows.checked_mul(cols).ok_or(MabError::InvalidRange {
        start: sweep.fast_period.0,
        end: sweep.fast_period.1,
        step: sweep.fast_period.2,
    })?;

    if upper_out.len() != expected || middle_out.len() != expected || lower_out.len() != expected {
        return Err(MabError::OutputLengthMismatch {
            upper_len: upper_out.len(),
            middle_len: middle_out.len(),
            lower_len: lower_out.len(),
            expected,
        });
    }

    if !combos.is_empty() {
        let p0 = &combos[0];
        let all_same_ma = combos.iter().all(|p| {
            p.fast_period == p0.fast_period
                && p.slow_period == p0.slow_period
                && p.fast_ma_type == p0.fast_ma_type
                && p.slow_ma_type == p0.slow_ma_type
        });

        if all_same_ma {
            use crate::indicators::ema::{ema, EmaInput, EmaParams};
            use crate::indicators::sma::{sma, SmaInput, SmaParams};

            let n = input.len();
            let first = input.iter().position(|x| !x.is_nan()).unwrap_or(0);
            let fast = p0.fast_period.unwrap();
            let slow = p0.slow_period.unwrap();
            let fast_ma_type = p0.fast_ma_type.as_deref().unwrap_or("sma");
            let slow_ma_type = p0.slow_ma_type.as_deref().unwrap_or("sma");

            let fast_ma = match fast_ma_type {
                "ema" => {
                    let params = EmaParams { period: Some(fast) };
                    ema(&EmaInput::from_slice(input, params))
                        .map_err(|_| MabError::NotEnoughValidData {
                            needed: fast,
                            valid: n - first,
                        })?
                        .values
                }
                _ => {
                    let params = SmaParams { period: Some(fast) };
                    sma(&SmaInput::from_slice(input, params))
                        .map_err(|_| MabError::NotEnoughValidData {
                            needed: fast,
                            valid: n - first,
                        })?
                        .values
                }
            };

            let slow_ma = match slow_ma_type {
                "ema" => {
                    let params = EmaParams { period: Some(slow) };
                    ema(&EmaInput::from_slice(input, params))
                        .map_err(|_| MabError::NotEnoughValidData {
                            needed: slow,
                            valid: n - first,
                        })?
                        .values
                }
                _ => {
                    let params = SmaParams { period: Some(slow) };
                    sma(&SmaInput::from_slice(input, params))
                        .map_err(|_| MabError::NotEnoughValidData {
                            needed: slow,
                            valid: n - first,
                        })?
                        .values
                }
            };

            let need_total = fast.max(slow) + fast - 1;
            let warmup = first + need_total - 1;
            let first_output = warmup + 1;

            if first_output < n {
                let mut dev: AVec<f64> = AVec::with_capacity(CACHELINE_ALIGN, n);
                unsafe {
                    dev.set_len(n);
                }

                unsafe {
                    let f_ptr = fast_ma.as_ptr();
                    let s_ptr = slow_ma.as_ptr();
                    let d_ptr = dev.as_mut_ptr();

                    let start = first_output + 1 - fast;
                    let mut sum_sq = 0.0f64;
                    let mut k = 0usize;
                    while k < fast {
                        let idx = start + k;
                        let diff = *f_ptr.add(idx) - *s_ptr.add(idx);
                        sum_sq += diff * diff;
                        k += 1;
                    }

                    *d_ptr.add(first_output) = (sum_sq / fast as f64).sqrt();

                    let mut i = first_output + 1;
                    while i < n {
                        let old_idx = i - fast;
                        let old = *f_ptr.add(old_idx) - *s_ptr.add(old_idx);
                        let new = *f_ptr.add(i) - *s_ptr.add(i);
                        sum_sq += new * new - old * old;
                        *d_ptr.add(i) = (sum_sq / fast as f64).sqrt();
                        i += 1;
                    }
                }

                let fill_row = |row: usize, u: &mut [f64], m: &mut [f64], l: &mut [f64]| {
                    let pr = &combos[row];
                    let devup = pr.devup.unwrap();
                    let devdn = pr.devdn.unwrap();
                    for i in first_output..n {
                        let d = dev[i];
                        m[i] = fast_ma[i];
                        u[i] = slow_ma[i] + devup * d;
                        l[i] = slow_ma[i] - devdn * d;
                    }
                    Ok(())
                };

                if parallel {
                    #[cfg(not(target_arch = "wasm32"))]
                    {
                        use rayon::prelude::*;
                        upper_out
                            .par_chunks_mut(cols)
                            .zip(middle_out.par_chunks_mut(cols))
                            .zip(lower_out.par_chunks_mut(cols))
                            .enumerate()
                            .try_for_each(|(row, ((u, m), l))| fill_row(row, u, m, l))?;
                    }
                    #[cfg(target_arch = "wasm32")]
                    {
                        for row in 0..rows {
                            let s = row * cols;
                            fill_row(
                                row,
                                &mut upper_out[s..s + cols],
                                &mut middle_out[s..s + cols],
                                &mut lower_out[s..s + cols],
                            )?;
                        }
                    }
                } else {
                    for row in 0..rows {
                        let s = row * cols;
                        fill_row(
                            row,
                            &mut upper_out[s..s + cols],
                            &mut middle_out[s..s + cols],
                            &mut lower_out[s..s + cols],
                        )?;
                    }
                }

                return Ok(combos);
            }
        }
    }

    let process_row = |row: usize, u: &mut [f64], m: &mut [f64], l: &mut [f64]| {
        let p = &combos[row];
        let in_row = MabInput::from_slice(
            input,
            MabParams {
                fast_period: p.fast_period,
                slow_period: p.slow_period,
                devup: p.devup,
                devdn: p.devdn,
                fast_ma_type: p.fast_ma_type.clone(),
                slow_ma_type: p.slow_ma_type.clone(),
            },
        );
        mab_into_slice(u, m, l, &in_row, kernel)
    };

    #[cfg(not(target_arch = "wasm32"))]
    {
        if parallel {
            use rayon::prelude::*;
            upper_out
                .par_chunks_mut(cols)
                .zip(middle_out.par_chunks_mut(cols))
                .zip(lower_out.par_chunks_mut(cols))
                .enumerate()
                .try_for_each(|(row, ((u, m), l))| process_row(row, u, m, l))?;
        } else {
            for row in 0..rows {
                let s = row * cols;
                process_row(
                    row,
                    &mut upper_out[s..s + cols],
                    &mut middle_out[s..s + cols],
                    &mut lower_out[s..s + cols],
                )?;
            }
        }
    }

    #[cfg(target_arch = "wasm32")]
    {
        for row in 0..rows {
            let s = row * cols;
            process_row(
                row,
                &mut upper_out[s..s + cols],
                &mut middle_out[s..s + cols],
                &mut lower_out[s..s + cols],
            )?;
        }
    }

    Ok(combos)
}

#[cfg(feature = "python")]
#[pyfunction(name = "mab")]
#[pyo3(signature = (data, fast_period=10, slow_period=50, devup=1.0, devdn=1.0, fast_ma_type="sma", slow_ma_type="sma", kernel=None))]
pub fn mab_py<'py>(
    py: Python<'py>,
    data: numpy::PyReadonlyArray1<'py, f64>,
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
    kernel: Option<&str>,
) -> PyResult<(
    Bound<'py, PyArray1<f64>>,
    Bound<'py, PyArray1<f64>>,
    Bound<'py, PyArray1<f64>>,
)> {
    let slice_in = data.as_slice()?;
    let params = MabParams {
        fast_period: Some(fast_period),
        slow_period: Some(slow_period),
        devup: Some(devup),
        devdn: Some(devdn),
        fast_ma_type: Some(fast_ma_type.to_string()),
        slow_ma_type: Some(slow_ma_type.to_string()),
    };
    let input = MabInput::from_slice(slice_in, params);

    let chosen_kernel = validate_kernel(kernel, false)?;

    let result = py
        .allow_threads(|| match chosen_kernel {
            Kernel::Auto => mab(&input),
            k => mab_with_kernel(&input, k),
        })
        .map_err(|e| PyValueError::new_err(e.to_string()))?;

    Ok((
        result.upperband.into_pyarray(py),
        result.middleband.into_pyarray(py),
        result.lowerband.into_pyarray(py),
    ))
}

#[cfg(feature = "python")]
#[pyclass(name = "MabStream")]
pub struct MabStreamPy {
    stream: MabStream,
}

#[cfg(feature = "python")]
#[pymethods]
impl MabStreamPy {
    #[new]
    fn new(
        fast_period: usize,
        slow_period: usize,
        devup: f64,
        devdn: f64,
        fast_ma_type: &str,
        slow_ma_type: &str,
    ) -> PyResult<Self> {
        let params = MabParams {
            fast_period: Some(fast_period),
            slow_period: Some(slow_period),
            devup: Some(devup),
            devdn: Some(devdn),
            fast_ma_type: Some(fast_ma_type.to_string()),
            slow_ma_type: Some(slow_ma_type.to_string()),
        };
        let stream =
            MabStream::try_new(params).map_err(|e| PyValueError::new_err(e.to_string()))?;
        Ok(MabStreamPy { stream })
    }

    fn update(&mut self, value: f64) -> Option<(f64, f64, f64)> {
        self.stream.update(value)
    }
}

#[cfg(feature = "python")]
#[pyfunction(name = "mab_batch")]
#[pyo3(signature = (data, fast_period_range, slow_period_range, devup_range=(1.0, 1.0, 0.0), devdn_range=(1.0, 1.0, 0.0), fast_ma_type="sma", slow_ma_type="sma", kernel=None))]
pub fn mab_batch_py<'py>(
    py: Python<'py>,
    data: numpy::PyReadonlyArray1<'py, f64>,
    fast_period_range: (usize, usize, usize),
    slow_period_range: (usize, usize, usize),
    devup_range: (f64, f64, f64),
    devdn_range: (f64, f64, f64),
    fast_ma_type: &str,
    slow_ma_type: &str,
    kernel: Option<&str>,
) -> PyResult<Bound<'py, pyo3::types::PyDict>> {
    use numpy::{IntoPyArray, PyArray1, PyArrayMethods};

    let slice_in = data.as_slice()?;

    let sweep = MabBatchRange {
        fast_period: fast_period_range,
        slow_period: slow_period_range,
        devup: devup_range,
        devdn: devdn_range,
        fast_ma_type: (
            fast_ma_type.to_string(),
            fast_ma_type.to_string(),
            "".to_string(),
        ),
        slow_ma_type: (
            slow_ma_type.to_string(),
            slow_ma_type.to_string(),
            "".to_string(),
        ),
    };

    let combos = expand_grid(&sweep).map_err(|e| PyValueError::new_err(e.to_string()))?;
    let rows = combos.len();
    let cols = slice_in.len();
    if cols == 0 {
        return Err(PyValueError::new_err(MabError::EmptyInputData.to_string()));
    }
    let total = rows
        .checked_mul(cols)
        .ok_or_else(|| PyValueError::new_err("mab_batch: rows*cols overflow"))?;

    let upper_arr = unsafe { PyArray1::<f64>::new(py, [total], false) };
    let middle_arr = unsafe { PyArray1::<f64>::new(py, [total], false) };
    let lower_arr = unsafe { PyArray1::<f64>::new(py, [total], false) };

    let slice_upper = unsafe { upper_arr.as_slice_mut()? };
    let slice_middle = unsafe { middle_arr.as_slice_mut()? };
    let slice_lower = unsafe { lower_arr.as_slice_mut()? };

    let kern = validate_kernel(kernel, true)?;

    let first_valid = slice_in
        .iter()
        .position(|x| !x.is_nan())
        .ok_or_else(|| PyValueError::new_err(MabError::AllValuesNaN.to_string()))?;
    let valid = cols - first_valid;
    let warmup_prefixes: Vec<usize> = combos
        .iter()
        .map(|p| {
            let fast = p.fast_period.unwrap();
            let slow = p.slow_period.unwrap();
            if fast == 0 || slow == 0 || fast > cols || slow > cols {
                return Err(MabError::InvalidPeriod {
                    fast,
                    slow,
                    data_len: cols,
                });
            }
            let need_total = fast.max(slow) + fast - 1;
            if valid < need_total {
                return Err(MabError::NotEnoughValidData {
                    needed: need_total,
                    valid,
                });
            }
            Ok(first_valid + need_total)
        })
        .collect::<Result<Vec<_>, MabError>>()
        .map_err(|e| PyValueError::new_err(e.to_string()))?;

    let mu_upper: &mut [MaybeUninit<f64>] = unsafe {
        let ptr = upper_arr.as_array_mut().as_mut_ptr();
        std::slice::from_raw_parts_mut(ptr as *mut MaybeUninit<f64>, total)
    };
    let mu_middle: &mut [MaybeUninit<f64>] = unsafe {
        let ptr = middle_arr.as_array_mut().as_mut_ptr();
        std::slice::from_raw_parts_mut(ptr as *mut MaybeUninit<f64>, total)
    };
    let mu_lower: &mut [MaybeUninit<f64>] = unsafe {
        let ptr = lower_arr.as_array_mut().as_mut_ptr();
        std::slice::from_raw_parts_mut(ptr as *mut MaybeUninit<f64>, total)
    };
    init_matrix_prefixes(mu_upper, cols, &warmup_prefixes);
    init_matrix_prefixes(mu_middle, cols, &warmup_prefixes);
    init_matrix_prefixes(mu_lower, cols, &warmup_prefixes);

    let combos = py
        .allow_threads(|| {
            let kernel = match kern {
                Kernel::Auto => detect_best_batch_kernel(),
                k => k,
            };
            let simd = match kernel {
                Kernel::Avx512Batch => Kernel::Avx512,
                Kernel::Avx2Batch => Kernel::Avx2,
                Kernel::ScalarBatch => Kernel::Scalar,
                _ => unreachable!(),
            };

            mab_batch_inner_into(
                slice_in,
                &sweep,
                simd,
                true,
                slice_upper,
                slice_middle,
                slice_lower,
            )
        })
        .map_err(|e| PyValueError::new_err(e.to_string()))?;

    let dict = PyDict::new(py);
    dict.set_item("upperbands", upper_arr.reshape((rows, cols))?)?;
    dict.set_item("middlebands", middle_arr.reshape((rows, cols))?)?;
    dict.set_item("lowerbands", lower_arr.reshape((rows, cols))?)?;

    dict.set_item(
        "fast_periods",
        combos
            .iter()
            .map(|p| p.fast_period.unwrap() as u64)
            .collect::<Vec<_>>()
            .into_pyarray(py),
    )?;
    dict.set_item(
        "slow_periods",
        combos
            .iter()
            .map(|p| p.slow_period.unwrap() as u64)
            .collect::<Vec<_>>()
            .into_pyarray(py),
    )?;
    dict.set_item(
        "devups",
        combos
            .iter()
            .map(|p| p.devup.unwrap())
            .collect::<Vec<_>>()
            .into_pyarray(py),
    )?;
    dict.set_item(
        "devdns",
        combos
            .iter()
            .map(|p| p.devdn.unwrap())
            .collect::<Vec<_>>()
            .into_pyarray(py),
    )?;

    Ok(dict)
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[derive(Serialize, Deserialize)]
pub struct MabJsSingle {
    pub values: Vec<f64>,
    pub rows: usize,
    pub cols: usize,
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen(js_name = "mab")]
pub fn mab_wasm(
    data: &[f64],
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
) -> Result<JsValue, JsValue> {
    let params = MabParams {
        fast_period: Some(fast_period),
        slow_period: Some(slow_period),
        devup: Some(devup),
        devdn: Some(devdn),
        fast_ma_type: Some(fast_ma_type.to_string()),
        slow_ma_type: Some(slow_ma_type.to_string()),
    };
    let input = MabInput::from_slice(data, params);

    let mut upper = vec![0.0; data.len()];
    let mut middle = vec![0.0; data.len()];
    let mut lower = vec![0.0; data.len()];

    mab_into_slice(
        &mut upper,
        &mut middle,
        &mut lower,
        &input,
        detect_best_kernel(),
    )
    .map_err(|e| JsValue::from_str(&e.to_string()))?;

    let mut values = Vec::with_capacity(3 * data.len());
    values.extend_from_slice(&upper);
    values.extend_from_slice(&middle);
    values.extend_from_slice(&lower);

    serde_wasm_bindgen::to_value(&MabJsSingle {
        values,
        rows: 3,
        cols: data.len(),
    })
    .map_err(|e| JsValue::from_str(&format!("Serialization error: {}", e)))
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen]
pub fn mab_js(
    data: &[f64],
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
) -> Result<Vec<f64>, JsValue> {
    let params = MabParams {
        fast_period: Some(fast_period),
        slow_period: Some(slow_period),
        devup: Some(devup),
        devdn: Some(devdn),
        fast_ma_type: Some(fast_ma_type.to_string()),
        slow_ma_type: Some(slow_ma_type.to_string()),
    };
    let input = MabInput::from_slice(data, params);

    let mut upper = vec![0.0; data.len()];
    let mut middle = vec![0.0; data.len()];
    let mut lower = vec![0.0; data.len()];

    mab_into_slice(
        &mut upper,
        &mut middle,
        &mut lower,
        &input,
        detect_best_kernel(),
    )
    .map_err(|e| JsValue::from_str(&e.to_string()))?;

    let mut result = Vec::with_capacity(3 * data.len());
    result.extend_from_slice(&upper);
    result.extend_from_slice(&middle);
    result.extend_from_slice(&lower);

    Ok(result)
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[derive(Serialize, Deserialize)]
pub struct MabBatchConfig {
    pub fast_period_range: (usize, usize, usize),
    pub slow_period_range: (usize, usize, usize),
    pub devup_range: (f64, f64, f64),
    pub devdn_range: (f64, f64, f64),
    pub fast_ma_type: String,
    pub slow_ma_type: String,
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[derive(Serialize, Deserialize)]
pub struct MabBatchJsOutput {
    pub upperbands: Vec<f64>,
    pub middlebands: Vec<f64>,
    pub lowerbands: Vec<f64>,
    pub combos: Vec<MabParams>,
    pub rows: usize,
    pub cols: usize,
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen(js_name = mab_batch)]
pub fn mab_batch_unified_js(data: &[f64], config: JsValue) -> Result<JsValue, JsValue> {
    let config: MabBatchConfig = serde_wasm_bindgen::from_value(config)
        .map_err(|e| JsValue::from_str(&format!("Invalid config: {}", e)))?;

    let sweep = MabBatchRange {
        fast_period: config.fast_period_range,
        slow_period: config.slow_period_range,
        devup: config.devup_range,
        devdn: config.devdn_range,
        fast_ma_type: (
            config.fast_ma_type.clone(),
            config.fast_ma_type.clone(),
            "".to_string(),
        ),
        slow_ma_type: (
            config.slow_ma_type.clone(),
            config.slow_ma_type.clone(),
            "".to_string(),
        ),
    };

    let output = mab_batch_inner(data, &sweep, Kernel::Auto, false)
        .map_err(|e| JsValue::from_str(&e.to_string()))?;

    let js_output = MabBatchJsOutput {
        upperbands: output.upperbands,
        middlebands: output.middlebands,
        lowerbands: output.lowerbands,
        combos: output.combos,
        rows: output.rows,
        cols: output.cols,
    };

    serde_wasm_bindgen::to_value(&js_output)
        .map_err(|e| JsValue::from_str(&format!("Serialization error: {}", e)))
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen]
pub fn mab_alloc(len: usize) -> *mut f64 {
    let mut vec = Vec::<f64>::with_capacity(len);
    let ptr = vec.as_mut_ptr();
    std::mem::forget(vec);
    ptr
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen]
pub fn mab_free(ptr: *mut f64, len: usize) {
    if !ptr.is_null() {
        unsafe {
            let _ = Vec::from_raw_parts(ptr, len, len);
        }
    }
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen]
pub fn mab_into(
    in_ptr: *const f64,
    upper_ptr: *mut f64,
    middle_ptr: *mut f64,
    lower_ptr: *mut f64,
    len: usize,
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
) -> Result<(), JsValue> {
    if in_ptr.is_null() || upper_ptr.is_null() || middle_ptr.is_null() || lower_ptr.is_null() {
        return Err(JsValue::from_str("Null pointer provided"));
    }

    unsafe {
        let data = std::slice::from_raw_parts(in_ptr, len);
        let params = MabParams {
            fast_period: Some(fast_period),
            slow_period: Some(slow_period),
            devup: Some(devup),
            devdn: Some(devdn),
            fast_ma_type: Some(fast_ma_type.to_string()),
            slow_ma_type: Some(slow_ma_type.to_string()),
        };
        let input = MabInput::from_slice(data, params);

        let need_temp = in_ptr == upper_ptr || in_ptr == middle_ptr || in_ptr == lower_ptr;

        if need_temp {
            let mut temp_upper = vec![0.0; len];
            let mut temp_middle = vec![0.0; len];
            let mut temp_lower = vec![0.0; len];

            mab_into_slice(
                &mut temp_upper,
                &mut temp_middle,
                &mut temp_lower,
                &input,
                Kernel::Auto,
            )
            .map_err(|e| JsValue::from_str(&e.to_string()))?;

            let upper_out = std::slice::from_raw_parts_mut(upper_ptr, len);
            let middle_out = std::slice::from_raw_parts_mut(middle_ptr, len);
            let lower_out = std::slice::from_raw_parts_mut(lower_ptr, len);

            upper_out.copy_from_slice(&temp_upper);
            middle_out.copy_from_slice(&temp_middle);
            lower_out.copy_from_slice(&temp_lower);
        } else {
            let upper_out = std::slice::from_raw_parts_mut(upper_ptr, len);
            let middle_out = std::slice::from_raw_parts_mut(middle_ptr, len);
            let lower_out = std::slice::from_raw_parts_mut(lower_ptr, len);

            mab_into_slice(upper_out, middle_out, lower_out, &input, Kernel::Auto)
                .map_err(|e| JsValue::from_str(&e.to_string()))?;
        }
        Ok(())
    }
}

#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
#[wasm_bindgen]
pub fn mab_batch_into(
    in_ptr: *const f64,
    upper_ptr: *mut f64,
    middle_ptr: *mut f64,
    lower_ptr: *mut f64,
    len: usize,
    fast_period_start: usize,
    fast_period_end: usize,
    fast_period_step: usize,
    slow_period_start: usize,
    slow_period_end: usize,
    slow_period_step: usize,
    devup_start: f64,
    devup_end: f64,
    devup_step: f64,
    devdn_start: f64,
    devdn_end: f64,
    devdn_step: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
) -> Result<usize, JsValue> {
    if in_ptr.is_null() || upper_ptr.is_null() || middle_ptr.is_null() || lower_ptr.is_null() {
        return Err(JsValue::from_str("Null pointer passed to mab_batch_into"));
    }

    unsafe {
        let data = std::slice::from_raw_parts(in_ptr, len);

        let sweep = MabBatchRange {
            fast_period: (fast_period_start, fast_period_end, fast_period_step),
            slow_period: (slow_period_start, slow_period_end, slow_period_step),
            devup: (devup_start, devup_end, devup_step),
            devdn: (devdn_start, devdn_end, devdn_step),
            fast_ma_type: (
                fast_ma_type.to_string(),
                fast_ma_type.to_string(),
                "".to_string(),
            ),
            slow_ma_type: (
                slow_ma_type.to_string(),
                slow_ma_type.to_string(),
                "".to_string(),
            ),
        };

        let combos = expand_grid(&sweep).map_err(|e| JsValue::from_str(&e.to_string()))?;
        let rows = combos.len();
        let cols = len;
        if cols == 0 {
            return Err(JsValue::from_str(&MabError::EmptyInputData.to_string()));
        }
        let total = rows
            .checked_mul(cols)
            .ok_or_else(|| JsValue::from_str("mab_batch_into: rows*cols overflow"))?;

        let first_valid = data
            .iter()
            .position(|x| !x.is_nan())
            .ok_or(MabError::AllValuesNaN)
            .map_err(|e| JsValue::from_str(&e.to_string()))?;
        let valid = cols - first_valid;
        let warmup_prefixes: Vec<usize> = combos
            .iter()
            .map(|p| {
                let fast = p.fast_period.unwrap();
                let slow = p.slow_period.unwrap();
                if fast == 0 || slow == 0 || fast > cols || slow > cols {
                    return Err(MabError::InvalidPeriod {
                        fast,
                        slow,
                        data_len: cols,
                    });
                }
                let need_total = fast.max(slow) + fast - 1;
                if valid < need_total {
                    return Err(MabError::NotEnoughValidData {
                        needed: need_total,
                        valid,
                    });
                }
                Ok(first_valid + need_total)
            })
            .collect::<Result<Vec<_>, MabError>>()
            .map_err(|e| JsValue::from_str(&e.to_string()))?;

        let mu_upper: &mut [MaybeUninit<f64>] =
            std::slice::from_raw_parts_mut(upper_ptr as *mut MaybeUninit<f64>, total);
        let mu_middle: &mut [MaybeUninit<f64>] =
            std::slice::from_raw_parts_mut(middle_ptr as *mut MaybeUninit<f64>, total);
        let mu_lower: &mut [MaybeUninit<f64>] =
            std::slice::from_raw_parts_mut(lower_ptr as *mut MaybeUninit<f64>, total);
        init_matrix_prefixes(mu_upper, cols, &warmup_prefixes);
        init_matrix_prefixes(mu_middle, cols, &warmup_prefixes);
        init_matrix_prefixes(mu_lower, cols, &warmup_prefixes);

        let upper_out = std::slice::from_raw_parts_mut(upper_ptr, total);
        let middle_out = std::slice::from_raw_parts_mut(middle_ptr, total);
        let lower_out = std::slice::from_raw_parts_mut(lower_ptr, total);

        mab_batch_inner_into(
            data,
            &sweep,
            Kernel::Auto,
            false,
            upper_out,
            middle_out,
            lower_out,
        )
        .map_err(|e| JsValue::from_str(&e.to_string()))?;

        Ok(rows)
    }
}

#[inline]
pub unsafe fn mab_scalar_classic_sma(
    data: &[f64],
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    first_valid_idx: usize,
    upper: &mut [f64],
    middle: &mut [f64],
    lower: &mut [f64],
) -> Result<(), MabError> {
    let n = data.len();

    let mut fast_ma = vec![f64::NAN; n];
    if fast_period > 0 && first_valid_idx + fast_period <= n {
        let mut sum = 0.0;
        for i in 0..fast_period {
            sum += data[first_valid_idx + i];
        }
        fast_ma[first_valid_idx + fast_period - 1] = sum / fast_period as f64;

        for i in (first_valid_idx + fast_period)..n {
            sum = sum - data[i - fast_period] + data[i];
            fast_ma[i] = sum / fast_period as f64;
        }
    }

    let mut slow_ma = vec![f64::NAN; n];
    if slow_period > 0 && first_valid_idx + slow_period <= n {
        let mut sum = 0.0;
        for i in 0..slow_period {
            sum += data[first_valid_idx + i];
        }
        slow_ma[first_valid_idx + slow_period - 1] = sum / slow_period as f64;

        for i in (first_valid_idx + slow_period)..n {
            sum = sum - data[i - slow_period] + data[i];
            slow_ma[i] = sum / slow_period as f64;
        }
    }

    let need_total = slow_period.max(fast_period) + fast_period - 1;
    let warmup = first_valid_idx + need_total - 1;
    let first_output = warmup + 1;

    for i in 0..first_output.min(n) {
        upper[i] = f64::NAN;
        middle[i] = f64::NAN;
        lower[i] = f64::NAN;
    }

    if first_output >= n {
        return Ok(());
    }

    let start_idx = if first_output >= fast_period {
        first_output - fast_period + 1
    } else {
        0
    };

    let mut sum_sq = 0.0;
    for i in start_idx..(start_idx + fast_period).min(fast_ma.len()) {
        let diff = fast_ma[i] - slow_ma[i];
        if !diff.is_nan() {
            sum_sq += diff * diff;
        }
    }

    if first_output < fast_ma.len() {
        let dev = (sum_sq / fast_period as f64).sqrt();
        middle[first_output] = fast_ma[first_output];
        upper[first_output] = slow_ma[first_output] + devup * dev;
        lower[first_output] = slow_ma[first_output] - devdn * dev;
    }

    for i in (first_output + 1)..fast_ma.len() {
        let old_idx = i - fast_period;
        let old = fast_ma[old_idx] - slow_ma[old_idx];
        let new = fast_ma[i] - slow_ma[i];
        if !old.is_nan() && !new.is_nan() {
            sum_sq += new * new - old * old;
        }
        let dev = (sum_sq / fast_period as f64).sqrt();

        middle[i] = fast_ma[i];
        upper[i] = slow_ma[i] + devup * dev;
        lower[i] = slow_ma[i] - devdn * dev;
    }

    Ok(())
}

#[cfg(all(feature = "python", feature = "cuda"))]
use crate::cuda::{cuda_available, moving_averages::CudaMab};
#[cfg(all(feature = "python", feature = "cuda"))]
use crate::indicators::moving_averages::alma::{make_device_array_py, DeviceArrayF32Py};
#[cfg(all(feature = "python", feature = "cuda"))]
#[cfg(all(feature = "python", feature = "cuda"))]
use numpy::{PyReadonlyArray1, PyReadonlyArray2, PyUntypedArrayMethods};
#[cfg(all(feature = "python", feature = "cuda"))]
use pyo3::{pyfunction, PyResult, Python};

#[cfg(all(feature = "python", feature = "cuda"))]
#[pyfunction(name = "mab_cuda_batch_dev")]
#[pyo3(signature = (data_f32, fast_period_range, slow_period_range, devup_range=(1.0,1.0,0.0), devdn_range=(1.0,1.0,0.0), fast_ma_type="sma", slow_ma_type="sma", device_id=0))]
pub fn mab_cuda_batch_dev_py(
    py: Python<'_>,
    data_f32: PyReadonlyArray1<'_, f32>,
    fast_period_range: (usize, usize, usize),
    slow_period_range: (usize, usize, usize),
    devup_range: (f64, f64, f64),
    devdn_range: (f64, f64, f64),
    fast_ma_type: &str,
    slow_ma_type: &str,
    device_id: usize,
) -> PyResult<(DeviceArrayF32Py, DeviceArrayF32Py, DeviceArrayF32Py)> {
    if !cuda_available() {
        return Err(PyValueError::new_err("CUDA not available"));
    }
    let slice = data_f32.as_slice()?;
    let sweep = MabBatchRange {
        fast_period: fast_period_range,
        slow_period: slow_period_range,
        devup: devup_range,
        devdn: devdn_range,
        fast_ma_type: (
            fast_ma_type.to_string(),
            fast_ma_type.to_string(),
            String::new(),
        ),
        slow_ma_type: (
            slow_ma_type.to_string(),
            slow_ma_type.to_string(),
            String::new(),
        ),
    };
    let (up, mid, lo) = py.allow_threads(|| {
        let cuda = CudaMab::new(device_id).map_err(|e| PyValueError::new_err(e.to_string()))?;
        let (trip, _combos) = cuda
            .mab_batch_dev(slice, &sweep)
            .map_err(|e| PyValueError::new_err(e.to_string()))?;
        Ok::<_, pyo3::PyErr>((trip.upper, trip.middle, trip.lower))
    })?;

    let up_py = make_device_array_py(device_id, up)?;
    let mid_py = make_device_array_py(device_id, mid)?;
    let lo_py = make_device_array_py(device_id, lo)?;

    Ok((up_py, mid_py, lo_py))
}

#[cfg(all(feature = "python", feature = "cuda"))]
#[pyfunction(name = "mab_cuda_many_series_one_param_dev")]
#[pyo3(signature = (data_tm_f32, fast_period, slow_period, devup=1.0, devdn=1.0, fast_ma_type="sma", slow_ma_type="sma", device_id=0))]
pub fn mab_cuda_many_series_one_param_dev_py(
    py: Python<'_>,
    data_tm_f32: PyReadonlyArray2<'_, f32>,
    fast_period: usize,
    slow_period: usize,
    devup: f64,
    devdn: f64,
    fast_ma_type: &str,
    slow_ma_type: &str,
    device_id: usize,
) -> PyResult<(DeviceArrayF32Py, DeviceArrayF32Py, DeviceArrayF32Py)> {
    if !cuda_available() {
        return Err(PyValueError::new_err("CUDA not available"));
    }
    let flat: &[f32] = data_tm_f32.as_slice()?;
    let rows = data_tm_f32.shape()[0];
    let cols = data_tm_f32.shape()[1];
    let params = MabParams {
        fast_period: Some(fast_period),
        slow_period: Some(slow_period),
        devup: Some(devup),
        devdn: Some(devdn),
        fast_ma_type: Some(fast_ma_type.to_string()),
        slow_ma_type: Some(slow_ma_type.to_string()),
    };
    let (up, mid, lo) = py.allow_threads(|| {
        let cuda = CudaMab::new(device_id).map_err(|e| PyValueError::new_err(e.to_string()))?;
        let trip = cuda
            .mab_many_series_one_param_time_major_dev(flat, cols, rows, &params)
            .map_err(|e| PyValueError::new_err(e.to_string()))?;
        Ok::<_, pyo3::PyErr>((trip.upper, trip.middle, trip.lower))
    })?;

    let up_py = make_device_array_py(device_id, up)?;
    let mid_py = make_device_array_py(device_id, mid)?;
    let lo_py = make_device_array_py(device_id, lo)?;

    Ok((up_py, mid_py, lo_py))
}

#[cfg(test)]
mod tests {
    use super::*;
    use crate::utilities::data_loader::read_candles_from_csv;
    #[cfg(feature = "proptest")]
    use proptest::prelude::*;
    use std::error::Error;

    macro_rules! skip_if_unsupported {
        ($kernel:expr, $test_name:expr) => {
            #[cfg(all(target_arch = "wasm32", target_feature = "simd128"))]
            if matches!(
                $kernel,
                Kernel::Avx2 | Kernel::Avx512 | Kernel::Avx2Batch | Kernel::Avx512Batch
            ) {
                eprintln!(
                    "[{}] Skipping - {:?} not supported on WASM",
                    $test_name, $kernel
                );
                return Ok(());
            }
            #[cfg(not(all(feature = "nightly-avx", target_arch = "x86_64")))]
            if matches!(
                $kernel,
                Kernel::Avx2 | Kernel::Avx512 | Kernel::Avx2Batch | Kernel::Avx512Batch
            ) {
                eprintln!(
                    "[{}] Skipping - {:?} requires 'nightly-avx' feature",
                    $test_name, $kernel
                );
                return Ok(());
            }
        };
    }

    #[cfg(debug_assertions)]
    fn check_mab_no_poison(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);

        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;

        let test_params = vec![
            MabParams::default(),
            MabParams {
                fast_period: Some(2),
                slow_period: Some(3),
                devup: Some(1.0),
                devdn: Some(1.0),
                fast_ma_type: Some("sma".to_string()),
                slow_ma_type: Some("sma".to_string()),
            },
            MabParams {
                fast_period: Some(5),
                slow_period: Some(10),
                devup: Some(0.5),
                devdn: Some(0.5),
                fast_ma_type: Some("sma".to_string()),
                slow_ma_type: Some("sma".to_string()),
            },
            MabParams {
                fast_period: Some(15),
                slow_period: Some(30),
                devup: Some(2.0),
                devdn: Some(2.0),
                fast_ma_type: Some("ema".to_string()),
                slow_ma_type: Some("ema".to_string()),
            },
            MabParams {
                fast_period: Some(50),
                slow_period: Some(100),
                devup: Some(3.0),
                devdn: Some(3.0),
                fast_ma_type: Some("sma".to_string()),
                slow_ma_type: Some("sma".to_string()),
            },
            MabParams {
                fast_period: Some(10),
                slow_period: Some(20),
                devup: Some(1.5),
                devdn: Some(1.5),
                fast_ma_type: Some("sma".to_string()),
                slow_ma_type: Some("ema".to_string()),
            },
            MabParams {
                fast_period: Some(8),
                slow_period: Some(21),
                devup: Some(2.5),
                devdn: Some(2.5),
                fast_ma_type: Some("ema".to_string()),
                slow_ma_type: Some("sma".to_string()),
            },
            MabParams {
                fast_period: Some(12),
                slow_period: Some(26),
                devup: Some(2.0),
                devdn: Some(1.0),
                fast_ma_type: Some("ema".to_string()),
                slow_ma_type: Some("ema".to_string()),
            },
            MabParams {
                fast_period: Some(9),
                slow_period: Some(10),
                devup: Some(1.0),
                devdn: Some(2.0),
                fast_ma_type: Some("sma".to_string()),
                slow_ma_type: Some("sma".to_string()),
            },
            MabParams {
                fast_period: Some(30),
                slow_period: Some(200),
                devup: Some(1.0),
                devdn: Some(1.0),
                fast_ma_type: Some("ema".to_string()),
                slow_ma_type: Some("ema".to_string()),
            },
        ];

        for (param_idx, params) in test_params.iter().enumerate() {
            let input = MabInput::from_candles(&candles, "close", params.clone());
            let output = mab_with_kernel(&input, kernel)?;

            for (i, &val) in output.upperband.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Found alloc_with_nan_prefix poison value {} (0x{:016X}) at index {} in upperband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Found init_matrix_prefixes poison value {} (0x{:016X}) at index {} in upperband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Found make_uninit_matrix poison value {} (0x{:016X}) at index {} in upperband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }
            }

            for (i, &val) in output.middleband.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Found alloc_with_nan_prefix poison value {} (0x{:016X}) at index {} in middleband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Found init_matrix_prefixes poison value {} (0x{:016X}) at index {} in middleband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Found make_uninit_matrix poison value {} (0x{:016X}) at index {} in middleband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }
            }

            for (i, &val) in output.lowerband.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Found alloc_with_nan_prefix poison value {} (0x{:016X}) at index {} in lowerband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Found init_matrix_prefixes poison value {} (0x{:016X}) at index {} in lowerband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Found make_uninit_matrix poison value {} (0x{:016X}) at index {} in lowerband \
						 with params: fast_period={}, slow_period={}, devup={}, devdn={}, fast_ma_type={}, slow_ma_type={} (param set {})",
						test_name, val, bits, i,
						params.fast_period.unwrap_or(10),
						params.slow_period.unwrap_or(50),
						params.devup.unwrap_or(1.0),
						params.devdn.unwrap_or(1.0),
						params.fast_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						params.slow_ma_type.as_ref().unwrap_or(&"sma".to_string()),
						param_idx
					);
                }
            }
        }

        Ok(())
    }

    #[cfg(not(debug_assertions))]
    fn check_mab_no_poison(_test_name: &str, _kernel: Kernel) -> Result<(), Box<dyn Error>> {
        Ok(())
    }

    #[cfg(debug_assertions)]
    fn check_batch_no_poison(test: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test);

        let file = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let c = read_candles_from_csv(file)?;

        let test_configs = vec![
            ((2, 10, 2), (10, 20, 5), (1.0, 1.0, 0.0), (1.0, 1.0, 0.0)),
            ((5, 15, 5), (20, 40, 10), (0.5, 2.0, 0.5), (0.5, 2.0, 0.5)),
            (
                (20, 40, 10),
                (50, 100, 25),
                (1.0, 3.0, 1.0),
                (1.0, 3.0, 1.0),
            ),
            ((2, 5, 1), (6, 10, 1), (1.0, 2.0, 0.5), (1.0, 2.0, 0.5)),
            ((10, 10, 0), (20, 50, 10), (1.0, 1.0, 0.0), (1.0, 1.0, 0.0)),
            ((5, 20, 5), (50, 50, 0), (2.0, 2.0, 0.0), (2.0, 2.0, 0.0)),
            ((8, 12, 2), (26, 26, 0), (1.0, 3.0, 0.5), (0.5, 2.0, 0.5)),
        ];

        for (cfg_idx, &(fast_range, slow_range, devup_range, devdn_range)) in
            test_configs.iter().enumerate()
        {
            let sweep = MabBatchRange {
                fast_period: fast_range,
                slow_period: slow_range,
                devup: devup_range,
                devdn: devdn_range,
                fast_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
                slow_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
            };

            let output = mab_batch_inner(c.close.as_slice(), &sweep, kernel, false)?;

            for (idx, &val) in output.upperbands.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();
                let row = idx / output.cols;
                let col = idx % output.cols;
                let combo = &output.combos[row];

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Config {}: Found alloc_with_nan_prefix poison value {} (0x{:016X}) in upperbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Config {}: Found init_matrix_prefixes poison value {} (0x{:016X}) in upperbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Config {}: Found make_uninit_matrix poison value {} (0x{:016X}) in upperbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }
            }

            for (idx, &val) in output.middlebands.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();
                let row = idx / output.cols;
                let col = idx % output.cols;
                let combo = &output.combos[row];

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Config {}: Found alloc_with_nan_prefix poison value {} (0x{:016X}) in middlebands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Config {}: Found init_matrix_prefixes poison value {} (0x{:016X}) in middlebands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Config {}: Found make_uninit_matrix poison value {} (0x{:016X}) in middlebands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }
            }

            for (idx, &val) in output.lowerbands.iter().enumerate() {
                if val.is_nan() {
                    continue;
                }

                let bits = val.to_bits();
                let row = idx / output.cols;
                let col = idx % output.cols;
                let combo = &output.combos[row];

                if bits == 0x11111111_11111111 {
                    panic!(
						"[{}] Config {}: Found alloc_with_nan_prefix poison value {} (0x{:016X}) in lowerbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x22222222_22222222 {
                    panic!(
						"[{}] Config {}: Found init_matrix_prefixes poison value {} (0x{:016X}) in lowerbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }

                if bits == 0x33333333_33333333 {
                    panic!(
						"[{}] Config {}: Found make_uninit_matrix poison value {} (0x{:016X}) in lowerbands \
						 at row {} col {} (flat index {}) with params: fast_period={}, slow_period={}, devup={}, devdn={}",
						test, cfg_idx, val, bits, row, col, idx,
						combo.fast_period.unwrap_or(10),
						combo.slow_period.unwrap_or(50),
						combo.devup.unwrap_or(1.0),
						combo.devdn.unwrap_or(1.0)
					);
                }
            }
        }

        Ok(())
    }

    #[cfg(not(debug_assertions))]
    fn check_batch_no_poison(_test: &str, _kernel: Kernel) -> Result<(), Box<dyn Error>> {
        Ok(())
    }

    #[cfg(feature = "proptest")]
    #[allow(clippy::float_cmp)]
    fn check_mab_property(
        test_name: &str,
        kernel: Kernel,
    ) -> Result<(), Box<dyn std::error::Error>> {
        use proptest::prelude::*;
        skip_if_unsupported!(kernel, test_name);

        let strat = (2usize..=50).prop_flat_map(|slow_period| {
            (2usize..=slow_period).prop_flat_map(move |fast_period| {
                (
                    prop::collection::vec(
                        (1f64..1000f64).prop_filter("finite", |x| x.is_finite()),
                        slow_period..400,
                    ),
                    Just(fast_period),
                    Just(slow_period),
                    0.5f64..3.0f64,
                    0.5f64..3.0f64,
                    prop::bool::ANY,
                    prop::bool::ANY,
                )
            })
        });

        proptest::test_runner::TestRunner::default()
			.run(&strat, |(data, fast_period, slow_period, devup, devdn, fast_is_ema, slow_is_ema)| {
				let params = MabParams {
					fast_period: Some(fast_period),
					slow_period: Some(slow_period),
					devup: Some(devup),
					devdn: Some(devdn),
					fast_ma_type: Some(if fast_is_ema { "ema" } else { "sma" }.to_string()),
					slow_ma_type: Some(if slow_is_ema { "ema" } else { "sma" }.to_string()),
				};
				let input = MabInput::from_slice(&data, params.clone());


				let result = mab_with_kernel(&input, kernel).unwrap();


				let ref_params = params.clone();
				let ref_input = MabInput::from_slice(&data, ref_params);
				let ref_result = mab_with_kernel(&ref_input, Kernel::Scalar).unwrap();


				let first_valid_idx = data.iter().position(|x| !x.is_nan()).unwrap_or(0);
				let warmup_period = first_valid_idx + fast_period.max(slow_period) - 1;

				for i in 0..data.len() {
					let upper = result.upperband[i];
					let middle = result.middleband[i];
					let lower = result.lowerband[i];
					let ref_upper = ref_result.upperband[i];
					let ref_middle = ref_result.middleband[i];
					let ref_lower = ref_result.lowerband[i];


					if upper.is_nan() {
						prop_assert!(ref_upper.is_nan(),
							"[{}] NaN mismatch in upperband at idx {}: kernel={:?} has NaN but scalar doesn't",
							test_name, i, kernel);
					}
					if middle.is_nan() {
						prop_assert!(ref_middle.is_nan(),
							"[{}] NaN mismatch in middleband at idx {}: kernel={:?} has NaN but scalar doesn't",
							test_name, i, kernel);
					}
					if lower.is_nan() {
						prop_assert!(ref_lower.is_nan(),
							"[{}] NaN mismatch in lowerband at idx {}: kernel={:?} has NaN but scalar doesn't",
							test_name, i, kernel);
					}


					if upper.is_finite() && ref_upper.is_finite() {
						let ulp_diff = upper.to_bits().abs_diff(ref_upper.to_bits());
						prop_assert!(
							(upper - ref_upper).abs() <= 1e-9 || ulp_diff <= 8,
							"[{}] Upperband mismatch at idx {}: {} vs {} (ULP={})",
							test_name, i, upper, ref_upper, ulp_diff
						);
					}
					if middle.is_finite() && ref_middle.is_finite() {
						let ulp_diff = middle.to_bits().abs_diff(ref_middle.to_bits());
						prop_assert!(
							(middle - ref_middle).abs() <= 1e-9 || ulp_diff <= 8,
							"[{}] Middleband mismatch at idx {}: {} vs {} (ULP={})",
							test_name, i, middle, ref_middle, ulp_diff
						);
					}
					if lower.is_finite() && ref_lower.is_finite() {
						let ulp_diff = lower.to_bits().abs_diff(ref_lower.to_bits());
						prop_assert!(
							(lower - ref_lower).abs() <= 1e-9 || ulp_diff <= 8,
							"[{}] Lowerband mismatch at idx {}: {} vs {} (ULP={})",
							test_name, i, lower, ref_lower, ulp_diff
						);
					}
				}


				for i in 0..warmup_period.min(data.len()) {
					prop_assert!(
						result.upperband[i].is_nan(),
						"[{}] Expected NaN in upperband during warmup at idx {} (warmup={})",
						test_name, i, warmup_period
					);
					prop_assert!(
						result.middleband[i].is_nan(),
						"[{}] Expected NaN in middleband during warmup at idx {} (warmup={})",
						test_name, i, warmup_period
					);
					prop_assert!(
						result.lowerband[i].is_nan(),
						"[{}] Expected NaN in lowerband during warmup at idx {} (warmup={})",
						test_name, i, warmup_period
					);
				}



				let first_valid_output = warmup_period + fast_period - 1;
				if first_valid_output < data.len() {
					for i in first_valid_output..data.len() {
						prop_assert!(
							result.upperband[i].is_finite(),
							"[{}] Non-finite value in upperband at idx {} after warmup",
							test_name, i
						);
						prop_assert!(
							result.middleband[i].is_finite(),
							"[{}] Non-finite value in middleband at idx {} after warmup",
							test_name, i
						);
						prop_assert!(
							result.lowerband[i].is_finite(),
							"[{}] Non-finite value in lowerband at idx {} after warmup",
							test_name, i
						);
					}
				}


				for i in first_valid_output..data.len() {
					let upper = result.upperband[i];
					let middle = result.middleband[i];
					let lower = result.lowerband[i];

					if upper.is_finite() && middle.is_finite() && lower.is_finite() {
						prop_assert!(
							upper >= middle - 1e-10,
							"[{}] Band ordering violated: upper {} < middle {} at idx {}",
							test_name, upper, middle, i
						);
						prop_assert!(
							middle >= lower - 1e-10,
							"[{}] Band ordering violated: middle {} < lower {} at idx {}",
							test_name, middle, lower, i
						);
					}
				}


				if data.windows(2).all(|w| (w[0] - w[1]).abs() < 1e-10) && data.len() > first_valid_output {

					for i in first_valid_output..data.len() {
						let upper = result.upperband[i];
						let middle = result.middleband[i];
						let lower = result.lowerband[i];

						if upper.is_finite() && middle.is_finite() && lower.is_finite() {
							prop_assert!(
								(upper - middle).abs() <= 1e-9,
								"[{}] Constant data: upper {} != middle {} at idx {}",
								test_name, upper, middle, i
							);
							prop_assert!(
								(middle - lower).abs() <= 1e-9,
								"[{}] Constant data: middle {} != lower {} at idx {}",
								test_name, middle, lower, i
							);
						}
					}
				}



				for i in first_valid_output..data.len() {
					let upper = result.upperband[i];
					let middle = result.middleband[i];
					let lower = result.lowerband[i];

					if upper.is_finite() && middle.is_finite() && lower.is_finite() {
						let upper_spread = upper - middle;
						let lower_spread = middle - lower;


						if upper_spread > 1e-10 && lower_spread > 1e-10 {
							let spread_ratio = upper_spread / lower_spread;
							let multiplier_ratio = devup / devdn;

							prop_assert!(
								(spread_ratio - multiplier_ratio).abs() <= multiplier_ratio * 0.05,
								"[{}] Deviation multiplier ratio mismatch at idx {}: spread_ratio={} vs multiplier_ratio={}",
								test_name, i, spread_ratio, multiplier_ratio
							);
						}
					}
				}



				use crate::indicators::sma::{sma, SmaInput, SmaParams};
				use crate::indicators::ema::{ema, EmaInput, EmaParams};

				let fast_ma = if fast_is_ema {
					let ema_params = EmaParams { period: Some(fast_period) };
					let ema_input = EmaInput::from_slice(&data, ema_params);
					ema(&ema_input).unwrap().values
				} else {
					let sma_params = SmaParams { period: Some(fast_period) };
					let sma_input = SmaInput::from_slice(&data, sma_params);
					sma(&sma_input).unwrap().values
				};

				for i in first_valid_output..data.len() {
					if result.middleband[i].is_finite() && fast_ma[i].is_finite() {
						prop_assert!(
							(result.middleband[i] - fast_ma[i]).abs() <= 1e-9,
							"[{}] Middle band != fast MA at idx {}: {} vs {}",
							test_name, i, result.middleband[i], fast_ma[i]
						);
					}
				}



				for i in first_valid_output..data.len() {
					let upper = result.upperband[i];
					let middle = result.middleband[i];
					let lower = result.lowerband[i];

					if upper.is_finite() && middle.is_finite() && lower.is_finite() {
						let upper_spread = upper - middle;
						let lower_spread = middle - lower;

						prop_assert!(
							upper_spread >= -1e-10,
							"[{}] Negative upper spread at idx {}: {}",
							test_name, i, upper_spread
						);
						prop_assert!(
							lower_spread >= -1e-10,
							"[{}] Negative lower spread at idx {}: {}",
							test_name, i, lower_spread
						);


						let data_range = data.iter()
							.filter(|x| x.is_finite())
							.fold((f64::INFINITY, f64::NEG_INFINITY), |(min, max), &x| {
								(min.min(x), max.max(x))
							});
						let range_span = data_range.1 - data_range.0;



						if range_span > 0.0 {
							prop_assert!(
								upper_spread <= range_span * devup * 10.0,
								"[{}] Upper spread unreasonably large at idx {}: {} (range_span={})",
								test_name, i, upper_spread, range_span
							);
							prop_assert!(
								lower_spread <= range_span * devdn * 10.0,
								"[{}] Lower spread unreasonably large at idx {}: {} (range_span={})",
								test_name, i, lower_spread, range_span
							);
						}
					}
				}



				let invalid_params = vec![
					MabParams {
						fast_period: Some(0),
						slow_period: Some(10),
						..Default::default()
					},
					MabParams {
						fast_period: Some(10),
						slow_period: Some(0),
						..Default::default()
					},
					MabParams {
						fast_period: Some(data.len() + 1),
						slow_period: Some(10),
						..Default::default()
					},
				];

				for invalid_param in invalid_params {
					let invalid_input = MabInput::from_slice(&data, invalid_param);
					let invalid_result = mab_with_kernel(&invalid_input, kernel);
					prop_assert!(
						invalid_result.is_err(),
						"[{}] Expected error for invalid parameters but got Ok",
						test_name
					);
				}




				if fast_period == slow_period && data.len() > first_valid_output {

					let equal_params = MabParams {
						fast_period: Some(fast_period),
						slow_period: Some(fast_period),
						devup: Some(devup),
						devdn: Some(devdn),
						fast_ma_type: params.fast_ma_type.clone(),
						slow_ma_type: params.slow_ma_type.clone(),
					};
					let equal_input = MabInput::from_slice(&data, equal_params);
					let equal_result = mab_with_kernel(&equal_input, kernel).unwrap();



					if params.fast_ma_type == params.slow_ma_type {
						for i in first_valid_output..data.len().min(first_valid_output + 10) {
							if equal_result.upperband[i].is_finite() &&
							   equal_result.middleband[i].is_finite() &&
							   equal_result.lowerband[i].is_finite() {
								let upper_spread = equal_result.upperband[i] - equal_result.middleband[i];
								let lower_spread = equal_result.middleband[i] - equal_result.lowerband[i];


								prop_assert!(
									upper_spread <= 1e-6 || upper_spread <= equal_result.middleband[i].abs() * 1e-6,
									"[{}] Equal periods: upper spread too large at idx {}: {}",
									test_name, i, upper_spread
								);
								prop_assert!(
									lower_spread <= 1e-6 || lower_spread <= equal_result.middleband[i].abs() * 1e-6,
									"[{}] Equal periods: lower spread too large at idx {}: {}",
									test_name, i, lower_spread
								);
							}
						}
					}
				}

				Ok(())
			})?;

        Ok(())
    }

    fn check_mab_partial_params(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;
        let default_params = MabParams {
            fast_period: None,
            ..MabParams::default()
        };
        let input = MabInput::from_candles(&candles, "close", default_params);
        let output = mab_with_kernel(&input, kernel)?;
        assert_eq!(output.upperband.len(), candles.close.len());
        Ok(())
    }

    fn check_mab_accuracy(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;
        let params = MabParams::default();
        let input = MabInput::from_candles(&candles, "close", params);
        let result = mab_with_kernel(&input, kernel)?;

        let expected_upper_last_five = [
            64002.843463352016,
            63976.62699738246,
            63949.00496307154,
            63912.13708526151,
            63828.40371728143,
        ];
        let expected_middle_last_five = [
            59213.90000000002,
            59180.800000000025,
            59161.40000000002,
            59132.00000000002,
            59042.40000000002,
        ];
        let expected_lower_last_five = [
            59350.676536647945,
            59296.93300261751,
            59252.75503692843,
            59190.30291473845,
            59070.11628271853,
        ];

        let len = result.upperband.len();
        for i in 0..5 {
            let idx = len - 5 + i;
            assert!(
                (result.upperband[idx] - expected_upper_last_five[i]).abs() < 1e-4,
                "[{}] Upper band mismatch at index {}: {} vs expected {}",
                test_name,
                i,
                result.upperband[idx],
                expected_upper_last_five[i]
            );
            assert!(
                (result.middleband[idx] - expected_middle_last_five[i]).abs() < 1e-4,
                "[{}] Middle band mismatch at index {}: {} vs expected {}",
                test_name,
                i,
                result.middleband[idx],
                expected_middle_last_five[i]
            );
            assert!(
                (result.lowerband[idx] - expected_lower_last_five[i]).abs() < 1e-4,
                "[{}] Lower band mismatch at index {}: {} vs expected {}",
                test_name,
                i,
                result.lowerband[idx],
                expected_lower_last_five[i]
            );
        }
        Ok(())
    }

    fn check_mab_default_candles(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;
        let input = MabInput::with_default_candles(&candles);
        let output = mab_with_kernel(&input, kernel)?;
        assert_eq!(output.upperband.len(), candles.close.len());
        Ok(())
    }

    fn check_mab_zero_period(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let input_data = [10.0, 20.0, 30.0];
        let params = MabParams {
            fast_period: Some(0),
            slow_period: Some(5),
            ..MabParams::default()
        };
        let input = MabInput::from_slice(&input_data, params);
        let res = mab_with_kernel(&input, kernel);
        assert!(
            res.is_err(),
            "[{}] Expected error for zero fast period",
            test_name
        );

        let params2 = MabParams {
            fast_period: Some(5),
            slow_period: Some(0),
            ..MabParams::default()
        };
        let input2 = MabInput::from_slice(&input_data, params2);
        let res2 = mab_with_kernel(&input2, kernel);
        assert!(
            res2.is_err(),
            "[{}] Expected error for zero slow period",
            test_name
        );
        Ok(())
    }

    fn check_mab_period_exceeds_length(
        test_name: &str,
        kernel: Kernel,
    ) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let data_small = [10.0, 20.0, 30.0];
        let params = MabParams {
            fast_period: Some(2),
            slow_period: Some(10),
            ..MabParams::default()
        };
        let input = MabInput::from_slice(&data_small, params);
        let res = mab_with_kernel(&input, kernel);
        assert!(
            res.is_err(),
            "[{}] Expected error when period exceeds data length",
            test_name
        );
        Ok(())
    }

    fn check_mab_very_small_dataset(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let single_point = [42.0];
        let params = MabParams {
            fast_period: Some(10),
            slow_period: Some(20),
            ..MabParams::default()
        };
        let input = MabInput::from_slice(&single_point, params);
        let res = mab_with_kernel(&input, kernel);
        assert!(
            res.is_err(),
            "[{}] Expected error for insufficient data",
            test_name
        );
        Ok(())
    }

    fn check_mab_reinput(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;
        let params = MabParams::default();
        let first_input = MabInput::from_candles(&candles, "close", params.clone());
        let first_result = mab_with_kernel(&first_input, kernel)?;

        let second_input = MabInput::from_slice(&first_result.upperband, params);
        let second_result = mab_with_kernel(&second_input, kernel)?;
        assert_eq!(second_result.upperband.len(), first_result.upperband.len());

        let non_nan_count = second_result
            .upperband
            .iter()
            .skip(100)
            .filter(|x| !x.is_nan())
            .count();
        assert!(
            non_nan_count > 0,
            "[{}] Second calculation produced all NaN values",
            test_name
        );
        Ok(())
    }

    fn check_mab_nan_handling(test_name: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test_name);
        let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let candles = read_candles_from_csv(file_path)?;
        let input = MabInput::from_candles(&candles, "close", MabParams::default());
        let res = mab_with_kernel(&input, kernel)?;

        for i in 100..res.upperband.len().min(200) {
            assert!(
                !res.upperband[i].is_nan(),
                "[{}] Unexpected NaN in upper band at index {}",
                test_name,
                i
            );
            assert!(
                !res.middleband[i].is_nan(),
                "[{}] Unexpected NaN in middle band at index {}",
                test_name,
                i
            );
            assert!(
                !res.lowerband[i].is_nan(),
                "[{}] Unexpected NaN in lower band at index {}",
                test_name,
                i
            );
        }
        Ok(())
    }

    #[allow(dead_code)]
    fn check_mab_streaming(_test_name: &str, _kernel: Kernel) -> Result<(), Box<dyn Error>> {
        Ok(())
    }

    fn check_batch_default_row(test: &str, kernel: Kernel) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test);
        let file = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let c = read_candles_from_csv(file)?;

        let sweep = MabBatchRange {
            fast_period: (10, 10, 0),
            slow_period: (50, 50, 0),
            devup: (1.0, 1.0, 0.0),
            devdn: (1.0, 1.0, 0.0),
            fast_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
            slow_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
        };

        let output = mab_batch_inner(c.close.as_slice(), &sweep, kernel, false)?;

        assert_eq!(
            output.rows, 1,
            "[{}] Expected 1 row for default params",
            test
        );
        assert_eq!(
            output.cols,
            c.close.len(),
            "[{}] Cols should match input length",
            test
        );

        let expected_upper = [
            64002.843463352016,
            63976.62699738246,
            63949.00496307154,
            63912.13708526151,
            63828.40371728143,
        ];
        let expected_middle = [
            59213.90000000002,
            59180.800000000025,
            59161.40000000002,
            59132.00000000002,
            59042.40000000002,
        ];
        let expected_lower = [
            59350.676536647945,
            59296.93300261751,
            59252.75503692843,
            59190.30291473845,
            59070.11628271853,
        ];

        let start = output.cols - 5;
        for i in 0..5 {
            let idx = start + i;
            assert!(
                (output.upperbands[idx] - expected_upper[i]).abs() < 1e-4,
                "[{}] batch upper mismatch at idx {}: {} vs expected {}",
                test,
                i,
                output.upperbands[idx],
                expected_upper[i]
            );
            assert!(
                (output.middlebands[idx] - expected_middle[i]).abs() < 1e-4,
                "[{}] batch middle mismatch at idx {}: {} vs expected {}",
                test,
                i,
                output.middlebands[idx],
                expected_middle[i]
            );
            assert!(
                (output.lowerbands[idx] - expected_lower[i]).abs() < 1e-4,
                "[{}] batch lower mismatch at idx {}: {} vs expected {}",
                test,
                i,
                output.lowerbands[idx],
                expected_lower[i]
            );
        }
        Ok(())
    }

    fn check_batch_grid_varying_fast_period(
        test: &str,
        kernel: Kernel,
    ) -> Result<(), Box<dyn Error>> {
        skip_if_unsupported!(kernel, test);
        let file = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
        let c = read_candles_from_csv(file)?;

        let sweep = MabBatchRange {
            fast_period: (10, 12, 1),
            slow_period: (50, 50, 0),
            devup: (1.0, 1.0, 0.0),
            devdn: (1.0, 1.0, 0.0),
            fast_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
            slow_ma_type: ("sma".to_string(), "sma".to_string(), String::new()),
        };

        let output = mab_batch_inner(c.close.as_slice(), &sweep, kernel, false)?;

        assert_eq!(
            output.rows, 3,
            "[{}] Expected 3 rows for fast period 10-12",
            test
        );
        assert_eq!(
            output.combos.len(),
            3,
            "[{}] Expected 3 parameter combinations",
            test
        );

        assert_eq!(
            output.combos[0].fast_period,
            Some(10),
            "[{}] First combo fast period",
            test
        );
        assert_eq!(
            output.combos[1].fast_period,
            Some(11),
            "[{}] Second combo fast period",
            test
        );
        assert_eq!(
            output.combos[2].fast_period,
            Some(12),
            "[{}] Third combo fast period",
            test
        );

        for row in 0..3 {
            let row_start = row * output.cols;
            let row_data = &output.upperbands[row_start..row_start + output.cols];

            let valid_count = row_data.iter().skip(100).filter(|x| !x.is_nan()).count();
            assert!(
                valid_count > 0,
                "[{}] Row {} should have valid values",
                test,
                row
            );
        }
        Ok(())
    }

    macro_rules! generate_all_mab_tests {
		($($test_fn:ident),*) => {
			paste::paste! {
				$(
					#[test]
					fn [<$test_fn _scalar_f64>]() {
						let _ = $test_fn(stringify!([<$test_fn _scalar_f64>]), Kernel::Scalar);
					}
				)*
				#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
				$(
					#[test]
					fn [<$test_fn _avx2_f64>]() {
						let _ = $test_fn(stringify!([<$test_fn _avx2_f64>]), Kernel::Avx2);
					}
					#[test]
					fn [<$test_fn _avx512_f64>]() {
						let _ = $test_fn(stringify!([<$test_fn _avx512_f64>]), Kernel::Avx512);
					}
				)*
			}
		}
	}

    macro_rules! gen_batch_tests {
        ($fn_name:ident) => {
            paste::paste! {
                #[test] fn [<$fn_name _scalar>]() {
                    let _ = $fn_name(stringify!([<$fn_name _scalar>]), Kernel::ScalarBatch);
                }
                #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
                #[test] fn [<$fn_name _avx2>]() {
                    let _ = $fn_name(stringify!([<$fn_name _avx2>]), Kernel::Avx2Batch);
                }
                #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
                #[test] fn [<$fn_name _avx512>]() {
                    let _ = $fn_name(stringify!([<$fn_name _avx512>]), Kernel::Avx512Batch);
                }
                #[test] fn [<$fn_name _auto_detect>]() {
                    let _ = $fn_name(stringify!([<$fn_name _auto_detect>]), Kernel::Auto);
                }
            }
        };
    }

    generate_all_mab_tests!(
        check_mab_no_poison,
        check_mab_partial_params,
        check_mab_accuracy,
        check_mab_default_candles,
        check_mab_zero_period,
        check_mab_period_exceeds_length,
        check_mab_very_small_dataset,
        check_mab_reinput,
        check_mab_nan_handling
    );

    #[cfg(feature = "proptest")]
    generate_all_mab_tests!(check_mab_property);

    gen_batch_tests!(check_batch_no_poison);
    gen_batch_tests!(check_batch_default_row);
    gen_batch_tests!(check_batch_grid_varying_fast_period);
}