franken_ocr 0.8.0

//! AArch64 register-blocked int8 GEMM — the primary Apple-Silicon kernel
//! (AGENTS.md doctrine #3/#4; PROPOSED_ARCHITECTURE.md §6.6).
//!
//! This module provides the ARM acceleration tier for the pinned int8 GEMM
//! contract shared by every SIMD backend (`simd/scalar.rs` is the reference
//! oracle; `simd/dispatch.rs` picks the best available tier at runtime):
//!
//! ```ignore
//! // C[M,N] += A[M,K] (i8/u8, row-major) · B[N,K] (i8, OUTPUT-CHANNEL-major) -> i32[M,N]
//! pub fn igemm_s8s8(a: &[i8], b: &[i8], m: usize, k: usize, n: usize, out: &mut [i32]);
//! pub fn igemm_u8s8(a: &[u8], b: &[i8], m: usize, k: usize, n: usize, out: &mut [i32]);
//! ```
//!
//! Two accelerated paths, gated by runtime CPU feature detection:
//!
//! * **SDOT** (`FEAT_DotProd`, `vdotq_s32`): four int8 MACs per i32 lane.
//!   Register-blocked 4×4 output micro-tile; four i32x4 accumulators per output
//!   row reduce K in 16-wide steps. This is the decode-GEMV and small-`M`
//!   workhorse.
//! * **SMMLA** (`FEAT_MATMUL_INT8` / i8mm, `vmmlaq_s32`): eight int8 MACs per
//!   i32 lane, computing a 2×2 i32 tile from two `[2×8]` int8 operands. Doctrine
//!   #4: an *un-blocked* SMMLA is load-bound and SLOWER than SDOT, so we
//!   **pre-pack** both A and B into the SMMLA-friendly interleaved `[2 rows × 8
//!   cols]` panel layout and **register-block** an 8×8 output tile (sixteen
//!   `int32x4` accumulators, each fed by `2×2 = 4` SMMLA-tile contributions per
//!   K-step), reaching compute:load ≥ 2:1 on the inner loop.
//!
//! Both paths produce **bit-identical** output to the scalar reference (i32
//! accumulation is exact and order-independent for integer addition), proven by
//! the `#[cfg(test)]` cross-check below over randomized + all-127 / all-(-128)
//! adversarial operands at K including the doctrine-#6 worst case 6848.
//!
//! ## U8S8
//!
//! The i8mm/dotprod intrinsics are signed-by-signed (`vdotq_s32` /
//! `vmmlaq_s32`). U8S8 (`DynamicQuantizeLinear` asymmetric activation × signed
//! weight) is handled by the **+128 bias-correction** identity, computed
//! entirely in i32 (exact, no saturation):
//!
//! ```text
//!   u8 activation a_u = (s8 activation a_s) + 128   with a_s = a_u - 128
//!   sum_k a_u[k]·w[k] = sum_k (a_s[k]+128)·w[k]
//!                     = (sum_k a_s[k]·w[k])  +  128 · (sum_k w[k])
//! ```
//!
//! So we run the signed kernel on `a_s = a_u - 128` (a lossless `i16`-range
//! shift folded back into i8 by reinterpreting the byte, since `a_u - 128`
//! lands in `[-128, 127]`), then add `128 · rowsum(w)` per output channel. The
//! per-channel weight row-sum is computed once. This matches the scalar U8S8
//! oracle exactly (the overflow proof in `tests/int32_overflow_proof.rs` bounds
//! the U8S8 accumulator at 221.7M < i32::MAX at K=6848).
//!
//! ## Safety
//!
//! All raw NEON intrinsics live in the single `#[allow(unsafe_code,
//! unsafe_op_in_unsafe_fn)]` island below. Every vector load reads a contiguous
//! 16-byte window from a slice whose length the caller has already bounded; each
//! load carries a `// SAFETY:` note proving the read is in-bounds and that the
//! enclosing function is only reached when the corresponding CPU feature was
//! runtime-detected. NEON loads are alignment-agnostic (`vld1q` does not require
//! 16-byte alignment).
//!
//! On non-AArch64 targets the whole accelerated body is `#[cfg]`-compiled out
//! and the public entrypoints delegate to the scalar reference, so the file
//! compiles and is correct everywhere.

#![allow(unsafe_code, unsafe_op_in_unsafe_fn)]

use super::scalar;

// ─────────────────────────────────────────────────────────────────────────────
// Public entrypoints (pinned signature; identical on every arch).
// ─────────────────────────────────────────────────────────────────────────────

/// Runtime ISA tier this ARM backend can offer for the current CPU.
///
/// `dispatch.rs` consults the booleans; the kernels themselves re-check the
/// feature before entering their `#[target_feature]` island so the unsafe entry
/// is never reachable on a CPU lacking the instruction.
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub enum ArmTier {
    /// `FEAT_MATMUL_INT8` (i8mm / SMMLA) available — the top ARM tier.
    Smmla,
    /// `FEAT_DotProd` (SDOT) available, but not i8mm.
    Sdot,
    /// Neither — caller should fall back to the scalar reference.
    None,
}

/// The implementation actually executed by the dense int8 GEMM entrypoints.
/// This is deliberately separate from [`ArmTier`], which describes hardware
/// capability: Apple hosts can expose SDOT while the measured winner is the
/// LLVM-autovectorized scalar loop.
#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
pub(crate) enum DenseI8Route {
    Autovec,
    Scalar,
    Sdot,
    Smmla,
}

/// Detect the int8 tier this CPU should dispatch to (cached after first call).
///
/// On non-AArch64 builds this is always [`ArmTier::None`].
///
/// **Apple Silicon prefers SDOT over SMMLA.** Every Apple aarch64 core (M-series
/// on macOS, A-series on iOS)
/// issues `SMMLA`/i8mm at *half* the rate of `SDOT`, so SMMLA's 2× MACs per
/// instruction exactly cancel — measured on an Apple M4, SMMLA delivers 0.994× of
/// SDOT's int8 MACs/second (`vmmlaq_s32` 5.31 Ginstr/s vs `vdotq_s32`
/// 10.69 Ginstr/s) — and SMMLA *additionally* needs a 2×2 operand repack the dot
/// path skips, so SDOT is the strictly faster kernel here. On other aarch64 cores
/// (e.g. Neoverse) i8mm can be full-rate, so there SMMLA leads.
///
/// `FOCR_FORCE_ARCH=smmla|sdot|scalar` overrides the choice (benchmark/debug only;
/// a tier whose feature is absent is ignored). Read once and cached.
#[must_use]
/// Whether the DENSE int8 GEMM entrypoints should prefer the LLVM-autovec
/// scalar kernel over the hand-written SDOT micro-tile on Apple Silicon
/// (bd-2mo.5.1, measured 2026-07-09 in a quiet window): the autovectorized
/// loop beats the SDOT kernel across every model shape in the tier
/// microbench (~4x at m=1) and 20.5% on the REAL int8 decode e2e
/// (page_0009: 1.27s vs 1.59s, interleaved runs, identical 98-token
/// output). This is NE-INH-1's inherited prior ("hand-written wide-SIMD
/// int8 dot ~5x slower than autovec") catching up with the m=1 GEMV — the
/// SDOT tile pays per-call packing/setup the fused autovec loop never does.
///
/// SCOPED to `igemm_s8s8`/`igemm_u8s8` ONLY: `detect_tier()` still reports
/// SDOT as the hardware tier while the effective route is reported separately,
/// the int4 packed path keeps its SDOT kernels (its scalar fallback is the
/// ledgered 5.8x-slower unpack path), and the offline-SMMLA packed-B path
/// is unaffected. `FOCR_INT8_AUTOVEC=0` restores the SDOT dispatch.
///
/// Gated on `target_vendor = "apple"` rather than `target_os = "macos"`: the
/// measurement is about how an Apple core issues the SDOT micro-tile versus the
/// fused autovec loop, which is a core property shared by M-series and A-series.
/// iOS re-measurement is tracked separately; until it lands, A-series inherits
/// the M-series route rather than silently taking the un-measured SDOT branch.
fn autovec_preferred() -> bool {
    use std::sync::OnceLock;
    static PREF: OnceLock<bool> = OnceLock::new();
    *PREF.get_or_init(|| {
        cfg!(target_vendor = "apple")
            && !matches!(
                std::env::var("FOCR_INT8_AUTOVEC").ok().as_deref(),
                Some("0") | Some("off") | Some("false") | Some("no")
            )
            // An honored tier override (bench A/Bs, the selftest sweep) must
            // keep meaning what it says. Unknown or unavailable tags are
            // ignored without silently switching autovec off.
            && !force_arch_honored()
    })
}

fn force_arch_honored() -> bool {
    let Ok(force) = std::env::var("FOCR_FORCE_ARCH") else {
        return false;
    };
    matches!(
        (force.trim().to_ascii_lowercase().as_str(), detect_tier()),
        ("smmla", ArmTier::Smmla) | ("sdot", ArmTier::Sdot) | ("scalar", ArmTier::None)
    )
}

/// Effective implementation selected for ordinary row-major dense int8 GEMM.
/// Packed int4 and offline-SMMLA-panel entrypoints have separate dispatch and
/// must not use this label.
#[must_use]
pub(crate) fn effective_dense_route() -> DenseI8Route {
    match detect_tier() {
        ArmTier::Smmla => DenseI8Route::Smmla,
        ArmTier::Sdot if autovec_preferred() => DenseI8Route::Autovec,
        ArmTier::Sdot => DenseI8Route::Sdot,
        ArmTier::None => DenseI8Route::Scalar,
    }
}

