Skip to main content

Module serial

Module serial 

Source
Expand description

Serial-over-chunks SSD with plain autodiff backward.

§SingleSsd Serial (K1–K5) SSD

Chunk-serial counterpart to crate::mamba3::single_ssd::ssd::minimal. Whereas the Minimal variant uses a segsum-based quadratic state passing, this one reuses the K1–K4 helpers from crate::mamba3::double_ssd::ssd::serial (which run a sequential loop for K4) and supplies a new K5 that bakes in the single-ssd logic:

  • Strict lower-triangular intra-chunk path (the same-time-step block is excluded from the SSM sum; it is the “diagonal correction” territory).
  • K is scaled by scaleₜ = γₜ + (1−λₜ₊₁) Δₜ₊₁ per source-time column.
  • Same-time-step block contributes via an explicit γₜ · (C·Bᵀ at t) · Vₜ correction term, restoring the right diagonal weighting.

K1–K4 are identical to the double-SSD because:

  • K1 (da_cumsum, da_chunk_end) depends only on da = Δ·A.
  • K2 (cb = C · Bᵀ) is computed on unscaled B / C; the single-ssd algorithm wants the unscaled CB so it can apply scaleₜ per-column (lower triangular) and reuse the same-step block for the γ-correction.
  • K3 (chunk-end state from V·decay·K) is form-invariant: passing the scale-multiplied K (K_scaled = scaleₜ · B) recovers the single-ssd chunk state, with no other changes needed.
  • K4 (sequential state passing across chunks) operates on a [H, P, R] per-chunk state and a per-chunk decay total; both are mode-agnostic.

Reference kernels (same as single_ssd_minimal):

  • refs/state-spaces/mamba/mamba_ssm/ops/triton/mamba3/mamba3_siso_fwd.py
  • refs/state-spaces/mamba/mamba_ssm/ops/tilelang/mamba3/mamba3_mimo_fwd.py

Re-exports§

pub use crate::mamba3::double_ssd::ssd::serial::k1_ssd_chunk_cumsum;
pub use crate::mamba3::double_ssd::ssd::serial::k2_ssd_bmm;
pub use crate::mamba3::double_ssd::ssd::serial::k3_ssd_chunk_state;
pub use crate::mamba3::double_ssd::ssd::serial::k4_ssd_state_passing;

Functions§

k5_single_ssd_chunk_scan
SingleSsd chunk scan.