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).

dsv4_ffi

FFI declarations for the DeepSeek-V4-Flash GPU kernels (cu/dsv4_gpu.cu, lane 4).

dsv4_gpu

DeepSeek-V4-Flash GPU trunk forward (lane 4): 2-card layer-split placement, correctness bring-up gated against the lane-2/3 CPU oracle fixtures.

eagle

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

ep_map

Measured expert-placement map (MEMRA_EP_MAP; glm5 alias honored) — the fail-closed memra-ep-map-v1 reader every family’s EP shard builders consume (fleet-shared by design; glm5 is the first consumer). (LAW:coactivation-expert-placement; maps are minted by the shared fleet tool from MEMRA_MOE_WEIGHT_TRACE traces). No CUDA deps. Measured expert-placement map consumption (MEMRA_EP_MAP, lane/glm5-ep-place 2026-08-31, generalized lane/glm5-extract-general) — the fail-closed memra-ep-map-v1 reader the shard builders trust. FLEET-SHARED BY DESIGN: one flag, one parser, one validation law for every family that arms measured placement (glm5 today; hy3/qwen adopt the same seam). Family loaders add only their own geometry laws (rank count, entry rank, layer cover) on top of the parsed map.

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.

glm5_tp

glm5_next (GLM-5.3-Flash) TP-N seam — MEMRA_GLM5_TP (lane/glm5-tp2, 2026-08-31; rank-widened to TP-4 by lane/glm5-composition, 2026-09-01).

glm_spec

glm5_next T-parallel speculative verify: the rows-walk verify, per-step KDA state-column rollback, latent/kpool truncation, and the MEMRA_GLM5_SPEC-gated draft->verify->rollback loop over the native MTP head (lane/glm5-tparallel-verify). glm5_next T-parallel speculative verify: the acceptance/rollback machinery that turns the native MTP draft head (mtp_head_forward_mla_cached, lane/glm5-mtp-remint) into served speculative decoding on the HyperConnections trunk.

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.

hyper

mHC — manifold-constrained hyper-connections, the ResidualTopology::HyperConnections residual program (glm5_next / GLM-5.3-Flash, and the dsv4 class).

kda

Kimi Delta Attention (KDA) — the glm5_next (GLM-5.3-Flash) linear-attention mixer.

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.

mla_ffi

FFI declarations + safe Engine wrappers for the MLA CUDA forward (cu/mla_attn.cu).

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.

parallel

ModelPlan-driven parallel topology and artifact placement contracts.

plan_backend

Compatibility re-export for build-time execution manifests owned by the ModelPlan compiler.

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).

qwen4exp_gpu

qwen4_exp (Qwen3.8-Flash-Next) GPU eager forward — onboarding phase 7, correctness arm gated against memra-reference (research/qwen4exp-bringup-20260829/GPU-EAGER.md). qwen4_exp (Qwen3.8-Flash-Next) GPU EAGER forward — onboarding-ladder phase 7, eager arm.

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).

spec_phase

Per-burst spec-round phase attribution (MEMRA_SPEC_TRACE; glm5 alias honored) — the draft/verify/accept/rollback/maintenance split every spec family owns, with caller-tagged emit lines so banked receipts keep their grep shape. No CUDA deps beyond the stream drains at phase boundaries. Per-burst phase attribution for speculative rounds (MEMRA_SPEC_TRACE, generalized lane/glm5-extract-general from the glm5 loop’s MEMRA_GLM5_SPEC_TRACE — the alias stays honored). The draft / verify / accept / rollback / source-maintenance split is SPEC-FAMILY-GENERIC: any spec loop owns those five boundaries, and the level-2 verify sub-split buckets are MIXER-CLASS buckets (KDA, MLA — multi-family classes), not one model’s. The emit TAGS are the caller’s, so a family’s banked receipts keep their exact grep shape ([glm5-phase] / [glm5-phase-v] for the glm5 loop).

spill

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

tp

Tensor-parallel correctness runtime.

tp_transport

TP-N TRANSPORT — MEMRA_TP_TRANSPORT (born lane/glm5-tp-transport 2026-09-01 as MEMRA_GLM5_TP_TRANSPORT, rank-widened by lane/glm5-composition 2026-09-01, generalized lane/glm5-extract2; the family alias stays honored).