pub fn detect_tier() -> ArmTier {
    use std::sync::OnceLock;
    static TIER: OnceLock<ArmTier> = OnceLock::new();
    *TIER.get_or_init(detect_tier_uncached)
}

fn detect_tier_uncached() -> ArmTier {
    #[cfg(target_arch = "aarch64")]
    {
        let has_i8mm = std::arch::is_aarch64_feature_detected!("i8mm");
        let has_dotprod = std::arch::is_aarch64_feature_detected!("dotprod");

        // Optional override (benchmark/debug). Only honor a present feature.
        if let Ok(force) = std::env::var("FOCR_FORCE_ARCH") {
            match force.trim().to_ascii_lowercase().as_str() {
                "smmla" if has_i8mm => return ArmTier::Smmla,
                "sdot" if has_dotprod => return ArmTier::Sdot,
                "scalar" => return ArmTier::None,
                _ => {}
            }
        }

        // Apple Silicon: SDOT > SMMLA (i8mm is half-rate; see fn docs).
        // `target_vendor`, not `target_os`: the half-rate i8mm issue is a
        // property of the Apple core, so A-series (iOS) must land here too.
        #[cfg(target_vendor = "apple")]
        {
            if has_dotprod {
                return ArmTier::Sdot;
            }
            if has_i8mm {
                return ArmTier::Smmla;
            }
        }
        // Other aarch64: SMMLA > SDOT (i8mm may be full-rate).
        #[cfg(not(target_vendor = "apple"))]
        {
            if has_i8mm {
                return ArmTier::Smmla;
            }
            if has_dotprod {
                return ArmTier::Sdot;
            }
        }
        ArmTier::None
    }
    #[cfg(not(target_arch = "aarch64"))]
    {
        ArmTier::None
    }
}

/// `out[M,N] += A[M,K] (i8, row-major) · B[N,K] (i8, output-channel-major)`,
/// accumulating in i32 and adding into the caller-provided `out` buffer,
/// matching the scalar oracle's contract.
///
/// Picks the best available ARM tier via [`detect_tier`] (SDOT > SMMLA > scalar
/// on Apple Silicon, SMMLA > SDOT > scalar elsewhere) and produces output
/// bit-identical to [`scalar::igemm_s8s8`] on every tier.
///
/// # Panics
/// Panics if `a.len() != m*k`, `b.len() != n*k`, or `out.len() != m*n` (a
/// shape/contract violation is a programming error, caught early — matches the
/// scalar reference's asserts).
pub fn igemm_s8s8(a: &[i8], b: &[i8], m: usize, k: usize, n: usize, out: &mut [i32]) {
    let _ = igemm_s8s8_with_route(a, b, m, k, n, out);
}

/// [`igemm_s8s8`] plus the implementation branch that actually produced the
/// result. Used by `robot selftest` so route evidence is execution-derived.
pub(crate) fn igemm_s8s8_with_route(
    a: &[i8],
    b: &[i8],
    m: usize,
    k: usize,
    n: usize,
    out: &mut [i32],
) -> DenseI8Route {
    let a_len = scalar::checked_len("igemm_s8s8", m, k, "m*k");
    let b_len = scalar::checked_len("igemm_s8s8", n, k, "n*k");
    let out_len = scalar::checked_len("igemm_s8s8", m, n, "m*n");
    assert_eq!(
        a.len(),
        a_len,
        "igemm_s8s8: a.len {} != m*k {}",
        a.len(),
        a_len
    );
    assert_eq!(
        b.len(),
        b_len,
        "igemm_s8s8: b.len {} != n*k {}",
        b.len(),
        b_len
    );
    assert_eq!(
        out.len(),
        out_len,
        "igemm_s8s8: out.len {} != m*n {}",
        out.len(),
        out_len
    );

    #[cfg(target_arch = "aarch64")]
    {
        match detect_tier() {
            ArmTier::Smmla => {
                // SAFETY: reached only when `is_aarch64_feature_detected!("i8mm")`
                // returned true in `detect_tier`, so SMMLA/i8mm is present.
                unsafe { aarch64_impl::igemm_s8s8_smmla(a, b, m, k, n, out) };
                return DenseI8Route::Smmla;
            }
            ArmTier::Sdot => {
                if !autovec_preferred() {
                    // SAFETY: reached only when `dotprod` was runtime-detected.
                    unsafe { aarch64_impl::igemm_s8s8_sdot(a, b, m, k, n, out) };
                    return DenseI8Route::Sdot;
                }
                // fall through: the autovec scalar loop wins on this host
                // (bd-2mo.5.1; bit-identical — integer math, exact).
            }
            ArmTier::None => {}
        }
    }

    scalar::igemm_s8s8(a, b, m, k, n, out);
    match detect_tier() {
        // The SDOT branch returned above unless autovec was selected, so this
        // fallthrough itself is the execution evidence.
        ArmTier::Sdot => DenseI8Route::Autovec,
        ArmTier::None => DenseI8Route::Scalar,
        ArmTier::Smmla => unreachable!("SMMLA branch must return before scalar fallback"),
    }
}

/// S8S8 GEMM whose B is an OFFLINE SMMLA panel stream (`focr convert --arch
/// aarch64-smmla`, bd-2mo.3): consumed with zero runtime shuffle on the SMMLA
/// tier; any other tier un-permutes and runs the ordinary row-major path
/// (bit-identical — the packing is a pure zero-padded permutation).
///
/// Unlike [`igemm_s8s8`], `out` is ZEROED here (the accumulate-vs-overwrite
/// contract is owned by this wrapper, not the caller).
///
/// # Panics
/// On length mismatches (`a != m*k`, `b_panels != ceil(n/2)*ceil(k/8)*16`,
/// `out != m*n`).
pub fn igemm_s8s8_packed_b(
    a: &[i8],
    b_panels: &[i8],
    m: usize,
    k: usize,
    n: usize,
    out: &mut [i32],
) {
    let a_len = scalar::checked_len("igemm_s8s8_packed_b", m, k, "m*k");
    let panels_len = crate::simd::pack::smmla_packed_len(n, k);
    let out_len = scalar::checked_len("igemm_s8s8_packed_b", m, n, "m*n");
    assert_eq!(
        a.len(),
        a_len,
        "igemm_s8s8_packed_b: a.len {} != m*k {}",
        a.len(),
        a_len
    );
    assert_eq!(
        b_panels.len(),
        panels_len,
        "igemm_s8s8_packed_b: b_panels.len {} != ceil(n/2)*ceil(k/8)*16 {}",
        b_panels.len(),
        panels_len
    );
    assert_eq!(
        out.len(),
        out_len,
        "igemm_s8s8_packed_b: out.len {} != m*n {}",
        out.len(),
        out_len
    );
    out.fill(0);

    #[cfg(target_arch = "aarch64")]
    if detect_tier() == ArmTier::Smmla {
        // SAFETY: reached only when `is_aarch64_feature_detected!("i8mm")`
        // returned true in `detect_tier`; slice lengths asserted above.
        unsafe { aarch64_impl::igemm_s8s8_smmla_packed_b(a, b_panels, m, k, n, out) };
        return;
    }

    // Degrade path (non-SMMLA tier handed a packed artifact): un-permute and
    // run the ordinary dispatch — never the hot path (the loader only keeps
    // panels when SMMLA is dispatched), always correct.
    let b = crate::simd::pack::smmla_unpack_panels(b_panels, n, k).expect("length asserted above");
    igemm_s8s8(a, &b, m, k, n, out);
}

/// U8S8 variant: `A` is `u8` (asymmetric activations), `B` is `i8` weights.
/// Output bit-identical to [`scalar::igemm_u8s8`].
///
/// Implemented via the +128 bias-correction identity (see module docs): run the
/// signed kernel on `a - 128`, then add `128 · rowsum(b)` per output channel.
///
/// # Panics
/// As [`igemm_s8s8`].
pub fn igemm_u8s8(a: &[u8], b: &[i8], m: usize, k: usize, n: usize, out: &mut [i32]) {
    let _ = igemm_u8s8_with_route(a, b, m, k, n, out);
}

