Crate memra_engine

Source

Expand description

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_decode_graph: glm5_next T=1 DECODE CUDA GRAPHS, per pipeline stage (door MEMRA_GLM5_DECODE_GRAPH, default ON since 2026-09-04, =0 disarms — lane/b200-glm5-graph-20260902).
glm5_sel_ledger: GATE-HARNESS SELECTION LEDGER for the glm5_next decode-graph door — never a serving flag.
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).
moe_sel_dump: MEMRA_MOE_SEL_DUMP=<path> (default OFF): the per-token routed-expert selection dump behind the expert co-activation question (lane/moe-coactivation-20260902, owner question 2026-09-02: can GLM-5.3-Flash’s 288 routed experts per MoE layer be split across two cards by co-activation, always-active experts replicated on both, so a token’s 8 selected experts rarely cross cards?). The host tool tools/moe_coact.py reads this file.
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).
progress: FORWARD-PROGRESS ODOMETER, the engine’s own answer to “is this worker busy, or hung?” (lane/health-busy-vs-hung, memra#50, 2026-09-03).
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_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).
F32_GEMV_KERNEL_DISPATCHES: Launches gemv_f32_rows took under MEMRA_F32_GEMV_KERNEL=1 (gate non-vacuity).
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 — the v3 arm 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.
GLM5_DECODE_GRAPH_CAPTURES
GLM5_DECODE_GRAPH_LAYERS
GLM5_DECODE_GRAPH_MLA_HALVES: MLA layers captured in halves (an eager middle between two graph pieces) by MEMRA_GLM5_GRAPH_MLA, summed over captures.
GLM5_DECODE_GRAPH_RECAPTURES: Stages torn down and rebuilt by the armed re-capture path (MEMRA_GLM5_GRAPH_RECAPTURE).
GLM5_DECODE_GRAPH_REPLAYS: Captured-run replays (one per graph launch), captures, and the layer count currently covered by captured runs — the door’s engagement receipt, read by the gate bin.
GLM5_Q8_FUSE_ATTN_DISPATCHES: Engagement counter for MEMRA_GLM5_Q8_FUSE_ATTN; gates take a delta.
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.
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).
MLA_WO_ZQ8_DISPATCHES: Launches whose fused pair the wo MMVQ consumed (gate non-vacuity, box engagement receipt).
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_DOWN_ILP2_DISPATCHES: Launches of the _ilp2 down twins (gate non-vacuity, box engagement receipt).
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_GATEUP_ILP2_DISPATCHES: Launches of the _ilp2 gate/up twins (gate non-vacuity, box engagement receipt).
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_SHEXP_OVERLAP_DISPATCHES: MoE layers whose shared expert ran on the side stream (MEMRA_MOE_SHEXP_OVERLAP=1).
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_ILP_DISPATCHES: Engagement counter for the ILP verify-rows MoE door (MEMRA_MOE_VROWS_ILP), both launches.
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).
NVFP4_ROW_ILP_DISPATCHES: Engagement counter for MEMRA_NVFP4_ROW_ILP; gates take a delta.
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).
Q8_ROW_ILP_DISPATCHES: Engagement counter for MEMRA_Q8_ROW_ILP (both launch sites); gates take a delta.
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§

