Crate memra_engine 

memra engine: Stage-1 correctness-first forward-pass kernels + ops, on sm_120 via cudarc.

Re-exports

pub use memra_gguf;

pub use memra_runtime;

pub use memra_sampling as sampler;

Modules

cache

The dual cache lives in the shared memra-kv crate (Phase D extraction); this re-export keeps every crate::cache:: / memra_engine::cache:: path unchanged.

decode

Incremental decode (T=1) with the dual cache + greedy generation loop. Serves end-to-end. Reuses the validated kernels; threads KV (full-attn) and conv/SSM state (linear-attn) across steps.

decode_batch

Batched decode step — B sequences share one fused pass (ARCHITECTURE-H100.md §3 B2’).

dflash

DFlash block-diffusion drafter (DFLASH-BRINGUP-PLAN.md, 2026-07-13).

eagle

EAGLE3.1 greedy-chain speculative decode (research/basics/EAGLE-PLAN.md, N1-N7).

f16_ffi

FP16-MIRROR PREFILL (MEMRA_PP_F16=1): cuBLASLt FP16 TN GEMM on a resident fp16 dequant mirror of the Q8_0 trunk weights.

forward

Dense forward pass (Stage-1, all f32, prefill of T tokens, batch=1). Matches llama.cpp qwen3 graph: embed → per layer {RMSNorm, QKV, QK-norm, RoPE, SDPA, O, residual, RMSNorm, SwiGLU, residual} → output_norm → lm_head.

fp8_ffi

FP8-ACT PREFILL (MEMRA_PP_FP8=1): cuBLASLt FP8-E4M3 TN GEMM for the F8-E4M3-origin projections.

gemma_spec

gemma4 MTP spec-decode: the “gemma4-assistant” drafter (4-layer, Q-only attention over the MAIN model’s KV cache — no draft KV, no trims) + the greedy draft/verify loop.

graph_update

CUDA-graph exec-update (shared, model-agnostic): capture a decode step ONCE, then re-tune individual kernel nodes’ launch geometry per token via cuGraphExecKernelNodeSetParams — the llama.cpp graph-serving mechanism (their decode replays one instantiated graph per token with exact per-token grid shapes; nsys shows zero launch gaps AND eager-exact grids, where a fixed-bucket replay wastes split blocks).

hybrid

Qwen3.5/3.6 hybrid model: linear-attention (Gated DeltaNet) layers + periodic full-attention layers + SwiGLU FFN. Loads weights, runs the forward, dual cache. Builds on the validated conv1d + gdn_scan kernels (M2/M3) and the dense full-attn path (M0).

hybrid_forward

Hybrid forward pass (Stage-1, f32, prefill, single sequence). Per layer dispatches to a linear-attention (Gated DeltaNet) or full-attention mixer, then SwiGLU FFN. Matches llama.cpp src/models/qwen35.cpp node-for-node.

mla

MLA (multi-head latent attention) CPU f32 reference — GLM-5.2 bring-up lane increment 1. Naive vs absorbed decode forms + NORM/NEOX rope permutation, unit-tested; the permanent oracle for the MLA kernel family (research/mla-bringup-20260801/DESIGN.md). No CUDA deps. MLA (multi-head latent attention, DeepSeek lineage / GLM-5 “MLA-256”) — CPU f32 reference.

mmq_ffi

FFI to the MMQ prefill GEMMs (cu/mmq_fp4.cu + cu/mmq_q45k.cu) — vendored floor kernels.

model

Dense transformer model: loads GGUF weights to GPU (Stage-1: dequant→f32), runs the shared full-attention + SwiGLU forward graph. Arch-agnostic via ModelConfig; this path is exactly the dense-transformer graph (qwen3) and the full-attention layers of hybrids.

moe_cache

EDGE-1 §B: SLRU GPU expert-residency cache (MOE-SLRU-PLAN §B).

moesd