/// [`igemm_u8s8`] plus the implementation branch that actually produced the
/// result. Used by `robot selftest` alongside the signed-domain trace.
pub(crate) fn igemm_u8s8_with_route(
    a: &[u8],
    b: &[i8],
    m: usize,
    k: usize,
    n: usize,
    out: &mut [i32],
) -> DenseI8Route {
    let a_len = scalar::checked_len("igemm_u8s8", m, k, "m*k");
    let b_len = scalar::checked_len("igemm_u8s8", n, k, "n*k");
    let out_len = scalar::checked_len("igemm_u8s8", m, n, "m*n");
    assert_eq!(
        a.len(),
        a_len,
        "igemm_u8s8: a.len {} != m*k {}",
        a.len(),
        a_len
    );
    assert_eq!(
        b.len(),
        b_len,
        "igemm_u8s8: b.len {} != n*k {}",
        b.len(),
        b_len
    );
    assert_eq!(
        out.len(),
        out_len,
        "igemm_u8s8: out.len {} != m*n {}",
        out.len(),
        out_len
    );

    #[cfg(target_arch = "aarch64")]
    {
        let tier = detect_tier();
        if matches!(tier, ArmTier::Smmla) || (matches!(tier, ArmTier::Sdot) && !autovec_preferred())
        {
            // Shift u8 activations into the signed domain: a_s = a_u - 128 lands
            // in [-128, 127], representable in i8. `wrapping_sub(128) as i8`
            // reinterprets the byte so that e.g. 0u8 -> -128i8, 255u8 -> 127i8.
            let a_signed: Vec<i8> = a.iter().map(|&x| x.wrapping_sub(128) as i8).collect();
            // Run the signed GEMM into `out`.
            match tier {
                ArmTier::Smmla => {
                    // SAFETY: i8mm runtime-detected (see `detect_tier`).
                    unsafe { aarch64_impl::igemm_s8s8_smmla(&a_signed, b, m, k, n, out) };
                }
                ArmTier::Sdot => {
                    // SAFETY: dotprod runtime-detected.
                    unsafe { aarch64_impl::igemm_s8s8_sdot(&a_signed, b, m, k, n, out) };
                }
                ArmTier::None => unreachable!(),
            }
            // Per-output-channel weight row-sum, then fold in the +128·rowsum(w)
            // correction (exact i32; the U8S8 overflow bound 221.7M holds).
            let mut rowsum = vec![0i32; n];
            for (oc, rs) in rowsum.iter_mut().enumerate() {
                let row = &b[oc * k..(oc + 1) * k];
                let mut s: i32 = 0;
                for &w in row {
                    s += i32::from(w);
                }
                *rs = s;
            }
            for r in 0..m {
                let orow = &mut out[r * n..(r + 1) * n];
                for (c, cell) in orow.iter_mut().enumerate() {
                    *cell += 128 * rowsum[c];
                }
            }
            return match tier {
                ArmTier::Smmla => DenseI8Route::Smmla,
                ArmTier::Sdot => DenseI8Route::Sdot,
                ArmTier::None => unreachable!(),
            };
        }
    }

    scalar::igemm_u8s8(a, b, m, k, n, out);
    match detect_tier() {
        // The intrinsic branch returned above unless autovec was selected.
        ArmTier::Sdot => DenseI8Route::Autovec,
        ArmTier::None => DenseI8Route::Scalar,
        ArmTier::Smmla => unreachable!("SMMLA branch must return before scalar fallback"),
    }
}

/// Native packed-int4 GEMM (bd-1azu.22) — the ARM tier dispatch.
///
/// Routes to the SDOT or SMMLA **packed** kernel (each consumes the packed int4
/// nibbles *directly* — mask/shift in-register, no dense-i8 materialization, the
/// distinction from the `unpack-then-int8` `int4::igemm_s4s8`) per [`detect_tier`]
/// (SDOT on Apple Silicon; SMMLA where i8mm is full-rate), or to the scalar
/// packed reference for [`ArmTier::None`]. Output is bit-identical to
/// [`crate::simd::int4::igemm_s4s8_packed_scalar`] and **overwrites** `out`
/// (`= C`, like `igemm_s4s8`, NOT the `+=` of the int8 kernels).
///
/// Shape/contract preconditions are asserted by the public caller
/// [`crate::simd::int4::igemm_s4s8_packed`]; this function only routes.
#[allow(clippy::too_many_arguments)]
pub fn igemm_s4s8_packed(
    a: &[i8],
    b_packed: &[u8],
    scales: &[f32],
    group: usize,
    m: usize,
    k: usize,
    n: usize,
    out: &mut [f32],
) {
    match detect_tier() {
        ArmTier::Sdot => {
            // SAFETY: reached only when `dotprod` was runtime-detected in
            // `detect_tier`; the buffers satisfy the packed GEMM shape (caller
            // asserted in `int4::igemm_s4s8_packed`).
            unsafe {
                aarch64_impl::igemm_s4s8_packed_sdot(a, b_packed, scales, group, m, k, n, out);
            }
        }
        ArmTier::Smmla => {
            // SAFETY: reached only when `i8mm` was runtime-detected; shapes valid.
            unsafe {
                aarch64_impl::igemm_s4s8_packed_smmla(a, b_packed, scales, group, m, k, n, out);
            }
        }
        ArmTier::None => {
            super::int4::s4s8_packed_kernel_scalar(a, b_packed, scales, group, m, k, n, out);
        }
    }
}

// ─────────────────────────────────────────────────────────────────────────────
// The audited unsafe SIMD island.
//
// Everything that touches a raw NEON intrinsic lives here. The crate root is
// `#![deny(unsafe_code)]`; this island re-enables it locally with an explicit
// allow + per-load SAFETY notes (AGENTS.md Toolchain rule).
// ─────────────────────────────────────────────────────────────────────────────
#[cfg(target_arch = "aarch64")]
#[allow(unsafe_code, unsafe_op_in_unsafe_fn)]
mod aarch64_impl {
    use core::arch::aarch64::{
        int8x8_t, int8x16_t, int32x4_t, vaddvq_s32, vand_u8, vcombine_s8, vdotq_s32, vdup_n_u8,
        vdupq_n_s32, vld1_u8, vld1q_s8, vld2_s8, vmmlaq_s32, vreinterpret_s8_u8, vshl_n_s8,
        vshr_n_s8, vshr_n_u8, vzip_s8,
    };

    /// Load a 16-byte window `[off, off+16)` of `s` as an `int8x16_t`.
    ///
    /// # Safety
    /// The caller MUST guarantee `off + 16 <= s.len()`. `vld1q_s8` performs an
    /// unaligned 16-byte read; NEON loads do not require 16-byte alignment.
    #[inline]
    #[target_feature(enable = "neon")]
    unsafe fn load16(s: &[i8], off: usize) -> int8x16_t {
        debug_assert!(off + 16 <= s.len(), "load16 out of bounds");
        // SAFETY: caller guarantees off+16 <= len; pointer is valid for a
        // 16-byte unaligned read within the slice's allocation.
        vld1q_s8(s.as_ptr().add(off))
    }

    /// Load a 16-byte window that may run off the end of `s`, zero-filling the
    /// tail. Used only on the K-remainder step so a partial 16-lane block reads
    /// no out-of-bounds memory yet contributes the correct (zero-padded) dot.
    ///
    /// Zero padding is exact for an int dot product: a zero lane contributes a
    /// zero product, so the accumulator is identical to summing only the real
    /// `len-off` elements.
    #[inline]
    #[target_feature(enable = "neon")]
    unsafe fn load16_tail(s: &[i8], off: usize, valid: usize) -> int8x16_t {
        let mut buf = [0i8; 16];
        let n = valid.min(16);
        // `off + n <= s.len()` by construction (valid = len - off, clamped).
        buf[..n].copy_from_slice(&s[off..off + n]);
        // SAFETY: `buf` is a fully-initialized 16-byte stack array; the read is
        // in-bounds and (stack) suitably accessible for an unaligned vld1q.
        vld1q_s8(buf.as_ptr())
    }

    // ── SDOT path ───────────────────────────────────────────────────────────
    //
    // `vdotq_s32(acc, a, b)` computes, for each of the 4 i32 lanes `j`:
    //     acc[j] += sum_{t=0..3} a[4j+t] * b[4j+t]
    // i.e. 4 int8 MACs per lane, 16 MACs per instruction. To get the full
    // 16-element dot of two int8x16 vectors we `vdotq_s32` into a 4-lane
    // accumulator and horizontally add the lanes (`vaddvq_s32`) at the end —
    // integer addition is associative, so this is bit-identical to the scalar
    // left-to-right sum.

    /// SDOT int8 GEMM. Register-blocked over a 4×4 output micro-tile: four
    /// activation rows × four weight rows share their 16-wide K loads, giving
    /// sixteen independent i32x4 accumulators per tile (compute:load ≈ 16 dots
    /// per 8 loads = 2:1).
    ///
    /// # Safety
    /// Caller MUST ensure `dotprod` (`FEAT_DotProd`) is available on this CPU
    /// and the slice lengths satisfy the GEMM shape (`a=m*k`, `b=n*k`,
    /// `out=m*n`).
    #[target_feature(enable = "neon,dotprod")]
    pub unsafe fn igemm_s8s8_sdot(
        a: &[i8],
        b: &[i8],
        m: usize,
        k: usize,
        n: usize,
        out: &mut [i32],
    ) {
        let k16 = k / 16 * 16; // K rounded down to a multiple of 16

        let mut r = 0;
        while r < m {
            let mr = (m - r).min(4);
            let mut c = 0;
            while c < n {
                let nr = (n - c).min(4);

                // 4×4 accumulator grid of i32x4 lanes.
                // SAFETY: `vdupq_n_s32` is a pure constant splat (no memory).
                let mut acc = [[vdupq_n_s32(0); 4]; 4];

                let mut t = 0;
                while t < k16 {
                    // Load up to 4 activation rows and 4 weight rows for this
                    // 16-wide K slice; these loads are SHARED across the tile
                    // (the register-blocking win).
                    let mut av = [vdupq_n_s32(0).into_i8(); 4];
                    let mut bv = [vdupq_n_s32(0).into_i8(); 4];
                    // indexed loop mirrors SIMD lane layout
                    #[allow(clippy::needless_range_loop)]
                    for i in 0..mr {
                        // SAFETY: t+16 <= k16 <= k, and (r+i)<m, so
                        // (r+i)*k + t + 16 <= (r+i+1)*k <= m*k = a.len().
                        av[i] = load16(a, (r + i) * k + t);
                    }
                    // indexed loop mirrors SIMD lane layout
                    #[allow(clippy::needless_range_loop)]
                    for j in 0..nr {
                        // SAFETY: same bound for the weight matrix (b.len = n*k).
                        bv[j] = load16(b, (c + j) * k + t);
                    }
                    for i in 0..mr {
                        for j in 0..nr {
                            // SAFETY: vdotq_s32 is register-only; inputs valid.
                            acc[i][j] = vdotq_s32(acc[i][j], av[i], bv[j]);
                        }
                    }
                    t += 16;
                }

                // K tail ( < 16 elements ): one zero-padded block.
                if k16 < k {
                    let valid = k - k16;
                    let mut av = [vdupq_n_s32(0).into_i8(); 4];
                    let mut bv = [vdupq_n_s32(0).into_i8(); 4];
                    // indexed loop mirrors SIMD lane layout
                    #[allow(clippy::needless_range_loop)]
                    for i in 0..mr {
                        // SAFETY: load16_tail reads only s[off..off+valid] into a
                        // local 16-byte buffer; off+valid = (r+i)*k + k <= a.len.
                        av[i] = load16_tail(a, (r + i) * k + k16, valid);
                    }
                    // indexed loop mirrors SIMD lane layout
                    #[allow(clippy::needless_range_loop)]
                    for j in 0..nr {
                        bv[j] = load16_tail(b, (c + j) * k + k16, valid);
                    }
                    for i in 0..mr {
                        for j in 0..nr {
                            acc[i][j] = vdotq_s32(acc[i][j], av[i], bv[j]);
                        }
                    }
                }

                // Horizontal-reduce each i32x4 accumulator into the output cell.
                for i in 0..mr {
                    for j in 0..nr {
                        // SAFETY: vaddvq_s32 is a register-only horizontal add.
                        out[(r + i) * n + (c + j)] += vaddvq_s32(acc[i][j]);
                    }
                }

                c += nr;
            }
            r += mr;
        }
    }