alloc_trace_on: MEMRA_ALLOC_TRACE=1 (gate-harness instrument, default OFF, never a serving flag): every Engine device allocation funnel (alloc_uninit and the zeroed / typed / host-upload wrappers) prints [alloc-trace] <bytes> bytes from <file>:<line> naming the CALLER (#[track_caller]). Why: the nsys trace of the B200 GLM-5.3-Flash decode (2026-09-03) counted ~1,545 cuMemAllocAsync / cuMemFreeAsync pairs per token (1.8 ms of host API time per token) with no host stacks; this names the Rust lines that churn the pool.
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.
dtoh_trace_on: MEMRA_DTOH_TRACE=1 (gate-harness instrument, default OFF, never a serving flag): every Engine device-to-host copy prints one line, [dtoh-trace] <bytes> bytes from <file>:<line>, naming the CALLER of the wrapper (#[track_caller]). Why: an nsys trace of the B200 decode (2026-09-03) showed two blocking DtoH calls per token (4 B after argmax_final_f32, 2112 B after moe_router_sigmoid_topk_f32) each costing ~1.3 ms of queue drain, and the trace has no host stacks; this names the Rust line that owns each drain.
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.
f32_gemv_kernel_on: MEMRA_F32_GEMV_KERNEL=1 (lane/f32-gemv-rows-20260905, default OFF pending its model-scale row): the f32-resident linear at the decode/verify tier (m <= 16) takes the native gemv_f32_rows kernel (one block per (row, token), fixed reduction tree) instead of cuBLASLt, whose m=1 path is a dot_kernel + reduce_1Block_kernel PAIR: two launches and ~9 us of host latency each, 33 pairs per token on the eager MLA layers (the DSA indexer’s wk, kpool_gate, weights_proj). NUMERIC CLASS (cuBLAS’s split is its own): tolerance + determinism + m-identity gate tests/f32_gemv_rows_gpu.rs. Shapes that do not fit (in_f % 1024 != 0, out_f > 65535, m > 16) keep cuBLASLt. Read per call.
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.
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. 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).
glm5_decode_graph_on: Door MEMRA_GLM5_DECODE_GRAPH (lane/b200-glm5-graph-20260902, DEFAULT ON since 2026-09-04): capture the glm5_next T=1 decode walk as replayable per-stage CUDA graphs instead of issuing every kernel per token. Unset/1 arms it, =0 is the eager walk. Read PER CALL so =0 is a live rollback seam, never a process-lifetime latch.
glm5_decode_graph_on_from: The pure parse behind glm5_decode_graph_on: only an explicit 0 disarms the door; unset, 1, and any other value arm it. Kept separate so the default can be unit-tested without mutating the process environment (the OFF arm of every gate sets =0 and this is the contract that makes that arm non-vacuous).
glm5_graph_host_moe: MEMRA_GLM5_GRAPH_HOST_MOE=1 — BISECT knob for MEMRA_GLM5_DECODE_GRAPH, gate harness only. The door has TWO enablers and box run 5 showed they fail independently: (1) the T=1 device-table MoE arm that removes the per-layer router readback, and (2) the capture/replay itself. With this set the door stays ON but the MoE arm stands down to the host oracle, and the capture then refuses BY NAME (a host readback inside a capture region is illegal), so a run isolates enabler 2’s absence from enabler 1’s behaviour instead of confounding them.
glm5_graph_mla_on: MEMRA_GLM5_GRAPH_MLA=1 (lane/mla-half-capture-20260905): the decode graph door captures MLA/DSA layers in HALVES. The attention-site hc pre + norm + the PRE segment (projections, norms, splits, rope: no position-derived launch geometry) close a run’s graph piece, an eager middle (append, k-pool selection, attention, decompress, wo) runs on the stage’s workspace, and the layer’s second half (hc post + FFN) opens the next graph piece of the same run. Needs MEMRA_MLA_SEG_WS=1 (the PRE outputs live in the session’s segment workspace); without it the plan is the KDA-only plan. Read per call. Default OFF pending its model-scale row.
glm5_graph_no_capture: MEMRA_GLM5_GRAPH_NO_CAPTURE=1 — THE OTHER HALF OF THE BISECT, and the half that was missing.
glm5_graph_recapture_on: MEMRA_GLM5_VROWS_T1_DEV=1 — the T=1 device-table MoE arm, forced ON with no capture, no graph, and no MEMRA_GLM5_DECODE_GRAPH anywhere in the run. Default OFF, gate harness only.
glm5_graph_trace_on: MEMRA_GLM5_GRAPH_TRACE=1 — GATE-HARNESS trace for MEMRA_GLM5_DECODE_GRAPH, never a serving flag. Prints one line per captured-run boundary per token, on BOTH arms and at the SAME layer boundaries, with a checksum of the stream state leaving that segment. Box run 4 produced token 0 at every step with the door running cleanly and no error anywhere: the only way to tell “the captured range wrote nothing the remainder reads” from “the state is wrong from layer N onward” is to compare the two arms segment by segment, and nz= in the line separates an all-zero hidden from a wrong-but-live one on sight.
glm5_q8_fuse_attn_on: MEMRA_GLM5_Q8_FUSE_ATTN=1 (lane/glm5-attn-norm-zq8-20260904, default OFF): the ATTENTION-input norm of a plain KDA layer in the glm5_next T=1 walk runs rms_norm_zq8_f32 and hands its q8_1 view to the fused six-projection launcher, which then skips its own quantize_q8_1_into. The FFN-input twin is MEMRA_GLM5_Q8_FUSE. Read PER CALL. Why and receipts: docs/FLAGS.md.
glm5_q8_fuse_on: Door MEMRA_GLM5_Q8_FUSE (lane/b200-q8-fuse-20260902, DEFAULT OFF pending the box A/B): on the glm5_next mHC decode trunk (hyper_range_decode / hyper_range_decode_ws_body), fold the FFN-input rms_norm and its consumer’s standalone quantize_q8_1 launch into ONE rms_norm_zq8_f32 launch. Byte-identical to the unfused chain (see that kernel’s header in cu/kernels.cu); this door only changes launch count. See docs/FLAGS.md and research/b200-q8-fuse-20260902/LANE.md.
glm5_trace_reset: Clear the trace budget. The gate calls this at every arm switch; without it the first arm’s exhaustion silences the second (take 10).
glm5_vrows_t1_dev_forced
glm5_vrows_t1_dev_forced_from: The pure parse behind glm5_vrows_t1_dev_forced (arch-keyed since 2026-09-04): 1 forces the T=1 device-table MoE arm on every eager layer, 0 leaves the eager layers on the host readback (the arm still engages inside a capture, which is keyed on the open region, not on this), and UNSET forces it on 100a builds. Receipt (darklanes research/glm5-b200-20260902/LANE.md, t1devab 2026-09-04, 2x B200 pair, composed defaults): host 79.60/79.81/79.31 -> 79.60 vs forced 84.11/84.19/83.29 -> 84.11, +5.67%, tape 9437b599f6b9d2a9 on all six boots; the 11 eager MLA-layer MoE calls each lose a pinned readback plus a device drain. no_capture_bisect is the MEMRA_GLM5_GRAPH_NO_CAPTURE knob’s own forcing, unchanged.
hc_pre_block_default: The build-arch default of MEMRA_HC_PRE_BLOCK: 512 on 100a, 128 elsewhere.
hc_pre_sink_reg_from: The pure parse behind [hc_pre_sink_reg] (arch-keyed since 2026-09-04): 1 arms, 0 disarms, unset = ON on 100a builds (receipt +7.86% alone, +10.25% with the 512-wide block, tape 9437b599f6b9d2a9, darklanes research/glm5-b200-20260902/LANE.md hcpreab), OFF elsewhere.
hc_pre_v4_on: MEMRA_HC_PRE_V4=1 (lane/hc-pre-phases-20260905): run the hc pre-chain on the v4 register schedule (dsv4_hc_pre_v4_*_kernel: x loaded once and kept, two barriers, warp 0’s Sinkhorn overlapped with the other warps’ combine) instead of _v3. Same arithmetic in the same order, bit-identical by construction (gate tests/hc_pre_v4_gpu.rs); the launcher refuses shapes the register schedule does not fit (40025) and the caller then runs v3 unchanged. Default OFF until its model-scale row (new-flags law).
hc_pre_v4z_on: MEMRA_HC_PRE_V4Z=1 (lane/hc-pre-v4z-20260905): under MEMRA_HC_PRE_ZQ8=1, the fused hc-pre + norm launch is the v4 register schedule with rms_norm_zq8_f32_v2 replayed inside the block (dsv4_hc_pre_v4z_e16_kernel), every norm operation pinned to the served kernel’s compiled form. Refuses (falls back to the zq8 kernel) off the served shape. Default OFF pending its model-scale row.
hc_pre_zq8_on: MEMRA_HC_PRE_ZQ8 (lane/hcpre-zq8-fusion-20260905): run each hc site’s pre-chain AND the rms_norm_zq8 that consumes its output as ONE launch (dsv4_hc_pre_zq8_kernel).
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 block geometry lives in the shared memra-kv crate (Phase D); re-exported so every existing crate::kv_blk_bytes() call site is unchanged. Per-32-element block bytes of the trunk KV cache: q8_0 K (34 B) and q5_1 V (24 B), the one validated config. The env-selected format arms (MEMRA_KV_K / MEMRA_KV_V: fp8 K, q4_0 / fp8 V) were removed 2026-09-05 (door sweep); gemma’s e4m3 layers size themselves in full_attention_kv_layout below.
mla_absorb_bf16_on: MEMRA_MLA_ABSORB_BF16=1 (lane/mla-absorb-bf16-20260905): the MLA absorb planes wk_b / wv_b are read as BF16 by the decode _wp kernels instead of the f32 copy the loader materializes (738 -> 369 MB per token on GLM-5.3-Flash). Exact where the source plane is BF16 (the B200 hybrid mint): the f32 copy is a widening, the kernel widens the same bits again, same products in the same order. The BF16 copy is built at load only if every element round-trips; otherwise the layer keeps the f32 path. Default OFF pending its row.
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.
mla_wo_zq8_on: MEMRA_MLA_WO_ZQ8=1 (lane/mla-wo-zq8-20260905, default OFF pending its model-scale row): on the decode MLA core (t=1, not the verify-rows arm) the coalesce-arm decompress_v launch emits wo’s q8_1 pair beside the f32 attention output (memra_mla_decompress_v_wp_zq8_kernel, and the BF16 plane twin under MEMRA_MLA_ABSORB_BF16), and the wo projection takes it through matmul_q8_fast: the standalone quantize_q8_1 before wo is gone (11 per token on GLM-5.3-Flash, in the eager MLA middle where each launch costs ~9 us of host latency). BIT-IDENTICAL (gate tests/mla_wo_zq8_gpu.rs). Shapes the epilogue cannot own (d_v / split not a whole number of q8 blocks, or no coalesce arm) keep the plain sequence. Read per call.
moe_down_ilp2_on: MEMRA_MOE_DOWN_ILP2=1 (lane/moe-down-ilp2-20260905, default OFF pending its model-scale row): the verify-rows MoE down/FMA launch takes the _ilp2 twins (moe_down8_fma_q8_rows_ilp2, _w4_ilp2) that walk two experts at once (8 groups in flight per lane instead of 4, one warp reduction per expert as before) instead of the _ilp twins that walk them one at a time. Rides on top of MEMRA_MOE_VROWS_ILP (interleaved NVFP4 only; a non-ILP or non-NVFP4 launch keeps its kernel). BIT-IDENTICAL by construction: each expert keeps its own accumulator and g-order, its own warp_reduce_sum, and the slot-ordered __fmaf_rn chain is unchanged (gate tests/moe_down_ilp2_gpu.rs). Read per call.
moe_expert_rp_on: MEMRA_MOE_EXPERT_RP=1 (default OFF, memra#147): the device-RESIDENT NVFP4 expert slabs are repacked at upload into the slot-major per-row layout the engine already names QT_NVFP4_V2 (per row: slot g’s 16 quant bytes at g16, its two UE4M3 scale bytes at nsb16 + g*2; nvfp4_expert_split_repack, the same bytes as tp.rs nvfp4_matrix_v2_permute) and DevExps::rp is set. Readers are told QT_NVFP4_V2 (rp_qt): every expert dot goes through expert_dot_g’s V2 case on the shared pinned core (one 16B window per lane-group at a 16B lane stride instead of five scattered 4B loads at a 36B stride: root ncu measured 24.97 sectors per warp request on moe_gate_up_preclamp8_q8_w4, 4 is coalesced), and the grouped prefill takes its existing V2 dequant / kq_fetch<V2> arms. Host bytes, the SLRU cache and the TP upload paths stay interleaved and untouched. A resident-slab reader not yet handed the V2 qtype refuses with a named error (moe_rp_refuse) rather than reading repacked bytes interleaved.
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). 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_gateup_ilp2_on: MEMRA_MOE_GATEUP_ILP2=1 (lane/moe-gateup-ilp2-20260905, default OFF pending its model-scale row): the verify-rows MoE gate/up launch takes the _ilp2 twins (moe_gate_up_preclamp8_q8_rows_ilp2, _w4_ilp2) that give a warp TWO pairs at the same expert-FFN row (16 groups in flight per lane; at t=1 both experts share the token’s activation loads) instead of one. Rides on MEMRA_MOE_VROWS_ILP (interleaved NVFP4 only) and yields to the _ord schedule (MEMRA_MOE_VROWS_ORD). BIT-IDENTICAL by construction: each pair keeps its own accumulators, g-order, reductions, SwiGLU and store (gate tests/moe_gateup_ilp2_gpu.rs). Read per call.
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_rp_refuse: A resident-slab reader that has no split-plane arm refuses, by name, instead of reading the repacked bytes with the interleaved walk (which would be a plausible-looking wrong answer).
moe_shexp_overlap_main_probe: MEMRA_MOE_SHEXP_OVERLAP=3: diagnostic arm, the fork’s program on the MAIN stream with no override and no side stream (still computed before the routed rows). Separates “the side stream / override changes the program” from “computing it earlier changes its inputs”.
moe_shexp_overlap_on: MEMRA_MOE_SHEXP_OVERLAP=1 (lane/moe-shexp-overlap-20260905): in the t=1 verify-rows MoE walk the SHARED expert (gate/up fused, SwiGLU, down) runs on Engine::side_stream while the routed experts’ rows run on the main stream; the two are forked and joined with events and the shared output is added with the same add_scaled_rows_ones as before. Same kernels, same operands, same add: bit-identical; ~17 us per MoE layer hidden behind ~60 us of routed work. Read per call (the gates flip it in-process). Default OFF pending its row.
moe_shexp_overlap_serial_probe: MEMRA_MOE_SHEXP_OVERLAP=2: diagnostic arm, the shared expert still runs on the side stream but the main stream joins it BEFORE the routed rows (no concurrency). Separates “the side stream program differs” from “the two run concurrently and race” when the identity gate fails.
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_ilp_dispatches: Snapshot of MOE_VROWS_ILP_DISPATCHES — gates take a before/after delta.
moe_vrows_ilp_on_from: The pure parse behind [moe_vrows_ilp_on]: 1 arms, 0 disarms, unset follows the BUILD ARCH (ON for 100a, OFF otherwise), the per-hardware arm selection law: the twins carry a 2x B200 receipt (+6.0% at c1, darklanes research/glm5-b200-20260902/LANE.md, ilpab) and no SM120 one, so an sm_120a build keeps its measured default until it has its own.
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.
nvfp4_row_ilp_dispatches: Snapshot of NVFP4_ROW_ILP_DISPATCHES.
nvfp4_row_ilp_on_from: The pure parse behind [nvfp4_row_ilp_on]: 1 arms, 0 disarms, unset follows the BUILD ARCH (ON for 100a, OFF otherwise): the twins carry a 2x B200 receipt (+1.98% alone, +2.55% with the grid fill, darklanes research/glm5-b200-20260902/LANE.md, nvab) and no SM120 one.
q8_row_ilp_dispatches: Snapshot of Q8_ROW_ILP_DISPATCHES.
q8_row_ilp_on_from: The pure parse behind [q8_row_ilp_on]: 1 arms, 0 disarms, unset follows the BUILD ARCH (ON for 100a, OFF otherwise): the twins carry a 2x B200 receipt (+2.16% at c1, darklanes research/glm5-b200-20260902/LANE.md, q8ab) and no SM120 one.
router_batch_on
rp_qt: The qtype a kernel is told for an expert slab: QT_NVFP4_V2 (the slot-major per-row layout, tp.rs nvfp4_matrix_v2_permute) when the slab it will read is a repacked resident slab, the tensor’s own qtype otherwise.
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.

Crate memra_engine

Crate memra_engine Copy item path

Re-exports§

Modules§

Structs§

Constants§

Statics§

Functions§

Crate memra_engine