Diagnostic-only expert-union capture for the MoESD target-efficiency harness.

pp

M2 pipeline-parallel N-stage runtime (generalizes the M1 2-stage seam).

prime_graph

PrimeGraph (task #14, design v3): a per-bucket CUDA graph of the FULL fresh-prime trunk, bound to a dedicated SCRATCH cache; serving replays it (one cuGraphLaunch, ~23ms vs ~26ms eager at bucket 512) and COPIES the outputs into the session’s cache (KV rows + conv rings + recurrent states, ~tens of us D2D — the copy-out beats both table-indirect kernels and graphExec node patching, ledger design v3).

round_stream

ROUND-STREAM: the model-generic device machinery for pre-issued M-round speculative bursts with zero per-round host readbacks (extracted from the qwen spec loop 2026-07-12 so the gemma/next-model loops reuse it instead of re-growing their own).

sigrouter_contract

Cross-surface contracts for DeepSeek-class sigmoid routing.

spec

Qwen3.5 MTP (NextN) greedy speculative decode (research/mtp/MTP-PLAN.md §A/§B/§C/§D).

spill

SPILLING-PLAN: full tiered spilling (VRAM ↔ pinned-host ↔ mmap-disk).

Structs

AttnPreVl

task #18 (attn pre-FA): per-seq split/norm/rope/append args (CUDA attnpre_t).

AttnPreVl8

Engine

Engine device context: CUDA context, stream, loaded kernel modules, cuBLASLt (via runtime::Gpu).

F32x8

STAGE-2 GROUPED DECODE: the 8 routed-expert weights by value (CUDA f32x8_t).

FaSeqVl

task #18 (attn side): per-seq varlen FA args (CUDA faseq_t/favl_t).

FaVl8

GdnChunkBufs

task #18 increment 2: one sequence’s FULL chunk-buffer set (alloc-only; the varlen K1-K5 chain fills them).

GdnPrepVl

task #18 increment 3: per-seq PREP/TAIL args (CUDA gdnprep_t/gdnprepvl_t).

GdnPrepVl8

GdnSeqVl

task #18 varlen GDN: per-seq args for gdn_chunk_{state,output}_mma_vl — one launch runs all B<=8 sequences’ K4/K5 (CUDA gdnseq_t/gdnvl_t, layout-identical repr(C)). Raw addresses are valid for the launch: every referenced buffer outlives the call and all work is on the single compute stream (same discipline as the f16 GEMM FFI).

GdnVl8

GdnWVl

task #22: per-seq wgmma-fused extras (CUDA gdnw_t/gdnwvl_t) — qb16 mirror + pre-masked Pb16, riding NEXT TO GdnSeqVl so the base struct stays untouched.

GdnWVl8

WPtr8

STAGE-2 GROUPED DECODE: 8 expert weight-block device pointers passed BY VALUE as one kernel param (matches the CUDA wptr8_t struct: 8x 64-bit pointers, #[repr(C)] => identical layout). The pointers are SLRU cache-slot base addresses — fixed for the engine’s lifetime (slots are never re-allocated), so passing raw values is stable across the launch.

Constants

ARGMAX_NB

Number of pass-1 blocks for the parallel argmax (fan-out across SMs to saturate HBM). 256 blocks x 256 threads = 65536 threads covering the 248K-vocab scan in ~4 strided loads/thread.

FA_DEEP_MIN_DEFAULT

FA-DEEP gate (2026-08-02, lane fa-decode-deep): deep-ctx v4 twins (fa_decode_vec_q_v4_deep / _deep_dc) — the depth-decode lane’s priced fix. Unlike v2/v3/v4 this is NOT a numeric config: the deep twins run the v4 program VERBATIM (same split partition, same softmax/accumulation order, same partials/combine) and only move the smem physical layout (bank de-conflict row pads) + the load schedule (next-tile L2 prefetch) — kernel-check pins bitdiff==0 vs the v4 twins across depths, so eager / rows-verify / graph / seqs stay mutually bit-identical wherever the threshold falls. Engages at t_kv >= MEMRA_FA_DEEP_MIN. The swept floor is 0 = ALWAYS ON where v4 ran (fa-deep-bench fine grid 96..6144, 2026-08-02: deep flat-or-better at EVERY depth, 1.01-1.26x, no losing cell — so there is no engagement boundary and no new capture-recapture edge; the env stays as a sweep/diagnostic seam only). MEMRA_FA_DEEP=0 is the rollback seam. Read per call so the battery + bench can A/B within one process (the v2/v3 pattern).

FA_VEC_MIN_TKV

TUNE SEAM: keys per FA-decode split (MEMRA_FA_SPLIT forces a fixed size; default 64). Smaller splits raise grid.y so grid = n_head_kv * n_splits fills the 82 SMs at short/mid ctx (vec path launches only n_head_kv=8 CTAs per split). Swept clock-locked 2026-07-03 (graph tg128): 32 beat 64 at ctx 128/512 (+0.5/+1.2%) and lost at 2048 (-3%) — BUT the adaptive 32/64 default BROKE the MTP spec-decode exact-match gate (run-spec K=1/2 self-consistency FAIL with 32; PASS with 64): the split count changes the combine’s FP summation order, and the spec verify’s batched forward only argmax-matches single-step decode under the 64-split order on real prompts. Spec exactness (the bigger lever) outranks a <=1.2% decode win -> default stays FIXED 64; sweeps use the env. Takes t_kv so eager, _dc capture, and fa_geom_eager stay signature-compatible for future adaptive retries (any retry MUST pass run-spec self-consistency first). Minimum t_kv for the warp-per-token vec FA path (below it the scalar path’s 4x-more-blocks hides latency better — measured crossover, see fa_decode). Shared by fa_decode / fa_decode_dc / fa_geom_eager / fa_decode_rows-eligibility (spec verify) so the kernel pick NEVER diverges between eager decode and the verify (the spec-exactness law).

QT_BF16

QT_F8_E4M3

Checkpoint-native FP8-E4M3 (MEMRA_ST_E4M3, lane e4m3dec): raw safetensors e4m3 weight bytes [out_f, in_f] row-major (row_bytes == in_f), per-tensor f32 weight_scale in GpuTensor scale (fused at the mmvq write / post-matmul scale_inplace). Decode = qmatvec_e4m3_mmvq (+ _b2/_b4/_b8 batched twins); prefill (m>=16) = the cuBLASLt FP8 GEMM on the SAME resident bytes (fp8_ffi.rs) — ONE weight copy total, no Q8_0 re-encode duplicate.

QT_F8_E4M3_BLK

Checkpoint-native FP8-E4M3 with a BLOCK-128 weight-scale GRID (lane/fp8-blk128-decode, 2026-08-05) — the Qwen-official FP8 / DeepSeek-V3 scale class. Same raw e4m3 bytes as QT_F8_E4M3 ([out_f, in_f] row-major, row_bytes == in_f), but the dequant scale is GpuTensor::Quant.blk (Fp8BlockScales, [ceil(out_f/128), ceil(in_f/128)] f32) and the scalar scale field is 1.0 by the layout contract.

QT_F32

Unquantized f32 weight (safetensors MoE Path A: experts dequantized to f32 host-resident).

QT_IQ3_S

QT_IQ4_XS

QT_NVFP4

QT_NVFP4_RP

Device-side tag for the A6 SPLIT-PLANE repacked NVFP4 layout (Stage-A generic kernel only; GpuTensor keeps qtype=QT_NVFP4 + an rp flag — this tag never lives in a GpuTensor).

QT_Q2_K

GGUF Q2_K. Appended after the existing Q4_0 code so kernel ABI values do not move. Mixed-expert artifacts use the generic f32-dequant staged kernel until a target-rig-gated dp4a/MMQ implementation exists.

QT_Q3_K

QT_Q4_0

QT_Q4_K

QT_Q5_K

QT_Q6_K

QT_Q8_0

Quant type codes matching qmatvec.cu QType enum.

ROUTER_BATCH_MIN_T

FAST-ROUTER batch twin (lane/fast-router, 2026-08-02). The concat-prime exactness fix (router_prefill_exact_on) routes prefill through router_gemv — m-invariant, but a per-(expert,token) GEMV program with zero operand reuse, so q35 board-2048 prefill paid -10% on the 5090. router_gemv_f32_w8_batch register-tiles (8x8 expert-x-token) the same per-row FP chains (BIT-IDENTICAL per row — kernel-check sweeps m=1..2048 on real router weights), so the t crossover below is pure perf, not a numeric config. Swept on-box (research/fast-router-20260802/crossover-router*.jsonl): plain wins t<=4, batch +7-9% at t=8, 1.9x at t=16 rising to 3.45x at t=2048 — MIN_T=8. Decode t=1 and spec verify t<8 keep the plain w8 form. MEMRA_ROUTER_BATCH=0 forces plain at every t (rollback seam, perf-only: bits are equal by the kernel-check gate). Killed arms (same sweep, JSONL is the record): the 8x16 tile lost to 8x8 at every t (128-accumulator register pressure beats the halved w-traffic), and the same-shape sigmoid_dot_rows twin (out_f=1) measured 0.62-0.89x at every prefill t (launch-latency-bound, ~7us/layer at m=2048) — both bit-identity-PASSED before dying.

Statics

FA_SMEM_TKV_DEFAULT

t_kv-conditional v4 pick (gemma depth lesson 2026-07-10: v4’s key-per-lane pipeline starves at the 1024-window with short splits — MEMRA_FA_V4=0 measured depth plain 158.0 vs 156.7). Threshold MEMRA_FA_V4_MAX (default usize::MAX = unchanged behavior; gemma sets 1024 at load via FA_V4_MAX_DEFAULT). Applied at EVERY dispatch site (eager, rows, rows_w, dc) so verify stays kernel-family-identical to decode at the same t_kv. Per-model deep-ctx smem floor default (MEMRA_FA_SMEM_TKV env overrides): gemma pushes it above the 1024 window so the windowed decode + verify rows share the REGISTER family.

FA_SP512_DEFAULT

FA_SPW_DEFAULT

Per-model windowed-split default (MEMRA_FA_SPW overrides): gemma MoE (26B, nkv=8) measured 32 (grid-limited t=1 under the raw-e4m3 sV ceiling, 2026-07-12); dense gemma (31B) measured 64 (37.13/37.12 vs 36.87/36.86 at 1.7k, N=2 — different attention geometry).

FA_SP_GEMMA

gemma4 fa split ladder switch (set at model load; see fa_split_keys).

FA_V4_MAX_DEFAULT

FA_VEC_MIN_DEFAULT

Per-model crossover default, set at model load BEFORE the first decode (per-model numeric-config adoption law). qwen keeps the measured 96; gemma4 (nkv=8 SWA) measured vec-always fastest: 119.9 (96) / 130.0 (48) / 133.2 (1) tok/s tg128-regime, 2026-07-10.

FUSED_MR1_DEFAULT

Per-model hd512 (gemma globals) split default (MEMRA_FA_SP512 overrides): 26B measured 16 (2026-07-11 N=2), dense 31B measured 32 (36.86/36.93 vs 36.73/36.73 at 1.7k, 2026-07-12). fused t=1 q4_0 pair/triple row mapping: true = mr1 (one row/warp). Per-model default (dense gemma wins +1.1% short / +0.6% depth on the 31B; MoE 26B REGRESSES −1.2% — its shared-expert fused2 shapes lose to the finer grid). MEMRA_Q40_MR env still wins.

KV_FP8_FORCE

Per-model FP8-KV door — lives in memra-kv next to the format policy it drives (re-export keeps crate::KV_FP8_FORCE setters in model.rs/hybrid.rs working). Per-model FP8-KV door (-1 = unset → env/default off; 0 = off; 1 = on). Set at qwen model load: the 2026-07-12 arc closed per-model — 9B +0.7-4% scaling with depth, 27B flat (weight-bound), 35B −2% (fp8 format-gates its v3 dp4a lane off). Explicit MEMRA_KV_FP8 wins. 9B adoption attempt REVERTED by measurement 2026-07-29 (−1% at 12k on the then-current build) — loaders currently store 0.

MMQ_SK_FORCE

Per-model stream-k override for SPEC serving (-1 = unset → env/default; 0 = force tiling; 1 = force sk). Set by generate_spec_gemma per model tier — the sk autotune’s per-process kernel coin made 12B-class spec cells bimodal, while the 26B’s drafter measures BETTER under sk’s fold order (2026-07-27). mmq_ffi reads this before the env.

PRIME_NANOS

Harness timing contract: wall nanos of the LAST generate/generate_spec prompt prime on this process. Bench binaries read it right after the call to print gen-only throughput without the prime-subtraction hack (which amplifies prime jitter into the gen number at long prompts).

RMS_BLOCK_DEFAULT

Per-model rms_norm block size (per-model numeric-config law: the per-thread partial-sum split changes with blockDim -> different FP order -> battery-arbitrated per model). qwen keeps the shipped 256; gemma4 adopts 1024 (single-row 2816-col norms are one-block latency-bound at 256 threads — 7us/launch measured).

ROUTER_W8_DEFAULT

Per-model router-GEMV form (2026-07-31): the 8-warp twin is +8.8% on the H100 q35 decode step (router was 14.8% of it) with argmax + spec self-consistency green on qwen-class MoE both rigs. The gemma-4 26B knife-edge block (2026-07-31, single synthetic prompt) was RE-ARBITRATED 2026-08-01 on 6 real prompts — gate outcomes identical to the lone-warp arm, +13% g26 decode — so gemma4 rides the default too (research/g26-decode-20260801/). MEMRA_ROUTER_V2 env overrides either way.

Functions

fa512_hp_on

hd512 head-pair arm (DEFAULT since stamp v4; MEMRA_FA512_HP=0 reverts to sp16): GQA ncols2=2 — 2 heads per CTA share each staged K/V tile, Q register-resident. Engages when n_head is even and the GQA group (n_head/n_head_kv) is even.

fa512_min_tkv

hd-512 vec crossover floor (MEMRA_FA512_MIN, default 512) — shared by fa_decode dispatch and the gemma global-layer rows/parity call sites.

fa512_wide_warps

4-warp sp16 experiment arm (MEMRA_FA512_W4=1, requires the f16pv door): GEMM0 split-K 4-way + GEMM1 4x128 O-dims. Own partial-sum order — oracle-band gated. Returns warp count (2 = base sp16). 8-warp arm measured NEGATIVE 2026-07-23 (jsonl) and removed.

fa_deep_at_pub

Public twin (kernel-check builds the deep-vs-v4 bit pin; bench sweeps the floor).

fa_f16pv_on

f16-P/V class (DEFAULT since 2026-07-23 stamp v4; MEMRA_FA_F16PV=0 = f32-class rollback): llama-fa=1-style f16 P + f16 P@V accumulation on the hd512/SWA prefill stamps (KQ/softmax/normalize stay f32). Laptop stamp: 12B 1.045x, 31B 0.979x vs llama.

fa_seqs_eligible

BATCHED-TICK increment 2 (2026-08-01): true iff a row at this t_kv would take the v4 eager arm in fa_decode_kvmod‘s dispatch — the exact precondition for the z-batched fa_decode_vec_q_seqs_v4 twin to reproduce its per-seq program bit-identically. Mirrors the kvmod predicates: vec on + above the vec floor + hd256 + inside the v4 window + the PRODUCTION v4 body (the noB3/stage phase probes are wrong-output) + the default flash module (no fp8-KV g-module). Callers must ALSO group rows on one fa_split_keys rung (the rows-twins’ straddle law) before batching.

fa_split_keys_pub

Public twin of the crate-private split ladder (kernel-check builds the seqs-vs-loop pin).

fa_v4_at_pub

fa_vec_min_tkv

Env-overridable crossover (MEMRA_FA_VEC_MIN, default FA_VEC_MIN_TKV). The 96 floor was measured on the qwen geometry (nkv=2); gemma4 SWA layers run nkv=8 = 4x the vec grid, which moves the crossover — sweep per model, adopt per the battery.

faw_hp_on

SWA head-pair arm (DEFAULT since stamp v4; MEMRA_FAW_HP=0 reverts to p1): llama-class windowed geometry — 32 q-rows x 2 heads per CTA sharing staged K/V, f16 P@V accumulation. Even n_head and even GQA group required (guarded per call).

kv_blk_bytes

KV format policy moved to the shared memra-kv crate (Phase D); re-exported so the fatbin router below and every existing crate::kv_blk_bytes() call site is unchanged. Per-32-element block bytes for the selected (K, V) formats.

kv_cache_formats

KV format policy moved to the shared memra-kv crate (Phase D); re-exported so the fatbin router below and every existing crate::kv_blk_bytes() call site is unchanged. Env-selected KV cache formats (MEMRA_KV_K / MEMRA_KV_V). The engine’s flash-fatbin router and the cache sizing below MUST agree — both read this one function.

moe_f16g_direct_on

DIRECT-FROM-QUANT sk tile loaders (lane/kquant-tile-loaders, 2026-08-02; IQ classes added by lane/iq-direct-loaders): Q4_K/Q6_K/IQ4_XS/IQ3_S expert projections on the mode-2/3 sk visitor forms dequant their weight tiles in-register from the quant superblocks instead of running the per-(layer,projection) dequant pass into an f16 workspace (41.8% of Ornith-35B t=512 kernel time — the pp512 wall, research/q4k-expert-prefill-20260802 §5; the IQ classes are 94.8% of q35’s bank bytes — the h100-sk-direct coverage pricing). Bit-identical to the workspace path by construction (kernel-check “f16g-kq-direct” gates it bitwise) — a data-movement change, not a numeric-class change. Default ON; MEMRA_F16G_DIRECT=0 reverts to the workspace path everywhere; MEMRA_F16G_DIRECT=kq keeps the k-quant loaders and reverts only the IQ classes (the iq-direct-loaders A/B seam — the pre-lane shipped config).

moe_f16g_gemma_on

Per-model door for the gemma-MoE (gelu) grouped path: round 49’s Hopper default REGRESSED g26 board-2048 prefill -8.3% interleaved x5 on-box (def median 10380, wild 8.9k-11.7k spread; off 11317, ±0.13%) — the +6-15% probe verdict didn’t survive the board workload (stale-verdict law, round 50). The silu/qwen class keeps the round-49 default (q35 +53% board-2048). Explicit MEMRA_MOE_F16G=1/2 still opens this door for A/B.

moe_f16g_mode

In-house MoE router GEMV on the spec-verify small-t path (DEFAULT ON since 2026-07-10: battery green on 35B p2/p3 K=1..8, acceptance bit-identical, +2-4% spec e2e — replaces ~240 per-column cuBLAS gemv launches/round). MEMRA_ROUTER_KERNEL=0 is the rollback seam. MoE grouped f16 GEMM door (experimental until gated), f16-mirror numeric class: per-layer expert dequant to f16 + one grouped f16 GEMM over the CSR groups. MEMRA_MOE_F16G=1 cublasGemmGroupedBatchedEx (round 46 arc 2). The grouped API issues through cublas-internal streams NOT ordered with ours — v1 pays a full stream sync per projection (round-47 ledgered defect). MEMRA_MOE_F16G=2 single-kernel grouped GEMM on the engine stream (round 49): ordered by construction, zero syncs, f32 C with the act row-scale folded in. DEFAULT (2026-08-01, round 49 promotion): mode 1 on the Hopper lane — with the 41/41 dequant coverage fix the q35 board-2048 prime measured 5490 (MMQ) / 8380 (mode 1, +53%) / 7990 (mode 2) x3 interleaved on the H100, argmax MATCH — the last board loss flips. The 5090 measured FLAT (858GB/s makes the dequant-workspace traffic cancel the GEMM win) — but that verdict is for expert banks the int8-MMA MMQ arm can take (IQ3_S/IQ4_XS/Q4_0). MEMRA_MOE_F16G=0 kills anywhere.

moe_f16g_sk_params

Mode-2 sk kernel form policy (round 51, lane/sk-bm128): the single-kernel grouped GEMM runs as a persistent problem-visitor over the real CSR tiles with two tile forms. Returns (shape_sel, cross) for the FFI: MEMRA_F16G_SK=0 -> (-1, _): the round-49 grid-scan kernel (rollback seam). MEMRA_F16G_SK=32 -> all groups on the 32x64x32 2-stage form (cross = i32::MAX). MEMRA_F16G_SK=128 -> all groups on the 128x64x64 3-stage form (cross = 1; groups fall back to 32x64 in-launcher when the device/in_f can’t take it). unset -> hybrid split: groups with m_e >= MEMRA_F16G_SK_CROSS ride the 128 form. Default cross = 64 (5090 sweep 2026-08-01, receipts research/sk-bm128-20260801/; H100 re-swept on the direct+tail form 2026-08-02, lane/h100-flip-full: {16,32,64} -> 12868/13192/13225 — 64 wins there too, the pre-direct 32 verdict was stale).

moe_f16g_tail_on

DEEP-TAIL sk form (lane/sk-tail-form, 2026-08-02): groups below the visitor crossover ride a 32x64x64 3-STAGE cp.async tile instead of the round-51 32x64x32 2-stage — the same 32-row tile (zero extra padding), 2 k-blocks in flight instead of 1 and half the syncs per k. The H100 ncu pricing (research/sk-bm128-20260801) put the 2-stage tail at 31% of the sk GEMM stage under q35’s routing skew. Bit-identical to every other sk form by construction (kernel-check “f16g-sk” gates all tail arms maxdiff==0); exists in both the workspace-f16 and direct-from-quant variants. Default ON; MEMRA_F16G_TAIL=0 = rollback to the 2-stage tail. in_f % 64 != 0 falls back in-launcher.

moe_fuse_actq_on

Fused act-epilogue (silu/gelu-mul + q8_1_mmq quantize in one launch) for the MoE prefill MMA arms. Byte-identical to the two-pass path (kernel-check gated) — default ON. MEMRA_MOE_FUSE_ACTQ=0 is the rollback/A-B seam.

router_batch_on

router_kernel_on

router_prefill_exact_on

PREFILL router m-invariance (lane/concat-prime-exact, 2026-08-02). The batched cuBLASLt router GEMM changes a row’s logits when OTHER rows join the call (probed: first change at m=65 on the Ornith-35B router, 3.9e-3 — while the MMQ/f16 trunk GEMMs are bit-identical across m). Feeding a top-k discontinuity, that made a served request’s expert selection a function of its CO-ARRIVALS under cross-request prime batching. The in-house router GEMV is m-invariant, so prefill uses it too and routing depends on a session’s own tokens only. DEFAULT ON: it is the serving isolation contract, and it is the same kernel decode and spec verify already use (dispatch parity, one router kernel for every t). MEMRA_ROUTER_PREFILL_EXACT=0 reverts to the batched GEMM.