    // ── SMMLA path ──────────────────────────────────────────────────────────
    //
    // `vmmlaq_s32(acc, a, b)` treats its int8x16 operands as 2×8 int8 matrices
    // (row-major: lanes 0..8 = row 0, lanes 8..16 = row 1) and computes the
    // 2×8 · 8×2 = 2×2 i32 product, ADDED into the int32x4 accumulator laid out
    //     [ c00, c01, c10, c11 ]
    // where c_xy = sum_{t=0..7} a_row_x[t] * b_row_y[t].
    //
    // 8 int8 MACs per i32 lane × 4 lanes = 32 MACs per instruction. Doctrine #4:
    // an un-blocked SMMLA is load-bound. We PRE-PACK A and B into contiguous
    // `[2 rows × 8 cols]` 16-byte panels (so each vld1q feeds exactly one SMMLA
    // operand) and register-block an 8×8 output tile: 4 A-panels (rows {0,1},
    // {2,3}, {4,5}, {6,7}) × 4 B-panels per 8-wide K-step drive 16 int32x4
    // accumulators with 16 SMMLA per 8 loads = 2:1 compute:load.

    /// Pre-pack a region of a row-major `[rows, k]` matrix into SMMLA panels.
    ///
    /// Delegates to the portable single source of truth
    /// ([`crate::simd::pack::smmla_pack_panels`]) so the runtime kernel, the
    /// offline `focr convert --arch aarch64-smmla` pre-packer, and the
    /// loader's un-permute fallback can never drift (bd-2mo.3).
    fn pack_panels(
        src: &[i8],
        base_row: usize,
        rows: usize,
        k: usize,
        src_k: usize,
    ) -> (Vec<i8>, usize, usize) {
        crate::simd::pack::smmla_pack_panels(src, base_row, rows, k, src_k)
    }

    /// SMMLA int8 GEMM with offline-style A/B pre-packing and 8×8 register
    /// blocking.
    ///
    /// # Safety
    /// Caller MUST ensure `i8mm` (`FEAT_MATMUL_INT8`) is available and the slice
    /// lengths satisfy the GEMM shape.
    #[target_feature(enable = "neon,i8mm")]
    pub unsafe fn igemm_s8s8_smmla(
        a: &[i8],
        b: &[i8],
        m: usize,
        k: usize,
        n: usize,
        out: &mut [i32],
    ) {
        // Process the output in 8×8 tiles. For each tile we pack the relevant
        // A rows and B rows into panel streams, then run the SMMLA micro-kernel.
        let mut r = 0;
        while r < m {
            let mr = (m - r).min(8);
            // Pack A rows [r, r+mr) -> ceil(mr/2) row-pairs, kb K-blocks.
            let (apack, a_pairs, kb) = pack_panels(a, r, mr, k, k);

            let mut c = 0;
            while c < n {
                let nr = (n - c).min(8);
                // Pack B rows [c, c+nr) (output-channel-major == row-major [n,k]).
                let (bpack, b_pairs, _kb2) = pack_panels(b, c, nr, k, k);

                // For each (A row-pair p, B row-pair q) we accumulate a 4×4
                // i32 region (2 A-rows × 2 B-rows) across all K-blocks. The
                // int32x4 result is [c(2p,2q), c(2p,2q+1), c(2p+1,2q), c(2p+1,2q+1)].
                for p in 0..a_pairs {
                    for q in 0..b_pairs {
                        // SAFETY: constant splat, no memory.
                        let mut acc = vdupq_n_s32(0);
                        for block in 0..kb {
                            let aoff = (p * kb + block) * 16;
                            let boff = (q * kb + block) * 16;
                            // SAFETY: apack/bpack lengths are pairs*kb*16 by
                            // construction; aoff+16 <= apack.len(), boff+16 <=
                            // bpack.len(). Reads are into our own packed Vecs.
                            let av = load16(&apack, aoff);
                            let bv = load16(&bpack, boff);
                            // SAFETY: vmmlaq_s32 is register-only; i8mm enabled.
                            acc = vmmlaq_s32(acc, av, bv);
                        }
                        // Scatter the 2×2 i32 tile into `out`, skipping padded
                        // rows/cols (row-pair p covers output rows r+2p, r+2p+1).
                        let tile = [
                            vgetq_lane0(acc),
                            vgetq_lane1(acc),
                            vgetq_lane2(acc),
                            vgetq_lane3(acc),
                        ];
                        let ar0 = 2 * p;
                        let ar1 = 2 * p + 1;
                        let bc0 = 2 * q;
                        let bc1 = 2 * q + 1;
                        if ar0 < mr && bc0 < nr {
                            out[(r + ar0) * n + (c + bc0)] += tile[0];
                        }
                        if ar0 < mr && bc1 < nr {
                            out[(r + ar0) * n + (c + bc1)] += tile[1];
                        }
                        if ar1 < mr && bc0 < nr {
                            out[(r + ar1) * n + (c + bc0)] += tile[2];
                        }
                        if ar1 < mr && bc1 < nr {
                            out[(r + ar1) * n + (c + bc1)] += tile[3];
                        }
                    }
                }
                c += nr;
            }
            r += mr;
        }
    }

    /// SMMLA int8 GEMM whose **B operand is already in offline panel layout**
    /// (`focr convert --arch aarch64-smmla`, bd-2mo.3): `b_panels` is the
    /// full-matrix `[ceil(n/2)][ceil(k/8)][16]` stream from
    /// [`crate::simd::pack::smmla_pack_panels`]. Identical arithmetic to
    /// [`igemm_s8s8_smmla`] — the ONLY difference is that the per-call
    /// `pack_panels(b, ..)` is replaced by direct panel loads (zero runtime
    /// shuffle), so the i32 output is bit-identical by construction.
    ///
    /// # Safety
    /// Caller MUST ensure `i8mm` is available, `b_panels.len() ==
    /// smmla_packed_len(n, k)`, and the remaining slices satisfy the GEMM
    /// shape.
    #[target_feature(enable = "neon,i8mm")]
    pub unsafe fn igemm_s8s8_smmla_packed_b(
        a: &[i8],
        b_panels: &[i8],
        m: usize,
        k: usize,
        n: usize,
        out: &mut [i32],
    ) {
        let kb = k.div_ceil(8);
        let mut r = 0;
        while r < m {
            let mr = (m - r).min(8);
            // Pack A rows [r, r+mr) -> ceil(mr/2) row-pairs, kb K-blocks
            // (activations change per call; only B's pack is offline).
            let (apack, a_pairs, _kb) = pack_panels(a, r, mr, k, k);

            let mut c = 0;
            while c < n {
                let nr = (n - c).min(8);
                // B panels for output rows [c, c+nr): pair-aligned because c
                // steps by 8 — the region IS a slice of the offline stream.
                let b_base = (c / 2) * kb;
                let b_pairs = nr.div_ceil(2);

                for p in 0..a_pairs {
                    for q in 0..b_pairs {
                        // SAFETY: constant splat, no memory.
                        let mut acc = vdupq_n_s32(0);
                        for block in 0..kb {
                            let aoff = (p * kb + block) * 16;
                            let boff = ((b_base + q * kb) + block) * 16;
                            // SAFETY: apack len is a_pairs*kb*16 by
                            // construction; b_panels len is
                            // ceil(n/2)*kb*16 (caller contract) and
                            // b_base + q*kb + block < ceil(n/2)*kb.
                            let av = load16(&apack, aoff);
                            let bv = load16(b_panels, boff);
                            // SAFETY: vmmlaq_s32 is register-only; i8mm enabled.
                            acc = vmmlaq_s32(acc, av, bv);
                        }
                        let tile = [
                            vgetq_lane0(acc),
                            vgetq_lane1(acc),
                            vgetq_lane2(acc),
                            vgetq_lane3(acc),
                        ];
                        let ar0 = 2 * p;
                        let ar1 = 2 * p + 1;
                        let bc0 = 2 * q;
                        let bc1 = 2 * q + 1;
                        if ar0 < mr && bc0 < nr {
                            out[(r + ar0) * n + (c + bc0)] += tile[0];
                        }
                        if ar0 < mr && bc1 < nr {
                            out[(r + ar0) * n + (c + bc1)] += tile[1];
                        }
                        if ar1 < mr && bc0 < nr {
                            out[(r + ar1) * n + (c + bc0)] += tile[2];
                        }
                        if ar1 < mr && bc1 < nr {
                            out[(r + ar1) * n + (c + bc1)] += tile[3];
                        }
                    }
                }
                c += nr;
            }
            r += mr;
        }
    }