vision

Vision tower for Qwen3.8-27B multimodal input (lane/vision, 2026-08-15).

vision_gemma

Vision tower for the gemma-4 family (lane/gemma-vision, 2026-08-16).

vision_glm5

Vision tower for GLM-5.3-Flash / glm5_next (lane/glm5-vision, 2026-08-30).

vision_pre

Host preprocessor for vision input (lane/vision): bytes -> ViT patch rows.

vision_step

Vision tower for StepFun Step-3.7-Flash (arch step35), lane/step37-vision 2026-08-30.

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).

ExactScope

RAII guard from Engine::exact_scope: restores the pre-scope verify_exact value on drop, so error propagation (?) can never leave the engine latched in the decode-exact matmul program (hermes finding, fixed 2026-08-23). Holds the flag, not the Engine, so the restoration contract is unit-testable without a 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

PinnedHostBuf

Owned page-locked CACHEABLE host buffer (flags=0, deliberately NOT write-combined) for the prefix-cache host tier (lane/kv-host-spill-20260830). Same allocation class as PinnedStage above and for the same reason: ctx().alloc_pinned is CU_MEMHOSTALLOC_WRITECOMBINED, which is right for H2D-only staging but pathologically slow for host READS (see the HostBuf CAVEAT in model.rs), and these bytes are CPU-read by the MEMRA_KV_HOST_VERIFY digest arm. Public because the server’s host-tier cache owns these buffers across requests.

VerifyWs

Size-keyed device-buffer free-lists for the verify walk (door W — see the field doc on Engine::verify_ws). Exact-length keying: the walk’s shapes quantize to a few classes per round (t in 2..=8 times fixed widths), so hit rates are structural, and an exact-size buffer keeps every debug_assert_eq!(len, ...) at the launchers intact.

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_NVFP4_V2

Slot-major v2 bank permutation of QT_NVFP4 (see tp.rs nvfp4_matrix_v2_permute) — only the grouped-prefill dequant consumes this tag; every direct/dp4a lane must keep refusing it.

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

BF16_TCOLS_RED_FUSED_DISPATCHES

Engagement counter for the fused-reduce-tail tcols door (MEMRA_BF16_TCOLS_RED_FUSED), incremented at the door’s own dispatch (LAW:wiring-assertions-match-prose).

BF16_TCOLS_WIDE_DISPATCHES

Engagement counter for the wide-t tcols door (MEMRA_BF16_TCOLS_WIDE), incremented at the door’s own dispatch (LAW:wiring-assertions-match-prose). Read via bf16_tcols_wide_dispatches.

BF16_TCOLS_X1_DISPATCHES

Engagement counter for the x1-grid tcols door (MEMRA_BF16_TCOLS_X1).

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.

HTOD_DIET_AVOIDED

HtoD calls avoided by door H (MEMRA_HTOD_DIET): the count receipt. A gate asserts it tracks the layer-call count on the ON arm and stays flat on the OFF arm.

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.

MLA_TC_PREFILL_DISPATCHES

Engagement counter for the glm5_next tensor-core MLA prefill chain (MEMRA_MLA_TC_PREFILL), incremented once per (layer, chunk) dispatch at the chain’s own invocation, AFTER the strided-batched GEMM decline check — a declined shape does not count. A gate that must prove “the TC arm ran N times for this workload” reads this delta; the once-per-boot announce line dedups and cannot carry a count (LAW:wiring-assertions-match-prose).

MMQ_SK_FORCE

Per-model stream-k entry override for SPEC serving (-1 = unset → env/default; 0 = force tiling; 1 = admit the deterministic form selector). The former timing selector made identical boots choose different fold orders; MEMRA_MMQ_SK_FORM is the explicit numerical-form seam. mmq_ffi reads this before the env.

MOE_FUSED_EPI_DISPATCHES

Fused MoE-epilogue dispatches taken since process start (MEMRA_MOE_FUSED_EPI), incremented once per (token, layer) that actually runs moe_fused_epi_token_q8.

MOE_GROUPED_PREFILL_DISPATCHES

