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 onda = Δ·A. - K2 (
cb = C · Bᵀ) is computed on unscaled B / C; the single-ssd algorithm wants the unscaled CB so it can applyscaleₜ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.pyrefs/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.