    // ── tiny lane-extraction helpers (register-only) ────────────────────────

    // These lane reads are register-only and side-effect-free. Inside a
    // `#[target_feature]` fn (edition 2024) the intrinsic call needs no inner
    // `unsafe` block — the function's target-feature contract is what gates it.
    #[inline]
    #[target_feature(enable = "neon")]
    fn vgetq_lane0(v: int32x4_t) -> i32 {
        core::arch::aarch64::vgetq_lane_s32::<0>(v)
    }
    #[inline]
    #[target_feature(enable = "neon")]
    fn vgetq_lane1(v: int32x4_t) -> i32 {
        core::arch::aarch64::vgetq_lane_s32::<1>(v)
    }
    #[inline]
    #[target_feature(enable = "neon")]
    fn vgetq_lane2(v: int32x4_t) -> i32 {
        core::arch::aarch64::vgetq_lane_s32::<2>(v)
    }
    #[inline]
    #[target_feature(enable = "neon")]
    fn vgetq_lane3(v: int32x4_t) -> i32 {
        core::arch::aarch64::vgetq_lane_s32::<3>(v)
    }

    // A tiny extension so we can write `vdupq_n_s32(0).into_i8()` for a zero
    // int8x16 splat without importing another intrinsic name.
    trait IntoI8 {
        fn into_i8(self) -> int8x16_t;
    }
    impl IntoI8 for int32x4_t {
        #[inline]
        fn into_i8(self) -> int8x16_t {
            // SAFETY: reinterpret-cast of an all-zero 128-bit register; the bit
            // pattern of a zero int32x4 is a zero int8x16. `vreinterpretq_s8_s32`
            // is a no-op bit cast.
            unsafe { core::arch::aarch64::vreinterpretq_s8_s32(self) }
        }
    }

    // ── native packed-int4 nibble unpack + SDOT/SMMLA kernels (bd-1azu.22) ────
    //
    // The packed-int4 path (`int4::igemm_s4s8_packed`) routes here. Unlike
    // `int4::igemm_s4s8` (which materializes the WHOLE B to a dense `[n,k]` i8
    // buffer — the ~5.8x-slower "unpack-then-int8" path in the negative-evidence
    // ledger), these kernels consume the packed nibbles IN-REGISTER: each 8-byte
    // chunk is masked/shifted to its two int8x8 nibble streams ([`unpack8`]) and
    // fed straight to SDOT/SMMLA. No dense B is ever built. Output is f32 (the
    // per-group scale is folded after the integer dot) and OVERWRITES `out`,
    // bit-identical to `int4::s4s8_packed_kernel_scalar`.

    /// Unpack 8 packed int4 bytes (16 nibbles) at `wptr` into `(lo, hi)` int8x8:
    /// `lo` = the even-K weights (each byte's low nibble, the first/even-index
    /// weight), `hi` = the odd-K weights (high nibble), both sign-extended from 4
    /// bits via the `(x << 4) >> 4` 8-bit-lane trick (identical codes to the
    /// scalar `sign_extend_nibble`, gathered 8 at a time).
    ///
    /// # Safety
    /// `wptr` must point at ≥ 8 readable bytes (`vld1_u8` is an unaligned 8-byte
    /// load). Reached only with `neon` enabled.
    #[inline]
    #[target_feature(enable = "neon")]
    unsafe fn unpack8(wptr: *const u8) -> (int8x8_t, int8x8_t) {
        // SAFETY: caller guarantees 8 readable bytes at `wptr`; unaligned u8 load.
        let v = vld1_u8(wptr);
        // low nibble = v & 0x0F, then sign-extend bit 3 (shift left 4, arith right 4).
        let lo_u = vand_u8(v, vdup_n_u8(0x0F));
        let lo_s = vshr_n_s8(vshl_n_s8(vreinterpret_s8_u8(lo_u), 4), 4);
        // high nibble = v >> 4 (logical), then sign-extend.
        let hi_u = vshr_n_u8(v, 4);
        let hi_s = vshr_n_s8(vshl_n_s8(vreinterpret_s8_u8(hi_u), 4), 4);
        (lo_s, hi_s)
    }

    /// Native packed-int4 GEMM via **SDOT**. Register-blocked over a 4×4 output
    /// tile (like `igemm_s8s8_sdot`); the K loop walks 16-element sub-blocks (8
    /// packed bytes), `vzip`-ing the two nibble streams back to natural K order so
    /// a single `vdotq_s32` consumes 16 weights against 16 contiguous activations.
    /// Per group the i32 lane-accumulator is horizontal-reduced (`vaddvq_s32`),
    /// dequantized by that group's f32 scale, and summed in increasing group order
    /// — bit-identical to `int4::s4s8_packed_kernel_scalar`. `out` is OVERWRITTEN.
    ///
    /// `group ∈ {16,32}` ⇒ `group/16 ∈ {1,2}` whole 16-K sub-blocks per group, and
    /// `k` is a multiple of `group` (hence of 16), so the K walk has no tail.
    ///
    /// # Safety
    /// Caller ensures `dotprod` is present and the buffers satisfy the packed
    /// GEMM shape (`a=m*k`, `b_packed=n*k/2`, `scales=n*(k/group)`, `out=m*n`,
    /// `group∈{16,32}` dividing `k`, `k` even).
    #[allow(clippy::too_many_arguments)]
    #[target_feature(enable = "neon,dotprod")]
    pub unsafe fn igemm_s4s8_packed_sdot(
        a: &[i8],
        b_packed: &[u8],
        scales: &[f32],
        group: usize,
        m: usize,
        k: usize,
        n: usize,
        out: &mut [f32],
    ) {
        let groups = k / group;
        let sub_per_group = group / 16; // 16→1, 32→2 sub-blocks per group
        let kbytes = k / 2;

        let mut r = 0;
        while r < m {
            let mr = (m - r).min(4);
            let mut c = 0;
            while c < n {
                let nr = (n - c).min(4);
                // f32 tile accumulator, summed across groups (overwrite semantics).
                let mut out_acc = [[0.0f32; 4]; 4];

                for g in 0..groups {
                    // 4×4 i32x4 accumulator grid for THIS group only.
                    // SAFETY: `vdupq_n_s32` is a pure constant splat.
                    let mut acc = [[vdupq_n_s32(0); 4]; 4];

                    for sub in 0..sub_per_group {
                        let sb = g * sub_per_group + sub; // global 16-K sub-block

                        // Unpack nr weight rows for this sub-block → natural K order.
                        let mut nat = [vdupq_n_s32(0).into_i8(); 4];
                        // indexed loop mirrors SIMD lane layout
                        #[allow(clippy::needless_range_loop)]
                        for j in 0..nr {
                            // SAFETY: (c+j)<n and sb*8+8 ≤ kbytes ⇒
                            //   (c+j)*kbytes + sb*8 + 8 ≤ (c+j+1)*kbytes ≤ b_packed.len().
                            let (lo, hi) =
                                unpack8(b_packed.as_ptr().add((c + j) * kbytes + sb * 8));
                            // zip(lo,hi) → [w(16sb+0), w(16sb+1), …, w(16sb+15)].
                            let zz = vzip_s8(lo, hi);
                            nat[j] = vcombine_s8(zz.0, zz.1);
                        }

                        // Load mr activation rows (16 contiguous int8 = natural K).
                        let mut act = [vdupq_n_s32(0).into_i8(); 4];
                        // indexed loop mirrors SIMD lane layout
                        #[allow(clippy::needless_range_loop)]
                        for i in 0..mr {
                            // SAFETY: (r+i)<m and sb*16+16 ≤ k ⇒
                            //   (r+i)*k + sb*16 + 16 ≤ (r+i+1)*k ≤ a.len().
                            act[i] = vld1q_s8(a.as_ptr().add((r + i) * k + sb * 16));
                        }

                        for i in 0..mr {
                            for j in 0..nr {
                                // SAFETY: vdotq_s32 is register-only; inputs valid.
                                acc[i][j] = vdotq_s32(acc[i][j], act[i], nat[j]);
                            }
                        }
                    }

                    // Dequant this group's i32 sums and fold into the f32 tile.
                    for i in 0..mr {
                        for j in 0..nr {
                            // SAFETY: vaddvq_s32 is a register-only horizontal add.
                            let gi = vaddvq_s32(acc[i][j]);
                            out_acc[i][j] += scales[(c + j) * groups + g] * gi as f32;
                        }
                    }
                }

                for i in 0..mr {
                    for j in 0..nr {
                        out[(r + i) * n + (c + j)] = out_acc[i][j];
                    }
                }
                c += nr;
            }
            r += mr;
        }
    }