Engagement counter for the glm5_next expert-grouped MoE PREFILL arm (MEMRA_MOE_GROUPED_PREFILL), incremented once per (layer, chunk) dispatch at the arm’s own call site. Same reason the fused-epilogue counter exists: the observation env vars divert dispatch, so a counter at the invocation is the only honest engagement receipt (LAW:wiring-assertions-match-prose). Read via moe_grouped_prefill_dispatches.

MOE_VROWS_DEDUP_ORDER_DISPATCHES

Engagement counter for the expert-major gate/up schedule (MEMRA_MOE_VROWS_DEDUP_ORDER).

MOE_VROWS_DEV_TABLES_DISPATCHES

Engagement counter for the device-side vrows table build (MEMRA_MOE_VROWS_DEV_TABLES).

MOE_VROWS_DISPATCHES

Verify-rows batched MoE dispatches taken since process start (lane/glm5-vrest): incremented once per (layer, verify-call) that runs the pairs-shaped routed-expert program (moe_gate_up_preclamp8_q8_rows + moe_down8_fma_q8_rows) instead of the per-(token,expert) sequential loop. Rides MEMRA_GLM5_VERIFY_BATCH’s arm — no flag of its own. Same rationale as MOE_FUSED_EPI_DISPATCHES: the observation envs divert dispatch, so gates anchor on the arm’s own invocation (LAW:wiring-assertions-match-prose).

MOE_VROWS_DOWN_TMAJ_DISPATCHES

Engagement counter for the token-major down schedule (MEMRA_MOE_VROWS_DOWN_TMAJ).

MOE_VROWS_PACK_DISPATCHES

Engagement counter for the warp-packed verify-rows MoE door (MEMRA_MOE_VROWS_PACK).

MOE_VROWS_PAIR_DISTINCT

DISTINCT experts in the pair union, summed over the same layer-calls. The dedup lever is 1 - distinct/visits; equal counters mean routing is disjoint across the verify rows and there is no byte to save.

MOE_VROWS_PAIR_VISITS

Expert VISITS (t x n_used) summed over vrows layer-calls under MEMRA_MOE_VROWS_DEDUP_STAT.

MOE_VROWS_ROUTER_SYNCS_AVOIDED

Router readbacks (one full cuStreamSynchronize + 2 DtoH each) that door D skipped. The count receipt for the host seam: a gate asserts it moves 1:1 with MOE_VROWS_DEV_TABLES_DISPATCHES on the ON arm and stays flat on the OFF arm.

MOE_VROWS_SLAB_READS_AVOIDED

AVOIDED SLAB READS — the box receipt for door E. Every layer-call adds visits - distinct, i.e. the expert-slab reads whose repeat visit the expert-major schedule places inside the reuse window. Multiply by the per-visit slab bytes (gate+up 9.4372 MB, down 4.7186 MB at the serving geometry) for the bytes the schedule makes avoidable; that product is the CEILING of the win, not the win (the realized share is a cache/scheduling property the box prices).

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.

SCRATCH_ALLOC_CALLS

Device-scratch allocation census (lane/glm5-decode-diet): bumped by every alloc_uninit and zeros call — the class the launch-diet census measured at 2,358 cuMemAllocAsync+Free calls/token. The decode-workspace gate reads deltas per step; the cost axis is the CALL COUNT (the box’s measured ~1.06 us/driver call), which is exactly what this counts. Relaxed atomic: one increment per allocation, noise-level.

TOPK_SHARDS_DISPATCHES

Engagement counter for the sharded top-k door (MEMRA_TOPK_SHARDS).

VERIFY_WS_HITS

Engagement counter for the verify-walk workspace (MEMRA_VERIFY_WS): incremented once per POOL HIT (a reused buffer = one avoided alloc + one avoided free). Gates anchor on the delta; SCRATCH_ALLOC_CALLS carries the complementary real-alloc count.

Functions

arm_step37_serving_defaults

bf16_tcols_red_fused_dispatches

Snapshot of BF16_TCOLS_RED_FUSED_DISPATCHES — gates take a before/after delta.

bf16_tcols_wide_dispatches

Snapshot of BF16_TCOLS_WIDE_DISPATCHES — gates take a before/after delta.