    /// Native packed-int4 GEMM via **SMMLA** / i8mm. Same packed-nibble
    /// consumption as the SDOT path, but the int8 MAC is `vmmlaq_s32` over 2×2
    /// output tiles: per 16-K sub-block the even-K nibbles (`lo`) pair with the
    /// even-indexed activations (`vld2_s8` de-interleave) in one SMMLA and the
    /// odd-K nibbles (`hi`) with the odd activations in a second — the SMMLA
    /// matched-lane contraction never needs the nibbles re-interleaved. The 2×2
    /// i32 tile is dequantized by the two B-rows' per-group scales (lane layout
    /// `[c00,c01,c10,c11]`: col `2q`=lanes{0,2}, col `2q+1`=lanes{1,3}) and summed
    /// in increasing group order — bit-identical to `s4s8_packed_kernel_scalar`.
    /// `out` is OVERWRITTEN.
    ///
    /// Doctrine #4: SMMLA is half-rate on Apple Silicon, so on macOS `detect_tier`
    /// prefers SDOT and this is reached only via `FOCR_FORCE_ARCH=smmla` or on
    /// i8mm-preferring cores. The 2×2 tile is intentionally un-deeply-blocked
    /// (correctness-first; the bench reports its throughput honestly). A missing
    /// row/col at an odd `m`/`n` edge is fed a zero panel (a zero nibble/activation
    /// contributes a zero product) and its output cell is skipped.
    ///
    /// # Safety
    /// Caller ensures `i8mm` is present and the buffers satisfy the packed GEMM
    /// shape (as [`igemm_s4s8_packed_sdot`]).
    #[allow(clippy::too_many_arguments)]
    #[target_feature(enable = "neon,i8mm")]
    pub unsafe fn igemm_s4s8_packed_smmla(
        a: &[i8],
        b_packed: &[u8],
        scales: &[f32],
        group: usize,
        m: usize,
        k: usize,
        n: usize,
        out: &mut [f32],
    ) {
        let groups = k / group;
        let sub_per_group = group / 16;
        let kbytes = k / 2;
        // SAFETY: register-only zero splat, padding a missing row/col panel.
        let zero8 = vreinterpret_s8_u8(vdup_n_u8(0));

        let mut rp = 0;
        while rp < m {
            // A-row pair (rp, rp+1); rp is always valid, rp+1 maybe past the edge.
            let mut cp = 0;
            while cp < n {
                // B-col pair (cp, cp+1); cp always valid, cp+1 maybe past the edge.
                // out_acc for [ (rp,cp), (rp,cp+1), (rp+1,cp), (rp+1,cp+1) ].
                let mut lane = [0.0f32; 4];

                for g in 0..groups {
                    // SAFETY: constant splat.
                    let mut acc = vdupq_n_s32(0);
                    for sub in 0..sub_per_group {
                        let sb = g * sub_per_group + sub;

                        // Two weight rows (cp, cp+1) → even (lo) / odd (hi) nibbles.
                        // SAFETY: cp<n and sb*8+8 ≤ kbytes ⇒ in-bounds 8-byte load;
                        // (cp+1)*kbytes+sb*8+8 ≤ b_packed.len() when cp+1<n.
                        let (lo0, hi0) = unpack8(b_packed.as_ptr().add(cp * kbytes + sb * 8));
                        let (lo1, hi1) = if cp + 1 < n {
                            unpack8(b_packed.as_ptr().add((cp + 1) * kbytes + sb * 8))
                        } else {
                            (zero8, zero8)
                        };

                        // Two activation rows (rp, rp+1), de-interleaved even/odd.
                        // SAFETY: rp<m and sb*16+16 ≤ k ⇒ rp*k+sb*16+16 ≤ a.len();
                        // (rp+1)*k+sb*16+16 ≤ a.len() when rp+1<m.
                        let t0 = vld2_s8(a.as_ptr().add(rp * k + sb * 16));
                        let (ev1, od1) = if rp + 1 < m {
                            let t1 = vld2_s8(a.as_ptr().add((rp + 1) * k + sb * 16));
                            (t1.0, t1.1)
                        } else {
                            (zero8, zero8)
                        };

                        // even-K SMMLA: A=[evenAct rp | rp+1], B=[lo cp | cp+1].
                        // SAFETY: vmmlaq_s32 is register-only; i8mm enabled.
                        acc = vmmlaq_s32(acc, vcombine_s8(t0.0, ev1), vcombine_s8(lo0, lo1));
                        // odd-K SMMLA: A=[oddAct rp | rp+1], B=[hi cp | cp+1].
                        acc = vmmlaq_s32(acc, vcombine_s8(t0.1, od1), vcombine_s8(hi0, hi1));
                    }

                    // Dequant by each B-row's per-group scale (cols 2q / 2q+1).
                    let s_c0 = scales[cp * groups + g];
                    let s_c1 = if cp + 1 < n {
                        scales[(cp + 1) * groups + g]
                    } else {
                        0.0
                    };
                    lane[0] += s_c0 * vgetq_lane0(acc) as f32;
                    lane[1] += s_c1 * vgetq_lane1(acc) as f32;
                    lane[2] += s_c0 * vgetq_lane2(acc) as f32;
                    lane[3] += s_c1 * vgetq_lane3(acc) as f32;
                }

                out[rp * n + cp] = lane[0];
                if cp + 1 < n {
                    out[rp * n + cp + 1] = lane[1];
                }
                if rp + 1 < m {
                    out[(rp + 1) * n + cp] = lane[2];
                }
                if rp + 1 < m && cp + 1 < n {
                    out[(rp + 1) * n + cp + 1] = lane[3];
                }
                cp += 2;
            }
            rp += 2;
        }
    }
}

// ─────────────────────────────────────────────────────────────────────────────
// Tests — bit-identical to the scalar reference on randomized + adversarial
// operands at K including the doctrine-#6 worst case 6848.
// ─────────────────────────────────────────────────────────────────────────────
#[cfg(all(test, target_arch = "aarch64"))]
mod tests {
    use super::*;

    /// A self-contained scalar S8S8 oracle (independent of `scalar.rs` so the
    /// test proves correctness even if both share a bug — this is the reference
    /// definition straight from the pinned contract).
    fn oracle_s8s8(a: &[i8], b: &[i8], m: usize, k: usize, n: usize) -> Vec<i32> {
        let mut out = vec![0i32; m * n];
        for r in 0..m {
            for col in 0..n {
                let mut acc: i32 = 0;
                for t in 0..k {
                    acc += i32::from(a[r * k + t]) * i32::from(b[col * k + t]);
                }
                out[r * n + col] = acc;
            }
        }
        out
    }

    fn oracle_u8s8(a: &[u8], b: &[i8], m: usize, k: usize, n: usize) -> Vec<i32> {
        let mut out = vec![0i32; m * n];
        for r in 0..m {
            for col in 0..n {
                let mut acc: i32 = 0;
                for t in 0..k {
                    acc += i32::from(a[r * k + t]) * i32::from(b[col * k + t]);
                }
                out[r * n + col] = acc;
            }
        }
        out
    }

    /// Deterministic xorshift PRNG so the test is reproducible without a dep.
    struct Rng(u64);
    impl Rng {
        fn next(&mut self) -> u64 {
            let mut x = self.0;
            x ^= x << 13;
            x ^= x >> 7;
            x ^= x << 17;
            self.0 = x;
            x
        }
        fn i8(&mut self) -> i8 {
            (self.next() & 0xff) as u8 as i8
        }
        fn u8(&mut self) -> u8 {
            (self.next() & 0xff) as u8
        }
    }

    fn rand_i8(rng: &mut Rng, len: usize) -> Vec<i8> {
        (0..len).map(|_| rng.i8()).collect()
    }
    fn rand_u8(rng: &mut Rng, len: usize) -> Vec<u8> {
        (0..len).map(|_| rng.u8()).collect()
    }

    /// Drive the SDOT kernel directly (bypassing dispatch) when available.
    fn run_sdot_s8s8(a: &[i8], b: &[i8], m: usize, k: usize, n: usize) -> Option<Vec<i32>> {
        if !std::arch::is_aarch64_feature_detected!("dotprod") {
            return None;
        }
        let mut out = vec![0i32; m * n];
        // SAFETY: feature just checked.
        unsafe { aarch64_impl::igemm_s8s8_sdot(a, b, m, k, n, &mut out) };
        Some(out)
    }

    fn run_smmla_s8s8(a: &[i8], b: &[i8], m: usize, k: usize, n: usize) -> Option<Vec<i32>> {
        if !std::arch::is_aarch64_feature_detected!("i8mm") {
            return None;
        }
        let mut out = vec![0i32; m * n];
        // SAFETY: feature just checked.
        unsafe { aarch64_impl::igemm_s8s8_smmla(a, b, m, k, n, &mut out) };
        Some(out)
    }

    /// bd-2mo.3: the OFFLINE-packed-B SMMLA kernel is bit-identical to the
    /// row-major SMMLA kernel and the scalar oracle over randomized +
    /// constant-extreme operands, including odd n, k off the 8-boundary, and
    /// the doctrine-#6 worst-case K=6848 overflow shape. The only difference
    /// between the two kernels is WHERE B's pack happens — offline vs per
    /// call — so equality here proves the zero-shuffle path exact.
    #[test]
    fn smmla_packed_b_matches_row_major_and_oracle() {
        if !std::arch::is_aarch64_feature_detected!("i8mm") {
            eprintln!(r#"{{"check":"smmla_packed_b_parity","result":"skip","reason":"no i8mm"}}"#);
            return;
        }
        let mut rng = Rng(0x0dd0_beef_cafe_f00d);
        let shapes = [
            (1usize, 16usize, 8usize),
            (1, 17, 5),
            (3, 5, 7),
            (4, 64, 8),
            (8, 96, 96),
            (1, 1280, 64),
            (1, 6848, 4),
        ];
        for &(m, k, n) in &shapes {
            let mut a = vec![0i8; m * k];
            let mut b = vec![0i8; n * k];
            if (m, k, n) == (1, 6848, 4) {
                // constant-extreme worst-case-K overflow stress
                a.fill(i8::MAX);
                b.fill(i8::MIN);
            } else {
                for v in a.iter_mut() {
                    *v = rng.i8();
                }
                for v in b.iter_mut() {
                    *v = rng.i8();
                }
            }
            let (panels, _, _) = crate::simd::pack::smmla_pack_panels(&b, 0, n, k, k);

            let mut packed_out = vec![0i32; m * n];
            // SAFETY: i8mm checked at test entry; panel len by construction.
            unsafe {
                aarch64_impl::igemm_s8s8_smmla_packed_b(&a, &panels, m, k, n, &mut packed_out);
            };
            let row_major = run_smmla_s8s8(&a, &b, m, k, n).expect("i8mm present");
            let mut oracle = vec![0i32; m * n];
            scalar::igemm_s8s8(&a, &b, m, k, n, &mut oracle);
            assert_eq!(
                packed_out, row_major,
                "packed-B vs row-major SMMLA [{m},{k},{n}]"
            );
            assert_eq!(
                packed_out, oracle,
                "packed-B vs scalar oracle [{m},{k},{n}]"
            );
            eprintln!(
                r#"{{"check":"smmla_packed_b_parity","m":{m},"k":{k},"n":{n},"result":"pass"}}"#
            );
        }
    }

    /// The SAFE public packed-B wrapper degrades correctly on this host's
    /// natural tier (un-permute + ordinary dispatch when SMMLA is not the
    /// selected tier) and owns the zeroing contract (a dirty `out` must not
    /// leak into the result).
    #[test]
    fn packed_b_public_wrapper_matches_dispatch_and_zeroes_out() {
        let mut rng = Rng(0x5eed_5eed_5eed_5eed);
        for &(m, k, n) in &[(1usize, 48usize, 96usize), (4, 33, 7), (1, 6848, 4)] {
            let mut a = vec![0i8; m * k];
            let mut b = vec![0i8; n * k];
            for v in a.iter_mut() {
                *v = rng.i8();
            }
            for v in b.iter_mut() {
                *v = rng.i8();
            }
            let (panels, _, _) = crate::simd::pack::smmla_pack_panels(&b, 0, n, k, k);
            let mut want = vec![0i32; m * n];
            igemm_s8s8(&a, &b, m, k, n, &mut want);
            let mut got = vec![0x5a5a_5a5ai32; m * n]; // deliberately dirty
            igemm_s8s8_packed_b(&a, &panels, m, k, n, &mut got);
            assert_eq!(got, want, "public packed-B wrapper [{m},{k},{n}]");
        }
    }

    #[test]
    fn sdot_and_smmla_match_oracle_randomized() {
        let mut rng = Rng(0x1234_5678_9abc_def0);
        // A spread of shapes that exercise tile remainders in M, N, and K
        // (non-multiples of 4/8/16).
        let shapes = [
            (1usize, 1usize, 1usize),
            (1, 16, 1),
            (1, 17, 1),
            (3, 5, 7),
            (4, 16, 4),
            (5, 23, 6),
            (8, 32, 8),
            (7, 31, 9),
            (10, 1280, 10),
            (6, 896, 13),
        ];
        for &(m, k, n) in &shapes {
            let a = rand_i8(&mut rng, m * k);
            let b = rand_i8(&mut rng, n * k);
            let want = oracle_s8s8(&a, &b, m, k, n);
            if let Some(got) = run_sdot_s8s8(&a, &b, m, k, n) {
                assert_eq!(got, want, "SDOT mismatch at shape ({m},{k},{n})");
            }
            if let Some(got) = run_smmla_s8s8(&a, &b, m, k, n) {
                assert_eq!(got, want, "SMMLA mismatch at shape ({m},{k},{n})");
            }
            // The public dispatcher must also match.
            let mut got = vec![0i32; m * n];
            igemm_s8s8(&a, &b, m, k, n, &mut got);
            assert_eq!(got, want, "dispatch S8S8 mismatch at shape ({m},{k},{n})");
        }
    }

    /// Batched-decode invariant (bd-1azu.2): an `M=B` GEMM row `r` is
    /// byte-identical to running that row alone as an `m=1` GEMV, on EACH
    /// accelerated tier directly — covering the M-block REMAINDER when `B` is not
    /// a multiple of the register-block height (SDOT MR=4, SMMLA 8-row panels).
    /// This is the property the Phase-6 batched spine rests on; checked here
    /// in-process per tier (the public-dispatch + `FOCR_FORCE_ARCH` re-exec sweep
    /// lives in `tests/batched_igemm_parity.rs`).
    #[test]
    fn batched_m_equals_per_row_gemv_per_tier() {
        let mut rng = Rng(0x0a1b_2c3d_4e5f_6071);
        let m_sweep = [1usize, 2, 3, 4, 5, 7, 8, 9, 16, 17, 33, 64, 65, 128, 129];
        let k = 320usize; // several 16-wide K-blocks; cheap.
        let n = 13usize; // straddles NR=4.
        let b = rand_i8(&mut rng, n * k);
        for &m in &m_sweep {
            let a = rand_i8(&mut rng, m * k);
            let want = oracle_s8s8(&a, &b, m, k, n);
            if let Some(batched) = run_sdot_s8s8(&a, &b, m, k, n) {
                assert_eq!(batched, want, "SDOT M=B vs oracle (m={m})");
                for r in 0..m {
                    let row = &a[r * k..(r + 1) * k];
                    let single = run_sdot_s8s8(row, &b, 1, k, n).expect("dotprod present");
                    assert_eq!(
                        &batched[r * n..(r + 1) * n],
                        &single[..],
                        "SDOT batched row {r} != standalone m=1 GEMV (m={m})"
                    );
                }
            }
            if let Some(batched) = run_smmla_s8s8(&a, &b, m, k, n) {
                assert_eq!(batched, want, "SMMLA M=B vs oracle (m={m})");
                for r in 0..m {
                    let row = &a[r * k..(r + 1) * k];
                    let single = run_smmla_s8s8(row, &b, 1, k, n).expect("i8mm present");
                    assert_eq!(
                        &batched[r * n..(r + 1) * n],
                        &single[..],
                        "SMMLA batched row {r} != standalone m=1 GEMV (m={m})"
                    );
                }
            }
        }
    }

    #[test]
    fn accelerated_paths_add_into_seeded_out() {
        let (m, k, n) = (2usize, 17usize, 3usize);
        let a = vec![
            1i8, -2, 3, -4, 5, -6, 7, -8, 9, -10, 11, -12, 13, -14, 15, -16, 17, -3, 4, -5, 6, -7,
            8, -9, 10, -11, 12, -13, 14, -15, 16, -17, 18, -19,
        ];
        let au = a
            .iter()
            .map(|&x| (x as i16 + 128) as u8)
            .collect::<Vec<_>>();
        let b = vec![
            2i8, 1, -1, 3, -3, 4, -4, 5, -5, 6, -6, 7, -7, 8, -8, 9, -9, -2, 3, -4, 5, -6, 7, -8,
            9, -10, 11, -12, 13, -14, 15, -16, 17, 1, -3, 5, -7, 9, -11, 13, -15, 17, -19, 21, -23,
            25, -27, 29, -31, 33, -35,
        ];
        let seed = (0..m * n)
            .map(|idx| (idx as i32 * 17) - 41)
            .collect::<Vec<_>>();

        let mut want_s8 = seed.clone();
        for (cell, dot) in want_s8.iter_mut().zip(oracle_s8s8(&a, &b, m, k, n)) {
            *cell += dot;
        }
        if std::arch::is_aarch64_feature_detected!("dotprod") {
            let mut got = seed.clone();
            // SAFETY: feature just checked.
            unsafe { aarch64_impl::igemm_s8s8_sdot(&a, &b, m, k, n, &mut got) };
            assert_eq!(got, want_s8, "SDOT must add into seeded out");
        }
        if std::arch::is_aarch64_feature_detected!("i8mm") {
            let mut got = seed.clone();
            // SAFETY: feature just checked.
            unsafe { aarch64_impl::igemm_s8s8_smmla(&a, &b, m, k, n, &mut got) };
            assert_eq!(got, want_s8, "SMMLA must add into seeded out");
        }
        let mut got = seed.clone();
        igemm_s8s8(&a, &b, m, k, n, &mut got);
        assert_eq!(
            got, want_s8,
            "public S8S8 dispatch must add into seeded out"
        );

        let mut want_u8 = seed.clone();
        for (cell, dot) in want_u8.iter_mut().zip(oracle_u8s8(&au, &b, m, k, n)) {
            *cell += dot;
        }
        let mut got = seed;
        igemm_u8s8(&au, &b, m, k, n, &mut got);
        assert_eq!(
            got, want_u8,
            "public U8S8 dispatch must add into seeded out"
        );
    }

    #[test]
    fn s8s8_adversarial_all_max_at_worst_case_k() {
        // Doctrine #6 worst case K plus the adversarial all-127 / all-(-128)
        // operands. These maximize the i32 accumulator (110.4M / 112.2M) and are
        // exactly the values the overflow proof bounds.
        let k = 6848usize;
        let m = 3usize;
        let n = 5usize;
        for &fill in &[127i8, -128i8] {
            let a = vec![fill; m * k];
            let b = vec![fill; n * k];
            let want = oracle_s8s8(&a, &b, m, k, n);
            // Spot-check the headline value.
            let expect = if fill == 127 {
                110_451_392
            } else {
                112_197_632
            };
            assert_eq!(want[0], expect, "oracle value at fill {fill}");
            if let Some(got) = run_sdot_s8s8(&a, &b, m, k, n) {
                assert_eq!(got, want, "SDOT all-{fill} @K=6848");
            }
            if let Some(got) = run_smmla_s8s8(&a, &b, m, k, n) {
                assert_eq!(got, want, "SMMLA all-{fill} @K=6848");
            }
        }
    }

    #[test]
    fn u8s8_matches_oracle_randomized_and_adversarial() {
        let mut rng = Rng(0xdead_beef_0bad_f00d);
        let shapes = [
            (1usize, 1usize, 1usize),
            (2, 15, 3),
            (4, 16, 4),
            (5, 17, 6),
            (8, 64, 8),
            (3, 896, 7),
            (4, 6848, 5),
        ];
        for &(m, k, n) in &shapes {
            let a = rand_u8(&mut rng, m * k);
            let b = rand_i8(&mut rng, n * k);
            let want = oracle_u8s8(&a, &b, m, k, n);
            let mut got = vec![0i32; m * n];
            igemm_u8s8(&a, &b, m, k, n, &mut got);
            assert_eq!(got, want, "U8S8 dispatch mismatch at shape ({m},{k},{n})");
        }

        // Adversarial U8S8: all 255 activations * all 127 weights -> the binding
        // worst case (221.7M) at K=6848.
        let (m, k, n) = (2usize, 6848usize, 3usize);
        let a = vec![255u8; m * k];
        let b = vec![127i8; n * k];
        let want = oracle_u8s8(&a, &b, m, k, n);
        assert_eq!(want[0], 221_772_480, "U8S8 worst-case oracle value");
        let mut got = vec![0i32; m * n];
        igemm_u8s8(&a, &b, m, k, n, &mut got);
        assert_eq!(got, want, "U8S8 adversarial @K=6848");

        // Adversarial negative: all 255 * all -128 (largest |negative|).
        let b_neg = vec![-128i8; n * k];
        let want_neg = oracle_u8s8(&a, &b_neg, m, k, n);
        let mut got_neg = vec![0i32; m * n];
        igemm_u8s8(&a, &b_neg, m, k, n, &mut got_neg);
        assert_eq!(got_neg, want_neg, "U8S8 negative adversarial @K=6848");
    }

    #[test]
    fn empty_and_unit_dims() {
        // K = 1 (below the 16/8 vector width) must work via the tail path.
        let a = vec![3i8, -2, 5];
        let b = vec![7i8, 11];
        let (m, k, n) = (3usize, 1usize, 2usize);
        let want = oracle_s8s8(&a, &b, m, k, n);
        let mut got = vec![0i32; m * n];
        igemm_s8s8(&a, &b, m, k, n, &mut got);
        assert_eq!(got, want);
    }

    // ── native packed-int4 SDOT/SMMLA kernels (bd-1azu.22) ───────────────────
    //
    // In-process per-tier coverage: drive the PRIVATE packed kernels directly
    // (feature-gated, bypassing dispatch) and assert bit-identical f32 to a
    // packed-int4 scalar oracle written straight from the pinned packing spec —
    // the same proof pattern `run_sdot_s8s8` / `run_smmla_s8s8` use for int8.
    // (The public-dispatch + `FOCR_FORCE_ARCH` sweep + cross-impl parity vs the
    // `unpack-then-int8` `igemm_s4s8` lives in `tests/int4_packed_parity.rs`.)

    /// Packed-int4 GEMM scalar oracle, written from the spec (low nibble = even-K
    /// weight, high = odd-K; per-group i32 dot, dequant-and-sum in f32 in
    /// increasing group order). Independent of the module's own reference.
    fn oracle_s4s8_packed(
        a: &[i8],
        b_packed: &[u8],
        scales: &[f32],
        group: usize,
        m: usize,
        k: usize,
        n: usize,
    ) -> Vec<f32> {
        let groups = k / group;
        let mut out = vec![0.0f32; m * n];
        for mi in 0..m {
            for ni in 0..n {
                let mut acc_f = 0.0f32;
                for g in 0..groups {
                    let mut gi: i32 = 0;
                    for kk in g * group..(g + 1) * group {
                        let byte = b_packed[ni * (k / 2) + kk / 2];
                        // even kk → low nibble, odd kk → high nibble; sign-extend.
                        let w = if kk % 2 == 0 {
                            ((byte << 4) as i8) >> 4
                        } else {
                            (byte as i8) >> 4
                        };
                        gi += i32::from(a[mi * k + kk]) * i32::from(w);
                    }
                    acc_f += scales[ni * groups + g] * gi as f32;
                }
                out[mi * n + ni] = acc_f;
            }
        }
        out
    }

    fn run_sdot_s4s8(
        a: &[i8],
        b_packed: &[u8],
        scales: &[f32],
        group: usize,
        m: usize,
        k: usize,
        n: usize,
    ) -> Option<Vec<f32>> {
        if !std::arch::is_aarch64_feature_detected!("dotprod") {
            return None;
        }
        let mut out = vec![0.0f32; m * n];
        // SAFETY: dotprod just runtime-detected; buffers match the packed shape.
        unsafe {
            aarch64_impl::igemm_s4s8_packed_sdot(a, b_packed, scales, group, m, k, n, &mut out);
        }
        Some(out)
    }

    fn run_smmla_s4s8(
        a: &[i8],
        b_packed: &[u8],
        scales: &[f32],
        group: usize,
        m: usize,
        k: usize,
        n: usize,
    ) -> Option<Vec<f32>> {
        if !std::arch::is_aarch64_feature_detected!("i8mm") {
            return None;
        }
        let mut out = vec![0.0f32; m * n];
        // SAFETY: i8mm just runtime-detected; buffers match the packed shape.
        unsafe {
            aarch64_impl::igemm_s4s8_packed_smmla(a, b_packed, scales, group, m, k, n, &mut out);
        }
        Some(out)
    }

    #[test]
    fn s4s8_packed_tiers_match_scalar_oracle_randomized() {
        let mut rng = Rng(0x5104_9ac4_ed00_0a1b); // deterministic seed
        // Shapes straddle the SDOT 4×4 and SMMLA 2×2 tile remainders in M and N,
        // exercise both group sizes (16/32), single & multi sub-block groups, and
        // the doctrine-#6 worst case K=6848 (dense layer-0 down_proj).
        let cases = [
            (1usize, 16usize, 1usize, 16usize),
            (1, 32, 1, 32),
            (2, 32, 3, 16),
            (4, 64, 5, 32),
            (3, 96, 2, 16),
            (5, 128, 7, 32),
            (8, 48, 9, 16),
            (7, 160, 6, 32),
            (9, 16, 4, 16),
            (16, 64, 13, 16),
            (6, 6848, 5, 16),
            (4, 6848, 5, 32),
        ];
        for (m, k, n, group) in cases {
            let groups = k / group;
            let a = rand_i8(&mut rng, m * k);
            let b_packed = rand_u8(&mut rng, n * (k / 2));
            // Mixed-sign scales incl. non-power-of-two; the f32 summation order is
            // identical on both sides, so any scale gives bit-identical results.
            let scales: Vec<f32> = (0..n * groups)
                .map(|i| ((i as f32 * 0.013) - 0.4) + if i % 3 == 0 { 0.125 } else { -0.0625 })
                .collect();
            let want = oracle_s4s8_packed(&a, &b_packed, &scales, group, m, k, n);
            if let Some(got) = run_sdot_s4s8(&a, &b_packed, &scales, group, m, k, n) {
                assert_eq!(
                    got, want,
                    "SDOT-packed != oracle (m={m},k={k},n={n},g={group})"
                );
            }
            if let Some(got) = run_smmla_s4s8(&a, &b_packed, &scales, group, m, k, n) {
                assert_eq!(
                    got, want,
                    "SMMLA-packed != oracle (m={m},k={k},n={n},g={group})"
                );
            }
        }
    }

    /// Adversarial: all weights = int4-min (-8, byte 0x88), activations ±127 — the
    /// operands that maximize the i32 per-group accumulator. Each group of ≤32
    /// terms is ≤ 32·8·127 = 32_512 ≪ i32::MAX (doctrine #6); the SIMD i32 lane
    /// reduce must equal the scalar oracle bit-for-bit (no overflow, exact dot).
    #[test]
    fn s4s8_packed_adversarial_max_operands() {
        for &(m, k, n, group) in &[
            (3usize, 6848usize, 5usize, 16usize),
            (4, 6848, 6, 32),
            (5, 64, 7, 16),
        ] {
            let groups = k / group;
            let b_packed = vec![0x88u8; n * (k / 2)]; // both nibbles = -8
            let a: Vec<i8> = (0..m * k)
                .map(|i| if i % 2 == 0 { 127 } else { -127 })
                .collect();
            let scales = vec![1.0f32; n * groups];
            let want = oracle_s4s8_packed(&a, &b_packed, &scales, group, m, k, n);
            if let Some(got) = run_sdot_s4s8(&a, &b_packed, &scales, group, m, k, n) {
                assert_eq!(got, want, "SDOT-packed adversarial (k={k},g={group})");
            }
            if let Some(got) = run_smmla_s4s8(&a, &b_packed, &scales, group, m, k, n) {
                assert_eq!(got, want, "SMMLA-packed adversarial (k={k},g={group})");
            }
        }
    }
}