bf16_tcols_x1_dispatches

Snapshot of BF16_TCOLS_X1_DISPATCHES — gates take a before/after delta.

ep_diet_on

MEMRA_EP_DIET=1 (default OFF; generalized from MEMRA_GLM5_EP_DIET, which stays honored per the flag-alias law — lane/glm5-ep-diet): the EP DISPATCH DIET door, general to any expert-parallel MoE walk. What the door names is a movement CLASS, not a family: one bulk peer activation fan-out per layer-call instead of per-token uploads, compact peer staging with one bulk return instead of a per-slot round-trip dribble, and one scatter launch instead of the t*n_used sequential axpy chain. The glm5 TP-2 walk is today’s CONSUMER (its kernels, its combine order, its counters in glm5_tp.rs); hy3/step EP walks arm the same door for their own walks.

ep_grouped_prime_on

MEMRA_EP_GROUPED_PRIME=1 (default OFF; generalized from MEMRA_GLM5_EP_GROUPED_PRIME, which stays honored per the flag-alias law — lane/glm5-ep-diet): the EP GROUPED-PRIME door, general to any expert-parallel MoE walk — “run the family’s own chunked grouped MoE prefill program per rank over each rank’s resident expert slab, then add the peer’s bulk-returned partial”. The glm5 TP-2 walk is today’s consumer.

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).

htod_diet_avoided

Snapshot of HTOD_DIET_AVOIDED — gates take a before/after delta.

htod_diet_on

MEMRA_HTOD_DIET=1 (default OFF; generalized from MEMRA_GLM5_HTOD_DIET, which stays honored per the flag-alias law above — door H, lane/glm5-moe-loc): ENGINE-GENERIC HtoD hygiene. Nothing in either class is family knowledge; both are “the host uploaded bytes the device already had”.

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.

mla_tc_prefill_dispatches

Snapshot of MLA_TC_PREFILL_DISPATCHES. Gates take a before/after pair around a workload and assert on the delta — including the DECODE byte-identity gate, whose assertion is that this stays FLAT across t=1 steps with the flag on.

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.

moe_fused_epilogue_dispatches

Snapshot of MOE_FUSED_EPI_DISPATCHES. Gates take a before/after pair around a workload and assert on the delta, anchoring on the arm’s own invocation rather than on a flag being set (LAW:wiring-assertions-match-prose).

moe_grouped_prefill_dispatches

Snapshot of MOE_GROUPED_PREFILL_DISPATCHES. Gates take a before/after pair around a workload and assert on the delta.

moe_vrows_dedup_order_dispatches

Snapshot of MOE_VROWS_DEDUP_ORDER_DISPATCHES — gates take a before/after delta.

moe_vrows_dev_tables_dispatches

Snapshot of MOE_VROWS_DEV_TABLES_DISPATCHES — gates take a before/after delta.

moe_vrows_dispatches

Snapshot of MOE_VROWS_DISPATCHES — gates take a before/after delta around a workload.

moe_vrows_down_tmaj_dispatches

Snapshot of MOE_VROWS_DOWN_TMAJ_DISPATCHES — gates take a before/after delta.

moe_vrows_pack_dispatches

Snapshot of MOE_VROWS_PACK_DISPATCHES — gates take a before/after delta.

moe_vrows_pair_overlap

Snapshot of the dedup instrument as (visits, distinct).

moe_vrows_router_syncs_avoided

Snapshot of MOE_VROWS_ROUTER_SYNCS_AVOIDED.

moe_vrows_slab_reads_avoided

Snapshot of MOE_VROWS_SLAB_READS_AVOIDED.

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.

topk_shards_dispatches

Snapshot of TOPK_SHARDS_DISPATCHES — gates take a before/after delta.

verify_ws_hits

Snapshot of VERIFY_WS_HITS — gates take a before/after delta.

vrows_expert_major_order_for_test

Gate hook for [vrows_expert_major_order] — the permutation is the whole door, so it is gated against the device build and on planted selections rather than inferred from a live tape.

vrows_overlap_counts_for_test

Gate hook for [vrows_overlap_counts] — the counting is the whole instrument, so it is gated on planted overlaps (disjoint / partial / identical) rather than inferred from a live tape.