Module ops

Expand description

2026-09-25: Kernel launch wrappers that the layers compose into forward passes, grouped into the submodules under ops/ and re-exported here (most by glob).

Owner: model-layers (ops). Invariants: none beyond the types.

Re-exports§

pub use model_stats::ModelStats;
pub use moe_lora_grouped::*;

Modules§

dense_gemv_tc: 2026-09-25: Routes the MTP drafter’s BF16 small-M GEMM rows to the tensor-core GEMV entries of dense_gemv_bf16_tc.cu.
gemv_tc: 2026-09-25: Routing of the narrow NVFP4 batched GEMV to its tensor-core kernels (w4a16_gemv_tc8 / w4a16_gemv_tc16 in kernels/gb10/common/w4a16_gemv_tc.cu).
lora_delta: 2026-09-25: Runtime LoRA delta y += scale * (x @ A^T) @ B^T on BF16 activations: the single-adapter path and the per-request routed path.
model_stats: 2026-09-25: ModelStats: diagnostic counters and one-shot latches owned by one model.
moe_lora_grouped: 2026-09-25: Launchers for the device-side MoE expert LoRA folds: moe_lora_grouped_down over the expert-sorted prefill rows and moe_lora_gather_bgmv over the slot-major decode rows, plus the host packing of their per-expert tables.
target_defaults: 2026-09-25: The compiled target’s serving defaults, resolved against the environment: the target’s declaration first, a METRALE_* override second.
token_overlay: 2026-09-25: Launchers for the LoRA token-overlay kernels in kernels/gb10/common/token_overlay.cu:
w4a4_proj: 2026-09-28: W4A4 small-M projection on FP4 tensor cores (w4a4_gemv_mx.cu, module w4a4_gemv_mx).

Structs§

AttnFa128Kernels: 2026-09-27: The two twins’ handles; 0 when the target lacks the module or METRALE_NO_ATTN_FA128 is present (the originals then run, for A/B measurement).
BaGatesPick: 2026-09-25: The kernel a launch runs, and why the other did not.
CublasScope: 2026-09-25: Which projection families take a cuBLASLt GEMM arm, so one family can be armed without the others.
DecodeW8a8Plan: 2026-09-25: The shape and output of one decode projection, as decode_w8a8_selected judges it. The route is taken only when all of these hold:
DecodeW8a8Scratch: 2026-09-25: The activation-quant scratch buffers with their capacities, and the quantizer and scale-layout adapter kernels.
DenseMmKernels: 2026-09-25: The three BF16 dense kernels dense_mm_bf16 chooses from, resolved once per projection site (the GLM-5.3 mixer and MLP sites bind one each).
DerivedWeights: 2026-09-25: Per-model memo of derived weight encodings, one map per Derivation.
E4m3Saturation: 2026-09-30: An unscaled E4M3 cast whose saturation is disclosed or known, with the largest |x| the disclosure covers.
E4m3SaturationScope: 2026-09-30: See allow_e4m3_saturation.
Fp8ActQuant: 2026-09-25: The FP8 activation quantizer a layer launches: the shared kernel, plus the Hopper twin when the image has one.
Fp8GroupedGeometry: 2026-09-28: The launch shape of a grouped gate+up or down kernel: output columns per CTA (grid.x = ceil(n / cols_per_cta)), rows per pass (the shared expert takes ceil(num_tokens / rows_per_pass) block rows) and threads per CTA.
Fp8GroupedSortOut: 2026-09-27: The outputs of moe_fp8_grouped_sort: moe_sort_by_expert’s four and the active-expert list with its length.
Fp8GroupedW8a8Layout: 2026-09-28: Byte offsets of the W8A8 activations inside the grouped decode’s buffers. The routed SiLU buffer (expert_gate_out, sized for [te, inter] FP32) holds the routed E4M3 products [te, inter], their scales [te, inter / 128] FP32, the quantized layer input [m, hidden] E4M3 and its scales [m, hidden / 128]; the shared SiLU buffer (logits, [m, inter] FP32) holds the shared products and their scales.
Fp8GroupedW8a8Rows: 2026-09-28: The routed expert tables and the grouped decode’s sort outputs the two W8A8 kernels read.
Fp8QuantPick: 2026-09-25: Which kernel a launch runs, on what grid, and why the other did not. super::Fp8ActQuant::pick_with sets the handle and the grid from one decision.
GdnCarrySizes: 2026-10-01: The device allocations TransformerModel::gdn_carry_bind makes for layers GDN layers and slots carry slots (the MTP slots plus the dummy), in bytes, one field per allocation. The binder allocates exactly these and the preflight reserve plans total().
GemmDispatch: 2026-09-25: Which GEMM implementation each projection takes, resolved by GemmDispatch::from_env.
ModelLevers: 2026-09-25: Kernel-path levers for one loaded model. TransformerModel::new resolves its own copy with ModelLevers::from_env; other readers use the process-wide ModelLevers::get.
MoeCutlassDownHostTables: 2026-09-25: The down-projection tables of MoeCutlassHostTables.
MoeCutlassHostTables: 2026-09-25: Host copies of the per-expert pointer and scale tables for the CUTLASS grouped path, which takes them on the host. MoeLayer::build_cutlass_grouped_sfb builds them at load into the layer’s cutlass_grouped_host, so they are dropped with the layer.
MoeE4m3Tile: 2026-09-27: One entry point of kernels/gb10/common/moe_w8a8_grouped_gemm_e4m3.cu: its work-list item shape (m_tile rows by n_tile columns, so the builder takes n_tiles = N / n_tile), its dynamic shared memory and how many of its 256-thread CTAs fit on one SM.
Nvfp4ExpertTables: 2026-09-27: One NVFP4 projection across the routed experts: device tables of each expert’s packed and block-scale pointers, and its per-tensor scales as f32.
W8a8Kernels: 2026-09-28: The quantizers and the ten GEMV entry points (2 scale layouts x the five row entries, mb1_ku8 to mb16). A handle is 0 when its module is not compiled.
W8a8Scratch: 2026-09-28: Device scratch for one quantized activation: [rows, k] E4M3 and its FP32 scales. One per model; every W8A8 projection on the stream quantizes into it immediately before its GEMV.
W8a8Weight: 2026-09-28: One W8A8 projection: up to three E4M3 weights with the same K and scale layout, stacked by rows into one output (attention Q|K|V, GDN QKV|Z, FFN gate|up). The checkpoint’s tensors are read in place.

Enums§

BaGatesLogSlot: 2026-09-25: The once-flag slot a verdict logs under in ba_gates_log; one slot per outcome, so an early small request cannot hide a later line about the other outcome.
Derivation: 2026-09-25: Which derivation a lookup is for.
F16ChainStates: 2026-10-01: Where the FP16 exact chain finds its states: per-sequence pointer tables (the verify’s WY tables: h, Hi0, Hi1, Hi2), or the bases of a single sequence.
Fp8QuantLogSlot: 2026-09-25: Which once-flag fp8_quant_log uses for a verdict. There is one per branch because the verdict changes with M, so one process can log both the twin line and a refusal.
GdnScalarSpine: 2026-09-28: The scalar spine a process binds (init_kernels::fused_spine_kernel) and launches (ops::gdn_prefill_fla, whose block size and shared memory depend on it). Both read it through gdn_scalar_spine, so they cannot disagree.
HcVariant: 2026-09-25: Which family a site’s weights select. Callers use it to decide whether the block’s input_norm runs on hc_pre’s output (HcVariant::applies_block_input_norm): in the low-rank family hc_norm inside hc_pre is the block’s input norm, and a second RMS pass with a ones-filled weight is not an identity.
ResidentMat: 2026-09-25: A resident [N, K] row-major projection from the GGUF path: expanded bf16, or the raw Q2_K / Q3_K blocks (WeightDtype::Q2K / Q3K) that the K-quant GEMV and MMQ kernels read.
W8a8Scale: 2026-09-28: How a W8A8 weight carries its scales, and so how the activation is quantized: one scale per output row with one per token, or one per 128x128 weight block with one per (token, 128-wide K group).

Constants§

BA_GATES_BLOCK: 2026-09-25: Threads per block, as in the parent (BAH_BLOCK in the twin). The reduction order depends on it, BA_GATES_LANES and BA_GATES_OUTS.
BA_GATES_FALLBACK_SM_COUNT: 2026-09-25: SM count used when GpuBackend::sm_count() fails: the sm_count in kernels/hopper/HARDWARE.toml. A wrong value moves only the guard’s threshold.
BA_GATES_GROUPS: 2026-09-25: BAH_GROUPS in the twin: output groups a thread accumulates per read of the activation row. The row is read ceil(ceil(N/4)/8) * 4 times per token, 12 at N = 96. The kernel source records the register and occupancy table behind the choice of 8.
BA_GATES_LANES: 2026-09-25: Threads that reduce one BA output: the parent’s threads_per_out.
BA_GATES_LOG_SLOTS: 2026-09-25: Once-flags ba_gates_log keeps: one per guard, plus the twin’s.
BA_GATES_OUTS
BA_GATES_REJECTS: 2026-09-25: Every string ssm_ba_gates_hopper_reject can return, in the order it tests them. ba_gates_log gives each its own once-flag by index here; the test every_reject_reason_has_its_own_log_slot checks that every returned reason has an index.
BA_GATES_TOO_FEW_TOKENS: 2026-09-25: The token-count guard’s string, named so tests can match it.
BA_GATES_WARPS: 2026-09-25: Warps per block: the width of one row of the twin’s cross-warp scratch red.
BF16_GROUPED_DOWN_TC: 2026-10-02: moe_expert_down_act_bf16_grouped_tc (BTC_DOWN_COLS columns per CTA).
BF16_GROUPED_GATE_UP_TC: 2026-10-02: moe_expert_gate_up_act_bf16_grouped_tc (BTC_GU_COLS columns per CTA; writes the SiLU product as BF16 hi + lo).
DECODE_W8A8_ROWS: 2026-09-25: The padded decode row counts this route takes. Any other count stays on the caller’s own kernels.
DENSE_GEMM_M16_BF16_K_STEP: 2026-09-25: K granularity: the kernel’s 4-stage cp.async pipeline advances 64 elements per step (DGM16_K_STEP), which also keeps every weight row 16-byte aligned.
DENSE_GEMM_M16_BF16_MAX_M: 2026-09-25: The kernel’s M tile. Rows past it are not computed, so the wrapper returns an error rather than leave part of the output stale.
DENSE_GEMM_M16_BF16_N_TILE: 2026-09-25: N columns one CTA owns in the default instantiation. It must equal the kernel’s DGM16_N_TILE.
DENSE_GEMM_M16_BF16_N_TILE_WIDE: 2026-09-25: N tile of the wide instantiation dense_gemm_m16_bf16_n64 (kernel DGM16_N_TILE_WIDE), selected by METRALE_LM_HEAD_M16_TC_NTILE=64. It launches half the CTAs for a given N, so the A tile is re-read from L2 half as many times.
DENSE_GEMV_BATCHM_DECODE_MAX_M: 2026-09-25: The widest batch that the MTP row dispatch (mtp_head/row_dispatch.rs) and the BF16 arm of the model’s lm_head_batched (impl_a3_lm_head.rs) hand to dense_gemv_bf16_batchm; target_defaults::BASELINE_BATCHM_MAX is set from it. Above it those paths choose other kernels, some of which accumulate in a different order, so this value decides which bits a decode of a given width produces.
DENSE_GEMV_BATCHM_MAX_M: 2026-09-25: Rows dense_gemv_bf16_batchm computes: MAX_M in kernels/gb10/common/dense_gemv_bf16_batchm.cu. The kernel clamps a larger M, so dense_gemv_batchm refuses above this value. Each row’s accumulation does not depend on M, so a row’s result is the same at every batch width.
E4M3_MAX: 2026-09-29: E4M3’s largest finite magnitude; an unscaled cast saturates above it.
FEED_HOST_BIT: 2026-09-25: Bit of a feed_resolve source word that marks an inline host token id in the low 31 bits; without it the word indexes the previous step’s feed cell (argmax_feed.cu).
FP8_GROUPED_DOWN_COLS_PER_CTA: 2026-09-25: Output columns per down CTA. Must equal DOWN_CTA_COLS in the .cu: 2026-09-27, four column groups of 32.
FP8_GROUPED_DOWN_ROWS_PER_PASS: 2026-09-26: Rows per down pass. Must equal GROUP_ROWS in the .cu.
FP8_GROUPED_DOWN_SCALAR: 2026-09-28: moe_expert_down_act_fp8_grouped (reads FP32 SiLU products).
FP8_GROUPED_DOWN_TC: 2026-09-28: moe_expert_down_act_fp8_grouped_tc (reads FP32 SiLU products as BF16 hi + lo).
FP8_GROUPED_DOWN_TC_W8A8: 2026-09-28: moe_expert_down_act_fp8_grouped_tc_w8a8 (TC8_DOWN_COLS).
FP8_GROUPED_GATE_UP_COLS_PER_CTA: 2026-09-26: Output columns per gate+up CTA. Must equal GU_COLS_PER_CTA in the .cu.
FP8_GROUPED_GATE_UP_ROWS_PER_PASS: 2026-09-26: Rows per gate+up pass. Must equal GU_GROUP_ROWS in the .cu: the shared expert takes ceil(num_tokens / this) block rows.
FP8_GROUPED_GATE_UP_SCALAR: 2026-09-28: moe_expert_gate_up_act_fp8_grouped (FP32 SiLU products).
FP8_GROUPED_GATE_UP_TC: 2026-09-28: moe_expert_gate_up_act_fp8_grouped_tc (FP32 SiLU products).
FP8_GROUPED_GATE_UP_TC_W8A8: 2026-09-28: moe_expert_gate_up_act_fp8_grouped_tc_w8a8 (TC8_GU_COLS: one 128-column quantization group of the down projection’s K per CTA).
FP8_GROUPED_SORT_MAX_EXPERTS: 2026-09-27: Experts moe_fp8_grouped_sort handles (PMS_MAX_EXPERTS).
FP8_GROUPED_TC_DOWN_COLS_PER_CTA: 2026-09-28: Output columns per tensor-core down CTA (TC_DOWN_COLS).
FP8_GROUPED_TC_GATE_UP_COLS_PER_CTA: 2026-09-28: Output columns per tensor-core gate+up CTA (TC_GU_COLS).
FP8_GROUPED_TC_ROWS_PER_PASS: 2026-09-28: Rows per pass of the tensor-core kernels (moe_fp8_grouped_tc.cu, TC_ROWS), gate+up and down alike.
FP8_QUANT_ENTRY: 2026-09-25: Entry point of the shared quantizer.
FP8_QUANT_HOPPER_ENTRY: 2026-09-25: Entry point of the Hopper twin. It differs from the shared entry point, so a Hopper image carries both kernels.
FP8_QUANT_HOPPER_GROUPS_PER_CTA: 2026-09-25: K-groups one Hopper CTA covers: 128 threads / 16 threads per group, as HQ_GROUPS_PER_CTA in the .cu. The kernel derives its group span from gridDim.y, so this value sets only the grid’s Y extent.
FP8_QUANT_HOPPER_MODULE: 2026-09-25: Module (file stem) of the Hopper twin.
FP8_QUANT_LOG_SLOTS: 2026-09-25: Once-flags fp8_quant_log keeps: one per refusal reason, plus the twin’s.
FP8_QUANT_MIN_CTAS_PER_SM: 2026-09-25: CTAs per SM the twin’s grid must reach before it takes a launch. The same value as ssm_ba_gates_hopper::MIN_CTAS_PER_SM.
FP8_QUANT_MODULE: 2026-09-25: Module of the shared quantizer; the file and the kernel share this name.
FP8_QUANT_REJECTS: 2026-09-25: Every reason fp8_act_quant_hopper_reject can return. fp8_quant_log keeps one once-flag per entry, by index.
FP8_QUANT_TOO_FEW_CTAS: 2026-09-25: The floor’s refusal reason.
GDN_CARRY_CAP: 2026-09-26: Rows the carry stash holds per (layer, slot): CARRY_CAP in gated_delta_rule_carry.cu. A verify whose K rows might not fit behind the pending ones (pending + K > GDN_CARRY_CAP) writes the state back and restarts the stash.
GDN_CARRY_EAGER_MAX_PENDING: 2026-09-26: Most pending rows an eager verify kernel folds; a slot holding more (only after lazy verifies) is folded by gdn_carry_flush first.
GDN_CARRY_LAZY_MIN_SEQS: 2026-09-26: Launch width (sequences) from which a run takes the lazy verify kernels, which write the state back only when the stash is full. Measured on dgx3 (k=1, –mtp-gate force): from 16 the lazy form is 2.5 % faster end to end at C=16; a threshold of 8 instead of 16 gives C=8 214.4 -> 216.2 tok/s (three reps each) with C=4 unchanged; at 2 and 4 sequences the eager form is ahead.
GDN_DECODE_STRIDED_SMEM_CAP_BYTES: 2026-09-29: Dynamic shared memory each gdn_decode_f32_strided CTA reserves and does not use, from GDN_DECODE_STRIDED_SMEM_CAP_MIN_BATCH sequences: it caps the CTAs resident per SM (3 on GB10’s 100 KB), so a CTA’s 64 KB head state is still in L2 when its update pass re-reads it; uncapped, the second pass of a wide launch misses L2. Measured 2026-09-29 on dgx2 (gdn_decode_strided_microtest, per layer, uncapped / capped / per-sequence launches), 32 value heads: B=16 365 / 362 / 428 us, B=64 1904 / 1362 / 1740, B=128 3857 / 2641 / 3541; 48 value heads: B=16 672 / 507 / 589, B=128 5737 / 3902 / 4766. The output and state bits do not depend on it.
GDN_DECODE_STRIDED_SMEM_CAP_MIN_BATCH: 2026-09-29: Smallest batch that takes the cap. Below it the cap costs more than the L2 it saves (32 value heads, B=4: 34.1 us uncapped, 43.7 us capped).
GDN_SCALAR_SPINE_PIPE: 2026-09-25: The scalar spine entries init_kernels::fused_spine_kernel binds, as gdn_scalar_spine picks them. The init route line and the handle are built from the same strings. 2026-09-28: The default: _vfused’s arithmetic with W, K and U double-buffered through cp.async (bit-identical to _vfused, 1.4x faster on GB10 at 1.6k-8k tokens per chunk).
GDN_SCALAR_SPINE_VFUSED: 2026-09-25: SPLIT=2, 256 threads, single-buffered; METRALE_GDN_PIPE=0.
GDN_SCALAR_SPINE_VTILE: 2026-09-25: SPLIT=4, 512 threads. Not the default; fused_spine_kernel’s doc says why.
GDN_TC_SPINE_ENTRY: 2026-09-25: The spine entry the gdn_prefill_tc lever launches: two bf16 limbs of S_c. Both entries take the same arguments, grid, block and shared memory, so nothing downstream depends on which one is bound.
GDN_TC_SPINE_MODULE: 2026-09-25: The module the entry lives in: kernels/hopper/common/gated_delta_rule_chunk_tc.cu.
KQUANT_MODULE
MAX_DYNAMIC_SMEM: 2026-09-25: Largest dynamic shared memory, in bytes, one block may opt into on the GB10 target (sm_121).
MIN_CTAS_PER_SM: 2026-09-25: Tokens per SM the twin requires before it takes a launch. It is a judgement, not a contract: set high, it only sends more launches to the parent, which computes the same bits.
MOE_E4M3_DN: 2026-09-27: moe_w8a8_grouped_gemm_e4m3_dn (down: N = hidden, K = inter).
MOE_E4M3_GU: 2026-09-27: moe_w8a8_grouped_gemm_e4m3_gu (gate/up: N = inter, K = hidden).
MOE_ROUTED_DOWN_KNOWN: 2026-09-30: The routed experts’ down projection reads the SiLU product. On Qwen3.6-35B-A3B-NVFP4 it reaches 510 at position 0 of the last layer (measured 2026-09-29, no effect on the output); every other certified model stays under 448. The scaled cast that removes it is follow-up work, so this covers exactly the known value.
MOE_ROUTER_GEMM_COLS: 2026-09-26: Output columns per block of moe_router_gemm_bf16 (RC_COLS in kernels/gb10/common/moe_router_gemm.cu).
MOE_ROUTER_GEMM_MAX_PREFILL_ROWS: 2026-10-01: Most rows for which the MoE prefill router takes moe_router_gemm (the decode routers’ kernel, the same bits as dense_gemm_router); the decode routers run it at 2 to 64 rows.
MOE_ROUTER_GEMM_ROWS: 2026-09-26: Activation rows per block of moe_router_gemm_bf16 (RC_ROWS).
MOE_ROUTER_RT_MIN_ROWS: 2026-09-28: Fewest rows for which moe_router_gemm_rt beats dense_gemm_router on GB10 (measured: 0.82x at 300 rows, 2.5x at 1000, 3.0-3.5x at 4k-32k).
MOE_TOPK_SIGMOID_MAX_EXPERTS: 2026-09-25: Largest num_experts the sigmoid routing kernel holds, from #define MAX_EXPERTS in the same file. Beyond it the kernel considers only the first MAX_EXPERTS experts (actual_n is a min).
MOE_TOPK_SIGMOID_MAX_TOP_K: 2026-09-25: Largest top_k the sigmoid routing kernel holds: #define MAX_TOP_K in kernels/gb10/common/moe_topk_sigmoid.cu, which sizes its shared s_top_vals / s_top_idxs. The kernel clamps a larger top_k to it, and NemotronMoeLayer construction refuses one (nemotron_moe.rs). crates/model-engine/tests/moe_topk_sigmoid_bounds.rs pins this mirror to the define and fails if a directory other than common/ holds a copy of the kernel.
NEMOTRON_SHARED_EXPERT_E4M3: 2026-09-30: --nemotron-shared-expert-e4m3: the Nemotron-H shared expert’s prefill tile GEMM casts its activations to E4M3 unscaled. The flag’s help discloses the saturation (1-30% local error measured on Nemotron-3-Nano), so any magnitude is covered.
NVFP4_BLOCK_BYTES
NVFP4_GROUPED_DOWN_COLS_PER_CTA: 2026-09-27: Output columns per down CTA. Must equal NG_DOWN_COLS_PER_CTA in the .cu.
NVFP4_GROUPED_DOWN_ROWS_PER_PASS: 2026-09-27: Rows per down pass. Must equal NG_DOWN_ROWS in the .cu.
NVFP4_GROUPED_DOWN_SCALAR: 2026-10-02: moe_expert_down_act_nvfp4_grouped (reads FP32 SiLU products).
NVFP4_GROUPED_DOWN_TC: 2026-10-02: moe_expert_down_act_nvfp4_grouped_tc (NTC_DOWN_COLS columns per CTA; reads the hi + lo SiLU rows).
NVFP4_GROUPED_GATE_UP_COLS_PER_CTA: 2026-09-27: Output columns per gate+up CTA. Must equal NG_GU_COLS_PER_CTA in the .cu.
NVFP4_GROUPED_GATE_UP_ROWS_PER_PASS: 2026-09-27: Rows per gate+up pass. Must equal NG_GU_ROWS in the .cu.
NVFP4_GROUPED_GATE_UP_SCALAR: 2026-10-02: moe_expert_gate_up_act_nvfp4_grouped (CUDA cores, FP32 SiLU products).
NVFP4_GROUPED_GATE_UP_TC: 2026-10-02: moe_expert_gate_up_act_nvfp4_grouped_tc (NTC_GU_COLS columns per CTA; writes the SiLU product as BF16 hi + lo).
NVFP4_GROUPED_TC_ROWS_PER_PASS: 2026-10-02: Rows per pass of the tensor-core kernels (moe_nvfp4_grouped_tc.cu, TC_ROWS), gate+up and down alike.
NVFP4_LEAN_REPACK_SMEM_PER_K: 2026-10-04: Shared bytes nvfp4_tc_lean_repack stages per unit of K: a 16-row tile’s packed (16 * K / 2) and scale (16 * K / 16) bytes.
NVFP4_MMQ_SMEM: 2026-09-25: Dynamic shared memory of the 128-wide M tile.
NVFP4_PIPE_SMEM: 2026-09-28: Dynamic shared memory of metrale_nvfp4_gemm_pipe: two stages of 128 weight rows and 192 activation rows, 144 bytes each.
Q2K_BLOCK_BYTES: 2026-09-25: Bytes per 256-value super-block: block_q2_K, block_q3_K, block_q6_K.
Q2K_MMQ_SMEM: 2026-09-25: Dynamic shared memory of the 128x128 MMQ tile on the MMA path, as the vendored mmq_get_nbytes_shared computes it with MMQ_TILE_NE_K = 32: ids 512 B + tile_x (128 rows x tile_x_k ints) + tile_y (128 x 144 B). MMQ_MMA_TILE_X_K_Q2_K = 232 + 32 + 4 = 100; MMQ_MMA_TILE_X_K_Q3_K = 232 + 16 + 4 = 84.
Q2_0_BLOCK_BYTES: 2026-09-25: sizeof(block_q2_0): an fp16 scale d and 128 codes at 4 per byte.
Q2_BATCHM_MAX_M: 2026-09-25: Row cap of the q2_0_gemv_vec_batchm kernel (MAX_M in the .cu). The kernel stages M activation rows into s_A[MAX_M * TILE_K], so M > MAX_M writes past that tile, and its compute and store loops stop at MAX_M, so rows from 8 on are never written. q2_0_gemv_vec_batchm chunks larger m itself.
Q3K_BLOCK_BYTES
Q3K_MMQ_SMEM
Q4K_BLOCK_BYTES: 2026-09-25: sizeof(block_q4_K).
Q4K_MMQ_SMEM: 2026-09-25: Dynamic shared memory of the Q4_K MMQ kernel at mmq_x = mmq_y = 128: mmq_get_nbytes_shared in q4k_vendor/mmq.cuh with an x tile of 76 ints per row. Above 48 KiB, so the registry raises the kernel’s dynamic shared memory limit (registry.rs).
Q6K_BLOCK_BYTES
Q8_1_BLOCK_BYTES: 2026-09-25: block_q8_1: 32 int8 + half2 (d, sum).
QK2_0: 2026-09-25: Weights per block_q2_0.
QK_K: 2026-09-25: Weights per block_q4_K.
QK_NVFP4: 2026-09-25: Weights per repacked NVFP4 block: 36 bytes, 32 of E2M1 nibbles and 4 UE4M3 scales.
SSD_L: 2026-09-25: SSD chunk length. Must equal SSD_L in mamba2_ssd_chunk.cu.
SSD_PT: 2026-09-25: head_dim rows per SSD scan block. Must equal SSD_PT in mamba2_ssd_chunk.cu.
W4A16_GEMV_OUTS_PER_BLOCK: 2026-09-25: Outputs per 256-thread block of w4a16_gemv; equals #define N_PER_BLOCK in kernels/gb10/common/w4a16_gemv.cu.
W4A16_GEMV_SW_OUTS_PER_BLOCK: 2026-09-25: Outputs per 256-thread block of the single-warp kernels (w4a16_gemv_sw and the _sw kernels in w4a16_gemv_fused.cu); equals #define N_PER_BLOCK_SW.
W4A16_TC_ROWS_COLS: 2026-10-02: Output columns per CTA (TR_COLS).
W4A16_TC_ROWS_MAX_M: 2026-10-02: Widest row count of one launch (8 * NT of w4a16_tc_rows_64).
W4A16_TC_ROWS_MODULE: 2026-10-02: The kernel module.
W4A16_VIA_FP8_MIN_M: 2026-09-28: Rows at or above which w4a16_t_via_fp8_ldmab beats w4a16_gemm_t_m128 on the dense 27B attention projections (dgx2: x1.23-1.35 at 1024 rows, x0.69 at 300, x2.15 at 8192): below it the per-call weight dequant dominates.
W8A8_K_UNIT: 2026-09-28: The K width of one scale group and of one kernel K unit.
W8A8_LAUNCH_ROWS: 2026-09-28: The most rows one GEMV launch serves (16 token tiles of 8).
W8A8_MAX_ROWS: 2026-09-28: The most rows one W8A8 projection serves: 128-row launches, the weight read once per launch (a C128 verify step has 256 rows).
W8A16_GEMM_M16_N_TILE: 2026-09-25: N columns one CTA owns on the default instantiation. The launch grid and m16_tc_kernel (dense_ffn_m16_tc.rs) both use it; the kernel’s M16_N_TILE must equal it.
W8A16_GEMM_M16_N_TILE_WIDE: 2026-09-25: The wide instantiation’s N tile (w8a16_gemm_m16_n64, kernel M16_N_TILE_WIDE), selected for the dense FFN by METRALE_FFN_M16_TC_NTILE=64. It halves the CTA count for a given N.
W8A16_M16_MAX_K: 2026-09-27: Largest K of the skinny body (PM16_MAX_K): one 1152-byte shared slot per 128-K step within 48 KiB.
W8A16_M16_MAX_ROWS: 2026-09-27: Rows the skinny body of the same kernel takes (PM16_MAX_M), when K is at most W8A16_M16_MAX_K. It gives each 8-column block the canonical tile’s per-output arithmetic with the 128-K steps’ products computed in parallel by 8 warps and added in step order.
W8A16_M32_K_BLOCK: 2026-09-25: K granularity: one FP8 scale block per K-step (PM32_K_STEP == PM32_FP8_BLOCK), which is why K % 128 == 0 is required rather than padded.
W8A16_M32_TILE_ROWS: 2026-09-25: Rows one M tile covers (PM32_M_TILE). Above it the kernel still runs, with grid.y = ceil(M / 32) and one weight pass per tile; w8a16_pipelined_prefers_m32 sends contiguous callers to the 128-tile kernel past this bound.
W8A16_M64_TILE_ROWS: 2026-09-26: Rows of the w8a16_gemm_pipelined_m64 tile.
W8A16_TC_ROWS_COLS: 2026-09-28: Output columns per CTA (TR_COLS).
W8A16_TC_ROWS_MAX_M: 2026-09-28: Widest row count of one row chunk (8 * NT of w8a16_tc_rows_64); more rows run as chunks of it (w8a16_tc_rows_64c).

Functions§

allow_e4m3_saturation: 2026-09-30: While the returned scope lives, a saturation on this thread up to saturation.bound is logged once and counted instead of panicking. Scopes nest; the innermost one applies.
argmax_bf16: 2026-09-25: Index of the largest of vocab_size BF16 logits, written as one u32 to out, by a single block.
argmax_bf16_batch: 2026-09-25: argmax_bf16 for n_rows rows row_stride apart, one block per row. Each block runs the per-row body of argmax_bf16 (argmax_bf16.cu), so a row’s index equals what argmax_bf16 returns for it.
argmax_bf16_batch_feed: 2026-09-25: Batched argmax into the feed cells (device u32 per row) with a per-row pair of masked ids masks ([m0, m1] per row, u32::MAX = none): when the plain argmax is a masked id, the row gets the argmax with both ids excluded. argmax_feed.cu states the tie rules. One block per row.
argmax_bf16_batch_lp: 2026-09-25: argmax_bf16_batch that also writes each row’s top-1 log-probability, out_logprob[row] = log softmax(row)[argmax] in FP32, from an online softmax in the same pass over the row. D-Cut ranks verification depths by prefix sums of these values.
ba_gates_cross_warp_pair: 2026-09-25: The two warp partials summed for output local_out, in order: the parent’s smem[local_out * 2] + smem[local_out * 2 + 1].
ba_gates_lane_kv: 2026-09-25: The kv indices lane lane accumulates, in order: the parent’s for (kv = lane; kv < K_VEC; kv += threads_per_out). The order is the reduction order.
ba_gates_log: 2026-09-25: Log which kernel runs and, when the lever asked for the twin and a guard refused, which guard, once per process per slot. The call runs once per layer per step and its verdict depends on m, so each outcome keeps its own flag and names the m it was first reached at.
ba_gates_log_slot: 2026-09-25: The slot a verdict belongs to, or None when the lever was not requested. Pure, so tests can replay a call sequence on the CPU.
ba_gates_min_tokens: 2026-09-25: The smallest token count the twin accepts with sm_count SMs (0 is treated as 1).
ba_gates_output: 2026-09-25: The BA output a thread owns: group * BA_GATES_OUTS + local_out. The parent’s group is blockIdx.x; the twin’s is g0 + g, walking 0..ceil(N/4) in tiles of BA_GATES_GROUPS.
ba_gates_pick: 2026-09-25: Choose between the gb10 parent and the Hopper twin. Its one caller is ssm_preproc::dense_gemm_ba_gates_prefill, which the five BA-gate call sites share, so the rule exists once.
ba_gates_slot: 2026-09-25: Where output n lands in the [gate(nv), beta(nv)] row: Ok(vh) for a gate (alpha) element, Err(vh) for a beta element, split as the parent splits them on within_group < vheads_per_group.
ba_gates_sm_count: 2026-09-25: The device’s SM count, or BA_GATES_FALLBACK_SM_COUNT when the query fails; at least 1.
ba_gates_warp: 2026-09-25: The block-wide warp that holds lane lane of output local_out. The parent writes its warp partial to smem[local_out * 2 + lane / 32] and the twin to red[g * BA_GATES_WARPS + threadIdx.x / 32]; they match because threadIdx.x / 32 == local_out * 2 + lane / 32.
batched_chunk_zero_admitted: 2026-09-28: The pure rule behind prefill_batched_chunk_zero_admitted.
batched_embed: 2026-09-25: Gather the BF16 embedding rows of num_tokens token ids (token_ids_dev, device [num_tokens] u32) into output, one block per token.
batched_embed_fp8: 2026-09-25: batched_embed over an FP8 E4M3 table with one f32 scale per row (the quantize_bf16_to_fp8 layout); rows are dequantized to BF16 on read.
bf16_absmax: 2026-09-25: Max absolute value of n_elems BF16 values, folded into the f32 at out_max with an atomic max. The caller zeroes out_max first; the FP8 KV calibration (fp8_calibration.rs) does so before each K/V pair.
bf16_concat: 2026-09-25: BF16 concatenation (kernel bf16_concat): out[0..N] = a[0..N], out[N..2N] = b[0..N].
bf16_grouped_tc_shape_ok: 2026-10-02: Whether the BF16 tensor-core kernels take an n-column, k-deep projection: k in whole load groups (two 32-K chunks; 128 for margin), n in whole CTAs of geometry.
bf16_to_fp8: 2026-09-25: Cast total_elements BF16 values to FP8 E4M3, two per thread. total_elements must be even: the kernel reads and writes pairs.
check_e4m3_activation_range: 2026-09-29: When kernel casts its A operand to E4M3 with no scale, check the [m, k] BF16 input it reads (see check_e4m3_range). A no-op otherwise, and in release builds.
check_e4m3_grouped: 2026-09-30: The grouped (pointer-table) MoE form of check_e4m3_activation_range. The kernel reads A row sorted_token_ids[p], or row p when sorted_token_ids is null, for every position p in expert_offsets[e]..expert_offsets[e + 1], e < num_experts; only those rows are checked, each k BF16 values wide. A no-op for a kernel that keeps A in BF16, in release builds, and under graph capture.
check_e4m3_range: 2026-09-29: Debug builds: copy the elements BF16 values at input to the host on stream and panic when one exceeds E4M3_MAX in magnitude, unless an allow_e4m3_saturation scope covers it. NaN is not checked here. Skipped under graph capture, where a copy to the host cannot run. Release builds: a no-op.
compute_gdn_gates: 2026-09-25: GDN gates from the interleaved BA projection and the learned A_log / dt_bias: FP32 gate (decay) and beta (write gate) per value head and token.
conv1d_fwd: 2026-09-25: Causal depthwise conv1d followed by SiLU over a full sequence (kernel causal_conv1d_fwd), launched without a bias. Input [batch, dim, seq_len] BF16 (channel-first), weight [dim, d_conv] BF16, output [batch, dim, seq_len] BF16.
conv1d_update: 2026-09-25: Causal conv1d update + SiLU for one decode token per sequence (causal_conv1d_update). conv_state, input and output are contiguous per sequence: [batch, d_inner, d_conv], [batch, d_inner], [batch, d_inner].
conv1d_update_chunk2: 2026-09-25: Conv1d update + SiLU for two tokens per sequence in one launch (causal_conv1d_update_chunk2), one thread per channel. The conv state after the first token goes to conv_state_intermediate, for rollback.
conv1d_update_l2norm: 2026-09-25: Conv1d update + SiLU, then an L2 norm per head_dim group over the Q and K channels (0..qk_channels), in one kernel (causal_conv1d_update_l2norm). The V channels get SiLU only.
conv1d_update_l2norm_strided: 2026-09-25: Conv1d update + SiLU + L2 norm for batch_size decode sequences in one launch, with FP32 output (causal_conv1d_update_l2norm_f32_strided). The input and output row strides are separate arguments because the concurrent-decode input rows are the QKVZ projection’s, qkvz_size apart, while the output rows are d_inner apart; a kernel that used d_inner for both would read sequence b >= 1 from the wrong row.
conv1d_update_prefill: 2026-09-25: Conv1d update + SiLU over seq_len prefill tokens, with separate per-token input and output strides. Runs the token-parallel kernel (conv1d_prefill_tp_k) when its handle is non-zero and METRALE_CONV1D_TP is not 0, else causal_conv1d_update_prefill, one thread per channel walking the tokens.
cublas_bf16_proj_dense: 2026-09-25: Route a projection out[M,N] = act[M,K] @ weightᵀ through cuBLASLt BF16 for a weight that is already BF16 [N,K].
cublas_bf16_proj_dense_f32_out: 2026-09-25: cublas_bf16_proj_dense writing an FP32 output buffer, for consumers whose next kernel reads FP32.
cublas_fp8_m_pad: 2026-09-25: The M extent cublas_fp8_proj_prequant hands cuBLASLt; callers bound their output buffers against it.
cublas_fp8_proj: 2026-09-25: Quantize the activation with per_token_group_quant_fp8 (FP8 plus one scale per token and 128 of K), then run cublas_fp8_proj_prequant. The FP8 weight and its 128x128 block scales go to cuBLASLt as they are, with no dequant.
cublas_fp8_proj_prequant: 2026-09-25: cublas_fp8_proj for an activation the caller already ran through per_token_group_quant_fp8, so one quantization can feed several projections (the dense FFN’s gate and up share one).
cublas_fp8_rowwise_proj: 2026-09-25: Route a projection through row-wise FP8 cuBLASLt. The weight pair comes from requant_weight_rowwise_fp8_cached; the activation is quantized per token on every call.
cublas_scale_layout_kmajor: 2026-09-25: Which VEC128 activation-scale layout the cuBLASLt block-scaled arm passes, from METRALE_CUBLAS_SCALE_LAYOUT, read once per process.
cutlass_bf16_proj: 2026-09-25: Route a projection out[M,N] = act[M,K] @ weightᵀ through CUTLASS BF16, against the FP8 weight’s memoized BF16 copy.
cutlass_nvfp4_proj: 2026-09-25: Route a projection out[M,N] = act[M,K] @ weightᵀ through CUTLASS NVFP4. weight_t is an NVFP4 weight with [K/2,N] packed bytes; the bytes are transposed once to [N,K/2] and memoized.
cutlass_nvfp4_proj_from_fp8: 2026-09-25: CUTLASS NVFP4 projection for an FP8 weight. On first use the weight is dequantized into a transient BF16 buffer, packed into NVFP4 data and scales, and the pair is memoized for later calls.
decode_w8a8_gemm: 2026-09-25: out[m_pad, n] = act_fp8[m_pad, k] @ weight[n, k]ᵀ at row pitch plan.ldc, with both sets of scales applied. The activation must already be through decode_w8a8_quant_act.
decode_w8a8_quant_act: 2026-09-25: Quantize BF16 act[rows, k] into the scratch: FP8 E4M3 bytes and one FP32 scale per token and 128 of K, with the pad rows rows..ceil16(rows) zeroed, and the scales transposed for cuBLASLt when the layout is K-major.
decode_w8a8_selected: 2026-09-25: Whether one decode projection takes the W8A8 cuBLASLt route; DecodeW8a8Plan lists the conditions.
deinterleave_qg: 2026-09-25: Deinterleave Q/gate from per-head interleaved to contiguous layout, in place, per token. Rows are stride elements apart.
deinterleave_qg_split: 2026-09-25: deinterleave_qg with Q written to q_out (contiguous [num_tokens, num_heads * head_dim]) instead of in place. The gate is written back into data at offset num_heads * head_dim of each row.
deinterleave_qg_split_qnorm: 2026-09-25: deinterleave_qg_split fused with the per-head Q RMS norm.
deinterleave_qg_split_qnorm_mrope: 2026-09-25: deinterleave_qg_split_qnorm plus rotate-half MRoPE on the first rotary_dim dimensions of each Q head, using the per-token positions pos_t / pos_h / pos_w.
deinterleave_qkvz: 2026-09-25: Deinterleave the QKVZ projection output from GQA-grouped to sequential layout, per token.
dense_gemm: 2026-09-25: Scalar BF16 GEMM, C = A @ B^T (kernel dense_gemm_bf16(A, B, C, M, N, K)): A [M, K] row-major, B [N, K] row-major (HuggingFace layout), C [M, N] row-major.
dense_gemm_ba_gates_prefill: 2026-09-25: BA GEMM and the GDN gate transforms for m tokens in one kernel, with no intermediate BA buffer.
dense_gemm_ba_gates_prefill_hopper: 2026-09-25: Launch dense_gemm_ba_gates_prefill_hopper, one block per token. It takes the parent’s arguments in the parent’s order; only the grid differs.
dense_gemm_bf16_pipelined: 2026-09-25: Tensor-core BF16 GEMM (kernel dense_gemm_bf16_pipelined): m16n8k16 MMAs fed by a 2-stage cp.async pipeline on a 128x128 tile, with the inputs and output layout of dense_gemm.
dense_gemm_m16_bf16: 2026-09-25: The default 32-wide CTA. input is [m, a_row_stride] BF16 with k used, weight.weight is the [n, k] BF16 weight as loaded, and output is [m, c_row_stride] BF16 with n used.
dense_gemm_m16_bf16_n64: 2026-09-25: The N_TILE = 64 instantiation of the same kernel template: the same arguments as dense_gemm_m16_bf16, ceil(n / 64) CTAs instead of ceil(n / 32).
dense_gemm_prefill: 2026-09-25: Dense BF16 prefill GEMM: dense_gemm_bf16_pipelined when pipelined_kernel is loaded, otherwise dense_gemm with fallback_kernel.
dense_gemm_router: 2026-09-25: Register-blocked BF16 GEMM (kernel dense_gemm_bf16_router) that keeps the scalar dense_gemm_bf16’s per-output FP32 accumulation order, strict k = 0..K-1, and with it the scalar kernel’s results under the --fmad=false build (kernels/gb10/common/KERNEL.toml). router_gate_gemm_dense uses it for the MoE router gate GEMM.
dense_gemm_splitk: 2026-09-25: Split-K BF16 GEMM: partial_kernel writes FP32 partial products for k_splits K chunks into workspace, and reduce_kernel sums them into BF16 output. workspace holds k_splits * m * n FP32 values.
dense_gemm_tc: 2026-09-25: Tensor-core BF16 GEMM, C = A @ B^T, with m16n8k16 MMAs (kernel dense_gemm_tc).
dense_gemm_tc_scaled_acc: 2026-09-25: output[m, n] += scale * bf16(input[m, k] @ weight[n, k]^T) in one pass: dense_gemm_tc with an accumulating epilogue, used for the LoRA expand+fold (lora_delta.rs). It saves the [m, n] scratch that a GEMM followed by bf16_scaled_add writes and reads back.
dense_gemv: 2026-09-25: Dense BF16 GEMV for one row, C = A @ B^T: A [1, K], B [N, K], C [1, N]. Each 256-thread block computes 4 outputs, 64 threads per output (kernel dense_gemv_bf16(A, B, C, N, K)).
dense_gemv_ba_gates: 2026-09-25: BA projection GEMV and the GDN gate/beta transforms in one kernel, for one token, with no intermediate BA buffer.
dense_gemv_batch2: 2026-09-25: Dense BF16 GEMV for two rows in one pass over the weight. Each row follows dense_gemv_bf16’s K order and reduction (gb10 kernels).
dense_gemv_batchm: 2026-09-25: Dense BF16 GEMV for M rows, C[t] = A[t] @ B^T for t in [0, M), reading the weight once for all rows. Each row follows dense_gemv_bf16’s K order and reduction (gb10 kernels).
dense_gemv_batchm_split: 2026-09-27: dense_gemv_batchm over m rows split into y_blocks block rows of ceil(m / y_blocks) rows each (at most DENSE_GEMV_BATCHM_MAX_M): more blocks for a narrow weight such as the MoE router, and one launch past 16 rows. Each row’s result is the same as in any other split.
dense_gemv_fp8w: 2026-09-25: FP8-weight GEMV for one row, C = A @ (dequant(B_fp8) * row_scale)^T: A [1, K] BF16, B [N, K] FP8 E4M3, row_scale [N] f32, C [1, N] BF16. One byte per weight, against two for dense_gemv. Each 256-thread block computes 4 outputs, 64 threads per output (kernel dense_gemv_fp8w(A, B, row_scale, C, N, K)).
dense_gemv_fp8w_batch2: 2026-09-25: FP8-weight GEMV for two rows in one pass over the weight: input [2, K] BF16, output [2, N] BF16, scaled by weight.row_scale. Not bit-identical to two dense_gemv_fp8w launches: the kernel adds four products to the accumulator at once, where dense_gemv_fp8w adds each product separately.
dense_mm_bf16: 2026-09-25: C[M, N] = A[M, K] @ B[N, K]^T, BF16 in and out, output row stride N, dispatched on M: dense_gemv_bf16 at M == 1, dense_gemv_bf16_batchm at 2..=DENSE_GEMV_BATCHM_MAX_M, and the tile GEMM otherwise or when the chosen GEMV kernel is missing.
dequant_fp8_bf16_bytes: 2026-09-25: BF16 bytes dequant_fp8_bf16_into writes for fp8w.
dequant_fp8_bf16_into: 2026-09-25: Dequantize a block-scaled or per-row FP8 weight [N,K] to BF16 into dst, a caller-owned buffer of dequant_fp8_bf16_bytes bytes.
dequant_nvfp4_to_bf16: 2026-09-25: Dequantize an NVFP4 weight [n, k] (packed E2M1, E4M3 group scales and the per-tensor scale2) to BF16 [n, k], one block per row.
dequant_q2_0_gn_to_bf16: 2026-09-25: Dequantize a packed Q2_0 weight [N, K] into the caller’s BF16 buffer [N, K] on stream, with no allocation and no host sync (kernel dequant_q2_0_gn_to_bf16 in kernels/gb10/common/dequant_gguf_bf16.cu). Packed-Q2 prefill uses it to feed a BF16 GEMM while the resident weight stays packed (dense_ffn.rs, qwen3_attention/prefill_weights.rs).
e4m3_saturations: 2026-09-30: Saturating launches counted under site in this process (debug builds only).
embed_from_argmax: 2026-09-25: Read the token id at argmax_out, copy its BF16 row of embed_table to embed_out, and copy the id to token_id_out, all on the device.
feed_resolve: 2026-09-25: ids_out[i] is the inline id when sources[i] has FEED_HOST_BIT set, else cells[sources[i]].
fill_slots_from_block_table: 2026-09-25: Fill count slot-mapping entries on the device from a block table: slots[i] = block_table[(start_pos + i) / block_size] * block_size + (start_pos + i) % block_size. count == 0 launches nothing.
fp4_act_scratch_bytes: 2026-09-25: block_fp4_mmq activation bytes for [m, k], with k rounded up to 256, plus 1 MiB of slack. It never exceeds q8_1_scratch_bytes(m, k), the size of the ffn_act_q8 scratch it is written into.
fp8_act_quant_hopper_enabled: 2026-09-25: Is the twin requested? The compiled target’s [defaults] fp8_act_quant_hopper, overridden by METRALE_FP8_ACT_QUANT_HOPPER when it is set (super::target_defaults).
fp8_act_quant_hopper_reject: 2026-09-25: Why the twin is not running; None means it runs.
fp8_act_scale_to_kmajor: 2026-09-25: Launch fp8_act_scale_to_kmajor: copy the quantizer’s row-major [M, K/128] FP32 activation scales into [K/128, M_pad], with zeros in the M..M_pad pad rows. a_scale is only read.
fp8_fp8_gemm_m128_mfast: 2026-09-25: FP8 x FP8 GEMM through fp8_fp8_gemm_t_m128_mfast: a 128-row M tile with M on the fast grid axis. A must already be FP8 E4M3 (bf16_to_fp8).
fp8_fp8_gemm_n128: 2026-09-25: FP8 x FP8 GEMM: A [M, K] FP8 E4M3 and B [N, K] FP8 E4M3 give C [M, N] BF16.
fp8_fp8_gemm_n128_m128: 2026-09-25: fp8_fp8_gemm_n128 with a 128-row CTA (two 64-row chunks); each CTA loads a B tile once for both chunks.
fp8_gemm_m128_mfast: 2026-09-25: FP8-weight GEMM through fp8_gemm_t_m128_mfast: a 128-row M tile (two 64-row chunks per CTA) with M on the fast grid axis.
fp8_gemm_n128: 2026-09-25: FP8-weight prefill GEMM, C = A @ B_fp8^T: A [M, K] BF16, B_fp8 [N, K] FP8 E4M3 (from predequant_nvfp4_to_fp8), C [M, N] BF16.
fp8_gemm_n128_m128: 2026-09-25: FP8-weight GEMM with a 128-row CTA (two 64-row chunks); each CTA loads a B tile once for both chunks.
fp8_gemm_n128_mfast: 2026-09-25: FP8-weight GEMM through the fp8_gemm_t_mfast kernel, with M on the fast grid axis (blockIdx.x = M block), so consecutive CTAs share a B panel.
fp8_gemm_n128_row_scaled: 2026-09-25: Row-scaled FP8 GEMM, C[M, N] = A[M, K] @ (dequant(B_fp8[N, K]) * row_scale[N])^T; the scale multiplies each output column before the BF16 store. Takes the Fp8DenseWeight from crate::weight_map::DenseWeight::quantize_to_fp8. The DFlash drafter calls it when its quantization is Fp8Weights (dflash_head/forward_block_layer_paged.rs).
fp8_gemm_n128_row_scaled_m16: 2026-09-25: fp8_gemm_n128_row_scaled through the single-warp kernel fp8_gemm_t_row_scaled_m16, whose M tile is 16 rows. The grid has no M dimension, so m must be at most 16: rows past the tile are not computed.
fp8_gemm_t_blockscaled: 2026-09-25: W8A8 GEMM with per-token activation scales and per-block weight scales, applied in an FP32 epilogue per 128-element K-group:
fp8_gemv_rowscale_batch8_rt2: 2026-09-25: Register-tiled row-scaled FP8 GEMV for up to 8 rows, two outputs per 64-lane group (kernel fp8_gemv_rowscale_batch8_rt2, module fp8_gemv_rt). input [M, K] BF16, output [M, N] BF16; the kernel applies row_scale[n] at write-out. The DFlash drafter’s FP8 path calls it (dflash_head/forward_block.rs). Refuses m outside 1..=8 and a k that is not a multiple of 16.
fp8_gemv_rowscale_batch16_rt2: 2026-09-25: The 16-row instantiation of the same kernel template (fp8_gemv_rowscale_batch16_rt2, module fp8_gemv_rt), with the launch geometry of fp8_gemv_rowscale_batch8_rt2. Refuses m outside 1..=16 and a k that is not a multiple of 16.
fp8_grouped_active_cap: 2026-09-25: Cap on active experts, which sizes the grouped grids’ Y extent: num_tokens * top_k rows can reach at most that many distinct experts. It does not depend on the routing, so a captured graph stays valid for every routing.
fp8_grouped_tc_shape_ok: 2026-09-28: Whether the tensor-core kernels take an n-column, k-deep projection: both in whole 128 blocks (the FP8 scale blocks) and n in whole CTAs of geometry.
fp8_quant_grid: 2026-09-25: Grid for one quantizer launch. Pure, so fp8_act_quant_tests.rs checks both arms without a GPU.
fp8_quant_hopper_ctas: 2026-09-25: CTAs the twin’s grid launches for (m, k): the product of super::fp8_quant_grid’s Hopper arm, so the floor measures the grid the launch uses.
fp8_quant_hopper_lane: 2026-09-25: The group slot thread tid of a Hopper CTA works on, and the half-open range [start, end) of 8 elements it owns within that 128-element group, as sub/lane in the .cu. None when the slot is past the CTA’s span.
fp8_quant_hopper_min_m: 2026-09-25: The smallest m the twin accepts at this k: the table in this module’s header, as a function.
fp8_quant_hopper_span: 2026-09-25: The half-open K-group range one Hopper CTA owns, computed as the .cu computes gpc/g0/g_end. fp8_act_quant_tests.rs uses it to check on the host that the grid’s Y extent and the kernel’s span partition the groups.
fp8_quant_log: 2026-09-25: Log which quantizer runs and, when the lever requested the twin and did not get it, which reason refused it: once per process per slot. per_token_group_quant_fp8 calls this on every launch.
fp8_quant_log_slot: 2026-09-25: The slot a verdict belongs to, or None when it is not logged. Pure, so tests replay a call order on the CPU.
fp8_quant_min_ctas: 2026-09-25: The smallest twin grid, in CTAs, the floor accepts for this SM count; an sm_count of 0 counts as 1.
fused_k_norm_rope_cache_write_bf16: 2026-09-25: K norm (x * rsqrt(mean(x^2) + eps) * (1 + w)), RoPE and the paged BF16 cache write in one kernel. K stays in FP32 from the BF16 load to the single BF16 rounding at the store. One block per (token, KV head) with head_dim threads; the kernel supports head_dim <= 256.
fused_k_norm_rope_cache_write_bf16_mrope: 2026-09-25: The interleaved-MRoPE variant (kernel fused_k_norm_rope_mrope_cache_write_bf16): rotation pair i takes its position from pos_t, pos_h or pos_w by i % 3. The rest of the kernel is the scalar-position one, so equal position streams give its results.
fused_k_norm_rope_cache_write_fp8: 2026-09-25: fused_k_norm_rope_cache_write_bf16 writing an FP8 E4M3 cache. The FP32 result times inv_scale (1 / k_scale) goes through one saturating FP8 cast, with no BF16 rounding.
fused_k_norm_rope_cache_write_fp8_kv: 2026-09-25: Decode-path fusion of rms_norm + rope_forward + reshape_and_cache_flash_fp8 for an FP8 KV cache: writes K (normed, rotated, quantized) and V (quantized) in one launch, leaving rope_forward for Q only.
gated_rms_norm: 2026-09-25: Gated RMS norm of num_tokens rows, one block per row; the gate row of token t starts at t * gate_stride. The order of norm and gate is the resolved kernel’s. gated_rms_norm and gated_rms_norm_f32_input in kernels/gb10/common/rms_norm.cu compute rms_norm(input) * weight * silu(gate) and ignore group_size. The per-group shadows, e.g. kernels/gb10/minimax-m2-229b/nvfp4/rms_norm.cu, gate first and normalise each group_size slice.
gated_rms_norm_prefill: 2026-09-25: Gated RMS norm of every (head, token) row of a prefill chunk in one launch. input and output rows of token t start at t * input_token_stride, and gate rows at t * gate_token_stride, both in BF16 elements.
gated_rms_norm_strided: 2026-09-25: Gated RMS norm of every (head, sequence) row of a multi-sequence decode in one launch: heads_per_seq rows per sequence, num_seqs sequences. Sequence strides count elements of each buffer’s own type: input_seq_stride in f32, gate_seq_stride and output_seq_stride in BF16.
gdn_carry_conv: 2026-09-26: Carried-state conv verify (gdn_carry_conv): the twin of gdn_verify_fused_conv_kn_batched that first shifts each sequence’s pending input rows into its window (writing it back when there were any), writes no snapshot and no final window, and stashes the position inputs at conv_stash + slot * stash_seq_elems. lazy must match the verify kernel of the same run (gdn_carry_wy{K}_lazy): the window is then written back only when the stash is full, else on every launch.
gdn_carry_conv_f32: 2026-10-01: Carried-state FP32 conv verify (gdn_carry_conv_f32): the twin of causal_conv1d_update_l2norm_f32_strided for all num_tokens positions of n_seq sequences, with gdn_carry_conv’s pending-row fold and stash. Sequence b’s window is at conv_state + b * dim * d_conv floats; row b * num_tokens + t reads new_input at input_stride and writes the FP32 output at output_stride. lazy must match the verify kernel of the same run.
gdn_carry_conv_flush: 2026-09-26: Standalone conv fold (gdn_carry_conv_flush) over layers GDN layers, laid out as in gdn_carry_flush with conv-state pointers and a BF16 stash.
gdn_carry_conv_seq_elems: 2026-09-26: Per-(layer, slot) conv stash length in BF16 elements: GDN_CARRY_CAP input rows.
gdn_carry_flush: 2026-09-26: Standalone fold (gdn_carry_flush) over layers GDN layers: layer l reads its h pointers at entry l * table_layer_entries of h_table, its stash at carry_base + l * carry_layer_floats floats and its counts at pend + l * pend_layer_entries. A count of 0 leaves H alone; pend is not cleared.
gdn_carry_seq_floats: 2026-09-26: Per-(layer, slot) stash width in floats: vn[CAP][nv][vd] | g[CAP][nv] | sk[CAP][nv][kd], the layout of the kernel’s CARRY_VN/CARRY_G/CARRY_SK macros.
gdn_carry_wy: 2026-09-26: Carried-state verify (gdn_carry_wy{2,3,4} or its _lazy form, chosen by the caller’s kernel). It applies each sequence’s pending rows (pend[slot_tab[b]]) to its copy of H and writes H back (the lazy form only when this verify’s rows might not fit behind them in the stash, pending + K > GDN_CARRY_CAP), writes the output, and stashes the rows it verified at carry_base + slot * seq_floats. Engaged word b becomes 2 after a write-back, else 1.
gdn_conv_chain_f32: 2026-10-01: The single-sequence exact verify’s FP32 conv (gdn_conv_chain_f32): num_tokens positions of causal_conv1d_update_l2norm_f32_strided for one sequence in one launch, the window after position t (t < num_tokens - 1) at conv_inter[t], the final window at conv_state. The unused conv_inter entries must be valid pointers; they are not written.
gdn_conv_chain_f32_batched: 2026-10-01: super::gdn_conv_chain_f32 for batch sequences on consecutive conv slots (dim * d_conv floats apart): row b * num_tokens + t reads new_input + row * input_stride and writes output + row * output_stride; the window after position t < num_tokens - 1 goes to conv_inter + b * inter_strides[1] + t * inter_strides[0] floats.
gdn_decode: 2026-09-25: Single-token GDN decode (gated_delta_rule_decode). The batch_size h-states are contiguous: [batch, num_v_heads, k_dim, v_dim] FP32.
gdn_decode_chunk2: 2026-09-25: Two tokens through the GDN recurrence in one launch (gated_delta_rule_chunk2). The state after the first token is written to h_state_intermediate, for rollback when the draft is rejected. The strides are in elements, not bytes.
gdn_decode_chunk3: 2026-09-25: Three tokens through the GDN recurrence in one launch (gated_delta_rule_chunk3). The states after the first and second tokens go to h_state_inter0 and h_state_inter1, for rollback when drafts are rejected.
gdn_decode_f16_strided_norm: 2026-09-25: FP16 h-state twin of gdn_decode_f32_strided_norm, used when ssm_h_fp16_enabled() is set. It adds one argument, h_seq_stride: the per-sequence slot pitch of the h-state pool in __half elements. The pitch depends on how the pool was sized, not on the head dims, so the caller passes it.
gdn_decode_f32_conv_norm: 2026-09-25: Conv1d update + L2 norm, the GDN recurrence and the gated RMS norm in one decode launch (gated_delta_rule_decode_f32_conv_norm).
gdn_decode_f32_norm: 2026-09-25: FP32 GDN decode fused with the gated RMS norm: the kernel normalises each v-head’s output in the same CTA and writes BF16, with no FP32 output buffer in global memory.
gdn_decode_f32_norm_snap: 2026-09-25: super::gdn_decode_f32_norm plus an inline h-state snapshot.
gdn_decode_f32_strided: 2026-09-25: Strided FP32 GDN decode for several sequences in one launch: Q/K/V, gate/beta and the output are rows qk_stride, v_stride, gb_stride and out_stride apart per sequence.
gdn_decode_f32_strided_norm: 2026-09-25: gdn_decode_f32_strided fused with the gated RMS norm: it writes the BF16 post-norm output directly, out_stride apart per sequence.
gdn_decode_f32_strided_norm_snap: 2026-09-25: super::gdn_decode_f32_strided_norm plus an inline h-state snapshot, for the batched-verify arm over batch_size sequences.
gdn_decode_strided_smem: 2026-09-29: The dynamic shared memory gdn_decode_f32_strided reserves for batch_size.
gdn_decode_wy2: 2026-09-25: WY-chunkwise 2-token verify decode (gated_delta_rule_wy2). One pass over H computes both tokens’ H^T k products, the WY correction is applied, and a second pass updates H and writes the intermediate state.
gdn_decode_wy3: 2026-09-25: WY-chunkwise 3-token verify decode (gated_delta_rule_wy3): the gdn_decode_wy2 scheme for three tokens, with two intermediate states.
gdn_decode_wy4: 2026-09-25: WY-chunkwise 4-token verify decode (gated_delta_rule_wy4): the gdn_decode_wy2 scheme for four tokens, with three intermediate states.
gdn_decode_wy4_woa: 2026-09-25: Write-on-accept K=4 verify (gated_delta_rule_wy4_woa). It writes the output but no state, and stashes the per-row update terms for gdn_wy4_fold. The state arguments are pointer tables.
gdn_decode_wyn: 2026-09-25: WY-chunkwise verify decode for K tokens, contiguous form: the launch for gated_delta_rule_wy5 .. wy16 (and their _f16 twins) from gated_delta_rule_wyn.cu, and for gated_delta_rule_wy17. K is compile-time in the kernel, so the kernel handle selects it. The kernel writes Hi_0..Hi_{K-2} and the final H.
gdn_decode_wyn_table: 2026-09-25: Cross-sequence batched-verify launch of the wyN pointer-table twins (gated_delta_rule_wy{5..16}_table and their _f16_table forms). state_is_table is compiled into the symbol, so the argument list is the same as gdn_decode_wyn, but the state arguments are tables:
gdn_exact_carry: 2026-10-01: Carried-state exact verify (gdn_exact_carry{2,3,4} or its _lazy form, chosen by the caller’s kernel; gdn_exact_carry.cu in the model directory): the stash protocol of gdn_carry_wy around the strided FP32 decode’s per-token chain. Rows are sequence-major, b * K + t: FP32 q and k at qk_stride, v at v_stride, gate and beta at gb_stride, the FP32 output at out_stride, all in floats.
gdn_exact_chain: 2026-10-01: The single-sequence exact verify (gdn_exact_chain{2,3,4}, chosen by the caller’s kernel): K rows of one sequence’s strided-decode chain in one launch. h_state is updated to the state after the last row, h_inter[t] receives the state after row t (t < K - 1; the other entries must still be valid pointers and are not written). Rows are strides = [qk, v, gate|beta, out] floats apart.
gdn_exact_chain_f16: 2026-10-01: The FP16 h-state exact verify, K = 2..4 rows per sequence (kernel is the gdn_exact_chain_f16_{K} entry): the decode’s FP16 fused-norm chain per row with the state in registers, the FP16 intermediates of rows 0..K-2 and the final state written, and the BF16 normed rows. Rows are b * K + t; strides is [qk, v, gb, z, out] per row.
gdn_init_spine_line: 2026-09-25: GDN state spine: …, the line qwen3_ssm::init prints for each layer as it binds the handles, before any prefill has run.
gdn_prefill: 2026-09-25: Token-sequential GDN prefill (gated_delta_rule_prefill): one block per (v-head, sequence) walks all seq_len tokens. Q/K/V and gate/beta are read at their own per-token strides, so the packed conv output needs no repacking.
gdn_prefill_fla: 2026-09-25: Chunked GDN prefill in three launches, ordered by stream:
gdn_prefill_persistent: 2026-09-25: Persistent GDN prefill (gated_delta_rule_prefill_persistent): each block loads its v-head’s H into shared memory once, runs every token against it, and writes it back to global memory once at the end.
gdn_prefill_persistent_batched: 2026-09-25: Batched persistent prefill with H plus double-buffered k and q in shared memory (gated_delta_rule_prefill_persistent_batched).
gdn_prefill_persistent_smem: 2026-09-25: Persistent GDN prefill whose shared-memory size the caller computes, for kernels with a larger layout (the WY4-persistent and WY64 prefill kernels).
gdn_prefill_persistent_smem_batched: 2026-09-25: Batched persistent prefill whose shared memory the caller computes (gated_delta_rule_prefill_wy64_batched or gated_delta_rule_prefill_persistent_wy4_batched).
gdn_prefill_regresident: 2026-09-25: Token-sequential prefill recurrence with H in registers (gated_delta_rule_prefill_regresident). One warp owns one v-column and holds its 128 k-rows (4 per lane); there is no shared-memory H and no block barrier. Requires k_dim == 128 and v_dim % 4 == 0 (4 warps per block).
gdn_prefill_split: 2026-09-25: gdn_prefill with v_dim split across 2 blocks per v-head (gated_delta_rule_prefill_split).
gdn_prefill_split4: 2026-09-25: gdn_prefill with v_dim split across 4 blocks of 32 threads per v-head (gated_delta_rule_prefill_split4).
gdn_prefill_split4_batched: 2026-09-25: Batched split4 prefill (gated_delta_rule_prefill_split4_batched).
gdn_scalar_spine: 2026-09-28: gdn_scalar_spine_from of this process’s environment, read once.
gdn_scalar_spine_from: 2026-09-28: From the values of METRALE_GDN_PIPE and METRALE_GDN_VTILE: PIPE=1 picks the pipe spine (as before), else VTILE=1 the vtile spine, else PIPE=0 the vfused spine, else the pipe spine.
gdn_tc_spine_route_line: 2026-09-25: GDN state spine: …, the line the prefill logs when the tensor-core spine runs, built from GDN_TC_SPINE_ENTRY. Pure and returning a String, so the tests below can check it.
gdn_verify_fused_conv_k2: 2026-09-25: K=2 MTP-verify conv1d + SiLU + L2 norm: both draft positions in one launch, with the position-0 conv-state snapshot written inline.
gdn_verify_fused_conv_kn: 2026-09-25: K-position verify conv1d + SiLU + L2 norm: all num_tokens draft positions in one launch. Position t’s conv-state snapshot is written to conv_state_inter + t * inter_stride.
gdn_verify_fused_conv_kn_batched: 2026-09-25: gdn_verify_fused_conv_kn for n_seq sequences in one launch, one grid row (blockIdx.y) per sequence.
gdn_verify_fused_conv_kn_f32: 2026-09-25: FP32-output twin of super::gdn_verify_fused_conv_kn: one launch runs conv1d + SiLU + L2 norm for all K verify positions, writes FP32 conv rows, and writes each per-token conv-state snapshot inline. output_stride is in FP32 elements.
gdn_verify_fused_norm_k2: 2026-09-25: K=2 MTP-verify gated RMS norm: both draft positions in one launch. The Z gate is read from the deinterleaved [Q|K|V|Z] buffer at z_offset within each position. The gdn_verify_fused_microtest example checks it against the per-token gated_rms_norm at cos >= 0.99999.
gdn_wy4_flag_clear: 2026-09-25: Clear a layer’s write-on-accept engaged word. The caller issues it at the start of a batched verify that requested write-on-accept, on the stream of the launches that follow.
gdn_wy4_fold: 2026-09-25: Post-verdict fold (gated_delta_rule_wy4_fold). When the engaged word is set, it applies stashed rows 0..na_tab[b] to H, reading and writing H once. Otherwise the parent kernel ran, and a partial accept (na < k_rows) restores Hi(na - 1) from hi_tables.
gemv_sw_from: 2026-09-25: Polarity of the METRALE_NO_GEMV_SW kill switch, given the variable’s value: the single-warp GEMV is on unless the value is exactly "1", so =0 and an empty value leave it on.
grouped_gemm_mla: 2026-09-25: g independent GEMMs C_g[m, n_g] = A_g[m, k_g] @ B_g[n_g, k_g]^T in one launch. A_g sits at column g * k_g of A, B_g at row g * n_g of B, and C_g at column g * n_g of C; one block per (row, group) computes 4 outputs.
hc_expand: 2026-09-25: Broadcast the BF16 hidden state into hc_mult identical FP32 streams: streams[t, i, d] = hidden[t, d].
hc_head: 2026-09-25: Final collapse before the LM head: one learned sigmoid-weighted sum over the hc_mult streams, written as BF16 [T, H].
hc_head_lowrank: 2026-09-25: The model-level mixer (use_combine=False in the reference bench/qwen4_exp/ref/modeling_qwen4_exp.py): the same collapse with no injection vector, on the same path choice as hc_pre_lowrank. It is also the model’s final normalization: the config parser sets final_norm_identity (config/src/parsers/qwen4_exp.rs).
hc_head_site: 2026-09-25: The model-level final collapse before the LM head. In the low-rank family this is also the model’s final normalization: the Qwen3.8-Flash-Next config parser sets final_norm_identity (config/src/parsers/qwen4_exp.rs), and the mixer’s hc_norm normalizes.
hc_post: 2026-09-25: Expand the sublayer output back into hc_mult streams: out[t,j,d] = post[t,j] * block_out[t,d] + sum_i comb[t,i,j] * residual[t,i,d]. out may alias residual: each element’s residual values are read before its outputs are written.
hc_post_lowrank: 2026-09-25: Inject the block output back into every stream: out[t, s*H + d] = residual[t, s*H + d] + block_out[t, d] * inj[t, s]. There is no comb argument: this family scales each stream by one scalar.
hc_post_site: 2026-09-25: Inject the block output back into every stream. out may alias residual.
hc_pre: 2026-09-25: Collapse the hc_mult streams to one BF16 row (an RMS-rescaled mix gives sigmoid pre weights for the weighted sum), and write post_out [T, hc] and the Sinkhorn-normalized comb_out [T, hc, hc] for the matching hc_post.
hc_pre_lowrank: 2026-09-25: Collapse the hc_mult streams to one, and emit the per-stream injection weights the matching hc_post_lowrank needs: streams [T, hc, H] -> y_out [T, H] BF16, inj_out [T, hc] FP32. Fails when w.inject_w is null (a site without an injection weight is the model-level mixer, which uses hc_head_lowrank).
hc_pre_site: 2026-09-25: Collapse the streams to one and emit what the matching hc_post needs: post and comb for Sinkhorn, the injection vector in post_out for low-rank.
int8_gemm_faith2_prefill: 2026-09-25: int8 prefill GEMM: requantize the BF16 activations to int8 with one F32 scale per 32 elements (requant_a_kernel), then multiply them with the int8 weight from requant_w_nvfp4_int8 (faith2_kernel). Both launches go on stream.
kquant_mmq_act_bytes: 2026-09-25: Bytes of the MMQ-layout q8_1 buffer for [m, k] activations: 144 B per 128 values, with m rounded up to the 128-row tile and k to 256.
kquant_mmq_gemm: 2026-09-25: Prefill GEMM on raw blocks through the vendored MMQ engine: out[m][n] (bf16) = A_q8[m][k] x W[n][k]. a_q8 must be in the type’s MMQ q8_1 layout (D2S6 for Q2_K, D4 for Q3_K), and smem is the matching *_MMQ_SMEM. kernel_nc serves n % 128 == 0, and kernel_wc (bounds-checked) any other n. The kernel’s arguments are (x = weights, y = activations, dst, nrows_x = n, ncols_dst = m, ncols_x = k, stride_row_x = k / 256, ncols_y = m, stride_col_dst = n).
kquant_mmvq: 2026-09-25: Decode GEMV on raw blocks: out[m][n] = sum_k xq8[m][k] * W[n][k] (kquant_mmvq_q2_k / kquant_mmvq_q3_k), one 4-warp block per output row. w_blocks is [n][k/256] raw super-blocks; y_q8 is the block_q8_1 [m][k/32] buffer; out_bf16 is [m][n].
kquant_mmvq_experts: 2026-09-25: kquant_mmvq_*_experts: one launch over n_experts experts whose block pointers sit in the device table vxs. The activation for expert e is y_q8 + e * y_stride_bytes (0 = shared), its output out_bf16 + e * m * n. Each expert’s rows go through the same device function as kquant_mmvq.
kquant_mmvq_experts_w: 2026-09-25: kquant_mmvq_*_experts_w: the warp-per-row expert batch, with the same arguments as kquant_mmvq_experts.
kquant_mmvq_experts_wn: 2026-09-25: kquant_mmvq_experts_w with nwarps rows a block (the _w2 / _w8 entries for 2 / 8; the plain _w entry is 4). One warp computes each row in the same way for any nwarps. Fails unless nwarps is 2, 4 or 8.
kquant_mmvq_groups_w: 2026-09-25: kquant_mmvq_q2_k_groups_w: n_groups independent [n, k] GEMVs in one launch. Group g takes weight rows [g*n, (g+1)*n) (w_blocks is [n_groups * n][k/256]) against activation columns [g*k, (g+1)*k) of the block_q8_1 [m][n_groups * k / 32] buffer y_q8, and writes columns [g*n, (g+1)*n) of out_bf16 ([m][n_groups * n]). Each group runs the per-row code of kquant_mmvq_w with offsets.
kquant_mmvq_pair_w: 2026-09-25: kquant_mmvq_q2_k_pair_w: two [n_i, k] Q2_K projections of one activation (y_q8, block_q8_1 [m][k/32]) in one launch, out_i = [m][n_i] bf16. Grid y picks the projection, and each runs the same device function as kquant_mmvq_w.
kquant_mmvq_w: 2026-09-25: kquant_mmvq_*_w: the warp-per-row GEMV (four rows a block, reduced by warp shuffle), with the same arguments as kquant_mmvq.
kquant_q8_1_rows: 2026-09-25: bf16 [m, k] -> block_q8_1 [m, k/32] (kquant_q8_1_rows_bf16), one warp per block. Fails unless k % 32 == 0.
kquant_q8_1_rows_bytes: 2026-09-25: Bytes of the plain block_q8_1 buffer for [m, k] activations.
kquant_swiglu_q8_1_rows: 2026-09-25: moe_v41_swiglu and kquant_q8_1_rows in one launch (kquant_swiglu_q8_1_rows_bf16): h[rows, inter] = bf16(swiglu(gate, up) * w[row]), then the q8_1 blocks of those bf16 rows. kquant_fold_tests compares both outputs byte for byte with the two launches. w may be null. Fails unless inter % 32 == 0.
kquant_weight_bytes: 2026-09-25: Raw bytes of an [n, k] K-quant weight.
l2_norm: 2026-09-25: In-place L2 normalization of each head, x / sqrt(sum(x^2) + eps) (kernel l2_norm_bf16). data is [num_tokens, stride] BF16 with num_heads * head_dim used per token; one block per (head, token).
log_cutlass_nvfp4_route
log_gemm_shape: 2026-09-25: With METRALE_GEMM_SHAPE_LOG=1, log each (kernel, M, N, K) GEMM shape once per backend at warn level, with its FLOP count.
mamba2_ssd_bmm: 2026-09-25: K2: CB[c][g][t][s] = C_t . B_s (raw, FP32).
mamba2_ssd_cumsum: 2026-09-25: K1: per-chunk dt (softplus and clamp) and the inclusive cumsum of the log-decay.
mamba2_ssd_scan: 2026-09-25: K3: chunk_state, state_passing and chunk_scan fused in one kernel. The running state h0 stays in shared memory, so no per-chunk state tensor is written to global memory.
mamba2_ssm_prefill: 2026-09-25: Mamba-2 SSM prefill (mamba2_ssm_prefill): the token-sequential recurrence over seq_len tokens in one launch. x_stride, bc_stride, dt_stride and y_stride are BF16 elements between consecutive tokens.
mamba2_ssm_prefill_persistent: 2026-09-25: Mamba-2 SSM prefill with each block’s H slice in shared memory for the whole token loop (mamba2_ssm_prefill_persistent). Same arguments as mamba2_ssm_prefill, but launched with head_dim * SUB threads and dynamic shared memory.
mla_batched_gemv: 2026-09-25: MLA batched GEMV over all heads in one launch: output[head, n] = sum_k(weight[head, n, k] * input[head, k]), BF16, with the per-head input and output at input_stride / output_stride and the weight contiguous per head.
mla_cache_assemble: 2026-09-25: Assemble one token’s cache entries: K = [kv_latent | k_rope] and V = [kv_latent | zeros], each mla_cache_dim wide. The block has one thread per element (max(mla_cache_dim, 256) threads).
mla_cache_assemble_batched: 2026-09-26: Per token, write the compressed MLA cache rows: K = [kv_latent | k_rope], and V = [kv_latent | zeros] in the mistral-small-4 kernel or [kv_latent | k_rope] in the deepseek-v4-flash kernel.
mla_fused_prefill: 2026-09-25: MLA prefill in one kernel: Q absorption through w_uk, attention over the latent cache, and V extraction through w_uv. One block per (head, query token).
mla_kv_assemble_batched: 2026-09-25: Per token, build K as [k_nope | k_rope] for each KV head (k_rope shared by all heads) and extract V from kv_expanded. blockIdx.y 0 builds K, 1 extracts V.
mla_paged_decode_fp8: 2026-09-25: MLA paged decode for DeepSeek-V4-Flash over an FP8 cache (kernel mla_paged_decode_fp8), with scalar dequant scales k_scale / v_scale and cache_stride in bytes. sinks (per-head attention sinks, FP32) may be null. The compressed-KV pool is attended only when comp_pool is non-null and comp_block_count > 0.
mla_paged_decode_nvfp4: 2026-09-25: MLA paged decode for DeepSeek-V4-Flash over an NVFP4 cache (kernel mla_paged_decode_nvfp4). A cache token is kv_lora_rank + qk_rope_head_dim = 512 + 64 = 576 values (MLA_CACHE_DIM in the kernel); each block holds packed E2M1 data (data_section_bytes) followed by one FP8 E4M3 scale per 16 values.
mla_prefill_attention_320: 2026-09-25: Absorbed MLA prefill attention with scalar BF16 dot products and a single KV head. The head dim is the kernel’s MLA_HDIM: 320 in the mistral-small-4 fork, 576 in the deepseek-v4-flash one; both keep the entry name mla_prefill_attn_320.
mla_q_final_assemble_batched: 2026-09-25: Build Q as [q_absorbed (kv_lora) | q_rope (rope)] per head and token, mla_cache_dim = kv_lora + rope wide.
mla_q_rope_extract_batched: 2026-09-25: Copy each head’s rope slice (offset nope in the head’s hd) out of Q, rows q_dim apart, into contiguous [num_tokens, nq, rope].
mla_q_rope_scatter: 2026-09-25: Copy each head’s rope part of q_full ([nq, hd], at offset nope) into q_absorbed_buf ([nq, mla_cache_dim], at offset kv_lora) and into the contiguous q_rope_contiguous ([nq, rope]) in one pass.
mla_q_rope_writeback: 2026-09-25: Write the rotated [nq, rope] Q rope parts back into q_absorbed_buf at offset kv_lora of each head’s mla_cache_dim row.
mla_q_rope_writeback_batched: 2026-09-25: Inverse of mla_q_rope_extract_batched: write [num_tokens, nq, rope] back into each head’s rope slice of Q.
moe_act_quant_e4m3: 2026-09-28: E4M3 per (row, 128-K group) of x [rows, k] BF16 into q / s.
moe_batched_blend: 2026-09-25: output[t] += sigmoid(dot(normed[t], gate_weight)) * shared_out[t], one block per token. A null gate_weight blends at weight 1.
moe_bf16_grouped_gemm: 2026-09-25: BF16 grouped GEMM for sorted MoE prefill, without scales: input [total_tokens, K] BF16, weight_ptrs[e] → [N, K] BF16, output [total_expanded, N] BF16, sorted_token_ids [total_expanded] or NULL. The MoE layer’s long-prefill path uses it when its experts are BF16, which the loader sets up when METRALE_FP8_DEQUANT_MOE_TO_BF16=1 dequantizes FP8 experts at load (qwen35/load_layers.rs).
moe_build_tile_worklist: 2026-09-25: Build the compacted work-list for the grouped-GEMM grid: one item per (expert, m-tile, n-tile) of every expert that has rows and a non-NULL weight pointer, experts in ascending order. One CTA; thread 0 writes worklist[2w] = expert, worklist[2w + 1] = (m-tile index << 6) | n-tile index, and total_tiles[0] = w.
moe_decode_atomic_c4_finalize
moe_decode_atomic_c4_silu_down_accum
moe_expert_down_act_bf16_grouped: 2026-10-02: Grouped BF16 down over the hi + lo SiLU rows. k is the intermediate width; rows of act and output are sorted positions, rows of sh_act and sh_down_out tokens.
moe_expert_down_act_fp8_grouped: 2026-09-26: Grouped FP8 down over the SiLU product, launched with geometry (FP8_GROUPED_DOWN_SCALAR reads FP32 products, FP8_GROUPED_DOWN_TC BF16). k is the intermediate width; rows of act and output are sorted positions, rows of sh_act and sh_down_out tokens.
moe_expert_down_act_fp8_grouped_tc_w8a8: 2026-09-28: W8A8 down over the quantized SiLU products: routed rows into output by sorted position, the shared rows into sh_down_out by token, both BF16.
moe_expert_down_act_nvfp4_grouped
moe_expert_gate_up_act_bf16_grouped: 2026-10-02: Grouped BF16 gate+up and SiLU. gate_ptrs / up_ptrs are device tables of each routed expert’s [n, k] BF16 weight; sh_gate / sh_up the shared expert’s. Writes the SiLU product as BF16 hi + lo rows ([positions, 2n] into act, [num_tokens, 2n] into sh_act).
moe_expert_gate_up_act_fp8_grouped: 2026-09-26: Grouped FP8 gate+up and SiLU. cap is fp8_grouped_active_cap; the first ceil(num_tokens / geometry.rows_per_pass) block rows are the shared expert. Writes the product silu(bf16(gate)) * bf16(up), [positions, n] for the routed experts into act and [num_tokens, n] for the shared expert into sh_act: FP32 from the scalar kernel (FP8_GROUPED_GATE_UP_SCALAR), BF16 from the tensor-core one (FP8_GROUPED_GATE_UP_TC, 2026-09-28).
moe_expert_gate_up_act_fp8_grouped_tc_w8a8: 2026-09-28: W8A8 gate+up and SiLU: (xq, xs) the quantized layer input; writes the quantized products (act_q, act_s) by sorted position and (sh_q, sh_s) by token.
moe_expert_gate_up_act_nvfp4_grouped: 2026-09-27: Grouped NVFP4 gate+up and SiLU over k-wide input rows. cap is fp8_grouped_active_cap; the first ceil(num_tokens / geometry.rows_per_pass) block rows are the shared expert. Writes the FP32 product silu(bf16(gate)) * bf16(up), [positions, n] for the routed experts into act and [num_tokens, n] for the shared expert into sh_act (2026-10-02: the tensor-core kernel writes it as BF16 hi + lo in the same bytes, which only its down kernel reads).
moe_expert_gate_up_shared: 2026-09-25: NVFP4 gate and up GEMVs of the routed experts and the shared expert in one launch.
moe_expert_gate_up_shared_batch2: 2026-09-25: Gate and up GEMVs of the routed and shared experts for two tokens in one launch (row layout in the module doc). The shared expert uses direct weight pointers, the routed experts the pointer tables. A block_size of 256 selects the kernel’s wide decomposition (blockDim.x == 256); any other width runs the 128-thread one.
moe_expert_gate_up_shared_batch3: 2026-09-25: moe_expert_gate_up_shared_batch2 for three tokens.
moe_expert_gate_up_shared_batch2_t: 2026-09-25: NVFP4 fused gate+up GEMV over transposed weights, for two tokens.
moe_expert_gate_up_shared_batch3_t: 2026-09-25: NVFP4 fused gate+up GEMV over transposed weights, for three tokens.
moe_expert_gate_up_shared_bf16: 2026-09-25: BF16 variant of moe_expert_gate_up_shared_fp8, with no scale tables, for experts installed by MoeLayer::set_bf16_experts.
moe_expert_gate_up_shared_bf16_batch2: 2026-09-25: moe_expert_gate_up_shared_bf16 for two tokens in one launch (called from MoeLayer::forward_k2). blockIdx.y in [0, 2*top_k) is a routed expert (token y / top_k, slot y % top_k, output row token * top_k + slot); blockIdx.y == 2*top_k is the shared expert, computed for both tokens into rows 0 and 1.
moe_expert_gate_up_shared_fp8: 2026-09-25: FP8 variant of moe_expert_gate_up_shared: two pointer tables per projection (weights, scales) and the shared expert as direct Fp8Weight pointers.
moe_expert_gate_up_shared_fp8_batch2: 2026-09-25: FP8 fused gate+up GEMV over the routed experts and the shared expert, for two tokens (moe/forward_k2.rs).
moe_expert_gate_up_shared_fp8_batch3: 2026-09-25: FP8 fused gate+up GEMV over the routed experts and the shared expert, for three tokens (moe/forward_k3.rs).
moe_expert_gate_up_shared_fp8_batch2_t: 2026-09-25: FP8 fused gate+up GEMV over transposed weights, for two tokens.
moe_expert_gate_up_shared_fp8_batch3_t: 2026-09-25: FP8 fused gate+up GEMV over transposed weights, for three tokens.
moe_expert_gate_up_shared_fp8_t: 2026-09-25: FP8 fused gate+up GEMV over transposed weights, for one token.
moe_expert_gate_up_shared_prefill: 2026-09-25: Gate and up GEMVs for num_tokens tokens, routed and shared experts in one launch. Row y < num_tokens * top_k is routed (token y / top_k, expert expert_indices[y]) and writes gate_out/up_out ([num_tokens * top_k, n] BF16); the remaining num_tokens rows run the shared expert and write sh_gate_out/sh_up_out ([num_tokens, n] BF16). blockIdx.z picks gate (0) or up (1). input is [num_tokens, k] BF16.
moe_expert_gate_up_shared_t: 2026-09-25: NVFP4 fused gate+up GEMV over transposed weights, for one token.
moe_expert_gemv: 2026-09-25: All top_k expert W4A16 GEMVs in one launch, output [top_k, N]. input_stride is 0 when every expert reads the same input row (gate/up) and K when slot e reads row e (down). A null weight pointer (a remote expert under EP) writes zeros.
moe_expert_gemv_gate_up: 2026-09-25: Gate and up expert GEMVs in one launch; blockIdx.z 0 is gate, 1 is up.
moe_expert_gemv_gate_up_2x: 2026-09-25: moe_expert_gemv_gate_up with two adjacent output rows per 32-thread group, 8 outputs per block.
moe_expert_gemv_silu_down: 2026-09-25: Down expert GEMV whose activation silu(gate) * up is computed inline from gate_out and up_out.
moe_expert_gemv_silu_down_2x: 2026-09-25: moe_expert_gemv_silu_down with two adjacent output rows per 32-thread group, 8 outputs per block.
moe_expert_silu_down_shared: 2026-09-25: NVFP4 SiLU+down GEMVs in one launch. Routed experts read gate_out/up_out; the shared expert reads sh_gate_in/sh_up_in. The kernel’s s_act is dynamic shared memory of K floats.
moe_expert_silu_down_shared_batch2: 2026-09-25: SiLU+down GEMVs of the routed and shared experts for two tokens in one launch. block_size as in moe_expert_gate_up_shared_batch2.
moe_expert_silu_down_shared_batch3: 2026-09-25: moe_expert_silu_down_shared_batch2 for three tokens.
moe_expert_silu_down_shared_batch2_t: 2026-09-25: NVFP4 fused SiLU+down GEMV over transposed weights, for two tokens.
moe_expert_silu_down_shared_batch3_t: 2026-09-25: NVFP4 fused SiLU+down GEMV over transposed weights, for three tokens.
moe_expert_silu_down_shared_bf16: 2026-09-25: BF16 variant of moe_expert_silu_down_shared_fp8, with no scale tables.
moe_expert_silu_down_shared_bf16_batch2: 2026-09-25: moe_expert_silu_down_shared_bf16 for two tokens, with the moe_expert_gate_up_shared_bf16_batch2 row layout.
moe_expert_silu_down_shared_fp8: 2026-09-25: FP8 variant of moe_expert_silu_down_shared: two pointer tables (weights, scales) and the shared down weight as a direct Fp8Weight pointer.
moe_expert_silu_down_shared_fp8_batch2: 2026-09-25: FP8 fused SiLU+down GEMV over the routed experts and the shared expert, for two tokens.
moe_expert_silu_down_shared_fp8_batch3: 2026-09-25: FP8 fused SiLU+down GEMV over the routed experts and the shared expert, for three tokens.
moe_expert_silu_down_shared_fp8_batch2_t: 2026-09-25: FP8 fused SiLU+down GEMV over transposed weights, for two tokens.
moe_expert_silu_down_shared_fp8_batch3_t: 2026-09-25: FP8 fused SiLU+down GEMV over transposed weights, for three tokens.
moe_expert_silu_down_shared_fp8_t: 2026-09-25: FP8 fused SiLU+down GEMV over transposed weights, for one token.
moe_expert_silu_down_shared_prefill: 2026-09-25: SiLU(gate) * up, then the down GEMV, for the routed and shared rows of moe_expert_gate_up_shared_prefill, with the same row split. Routed rows read gate_out/up_out and write output ([num_tokens * top_k, n] BF16); shared rows read sh_gate_in/sh_up_in and write sh_down_out ([num_tokens, n] BF16).
moe_expert_silu_down_shared_t: 2026-09-25: NVFP4 fused SiLU+down GEMV over transposed weights, for one token.
moe_fp8_grouped_gemm: 2026-09-25: FP8 grouped GEMM for sorted MoE prefill over the work-list from moe_build_tile_worklist, which must run first on the same stream.
moe_fp8_grouped_gemm_ptrtable_n128: 2026-09-25: Pointer-table grouped GEMM with FP8 E4M3 activations and transposed NVFP4 weights; a_fp8 must already be FP8. Same launch shape as moe_w4a16_grouped_gemm_ptrtable_n128.
moe_fp8_grouped_sort: 2026-09-27: Sorts the total_expanded slots of topk_ids by expert, as moe_sort_by_expert does, and writes active_experts[0..count] in ascending order with active_count[0] = count, in one launch on stream (the grouped kernels that read them must follow on it).
moe_gate_topk_fused: 2026-09-25: Fused gate GEMV and top-k for one token (kernel moe_gate_topk_fused): gate[num_experts] = A[K] @ B_gate[num_experts, K]^T with an NVFP4 gate weight, then the top-k expert indices and weights exp(logit - max) over those k, divided by their sum when normalize is nonzero. One 256-thread CTA; K BF16 values of dynamic shared memory hold the input.
moe_grouped_down_cutlass: 2026-09-25: Single-launch CUTLASS grouped NVFP4 down GEMM through metrale_gpu_runtime::cutlass::nvfp4_grouped_down. a is the expert-sorted post-SiLU intermediate [total_expanded, inter]; expert_offsets is the device prefix sum.
moe_grouped_gate_up_cutlass: 2026-09-25: Single-launch CUTLASS grouped NVFP4 gate_up GEMM through metrale_gpu_runtime::cutlass::nvfp4_grouped_gate_up_fused, with the weights from host. expert_offsets is the device i32 [num_experts + 1] prefix sum; the host entry needs it on the host, so it is copied and the stream synchronised first. Returns that host copy, which moe_grouped_down_cutlass accepts to skip its own copy and synchronise.
moe_hash_route: 2026-09-25: Hash router, one token: the experts are the static tid2eid[token_id] row, and the gate’s sqrt-softplus scores weight them. The token id is read on the device from token_id_ptr.
moe_hash_route_batched: 2026-09-25: moe_hash_route for n tokens, one block per token, reading token_ids[n].
moe_permute_tokens: 2026-09-25: Gather token rows into expert-sorted order: permuted[i] = hidden[sorted_token_ids[i]], permuted being [total_expanded, hidden]. One block per output row. No code calls it.
moe_router_gemm: 2026-09-26: The MoE router gate GEMM, C = A @ B^T with the operands of dense_gemm and its bits (kernel moe_router_gemm_bf16): each output keeps one FP32 accumulator over k = 0..K-1 in order, but a block stages a 512-wide K slice of 4 weight rows and up to 16 activation rows in shared memory at once instead of waiting on every 16-wide tile. k must be a multiple of 16 (below that dense_gemm adds zero padding terms).
moe_router_gemm_rt: 2026-09-28: dense_gemm_router on 64x128 tiles with a 4x8 register tile per thread (kernel moe_router_gemm_rt, kernels/gb10/common/moe_router_gemm_prefill.cu): the same per-output FP32 accumulation in ascending k, so the same bits. Needs k % 16 == 0.
moe_silu_mul: 2026-09-25: output[i] = silu(gate[i]) * up[i] over total_elements BF16 values.
moe_sort_by_expert: 2026-09-25: Counting sort of the total_expanded (token, slot) routes by expert, in one block: writes sorted_token_ids and sorted_expert_ids in expert order, expert_offsets (the [num_experts + 1] prefix sum) and token_to_perm (each route’s sorted position).
moe_topk_sigmoid: 2026-09-25: Sigmoid router, one token. The bias is added to the scores for selection only; the weights come from the pre-bias sigmoid scores. top_k and num_experts are bounded by MOE_TOPK_SIGMOID_MAX_TOP_K and MOE_TOPK_SIGMOID_MAX_EXPERTS.
moe_topk_sigmoid_batched: 2026-09-25: The sigmoid router with a selection-only bias for num_tokens tokens.
moe_topk_softmax: 2026-09-25: Top-k experts of one token’s gate logits, with softmax weights, in one block. The kernel handle decides the logit type (moe_topk_softmax BF16, moe_topk_softmax_f32).
moe_topk_softmax_batched: 2026-09-25: The top-k softmax router for num_tokens tokens. The handle picks the tie-break: moe_topk_softmax_batched keeps the lower lane on an equal logit, moe_topk_softmax_rows (2026-09-26) the lower expert index, as moe_topk_softmax does.
moe_topk_softmax_bias: 2026-09-25: Softmax router with a correction bias and the zero-expert fold, one token. Logits cover num_logits experts, of which ids >= num_routed are identity experts: a selected one adds its weight to zero_accum and its slot becomes expert 0 with weight 0.
moe_topk_softmax_bias_batched: 2026-09-25: moe_topk_softmax_bias for n tokens, one block per token.
moe_topk_sqrtsoftplus: 2026-09-25: Sqrt-softplus router (sqrt(log(1 + exp(logit)))), one token. As in moe_topk_sigmoid, the bias steers selection only and the weights come from the pre-bias scores.
moe_topk_sqrtsoftplus_batched: 2026-09-25: Same arguments as moe_topk_sigmoid_batched, scoring experts with sqrt(log(1 + exp(logit))), the score of the single-token moe_topk_sqrtsoftplus.
moe_transpose_u8_batched: 2026-09-25: Per-expert uint8 transpose for the MoE down_proj relayout: reads expert e’s [rows, cols] matrix through src_ptrs[e] and writes it as [cols, rows] through dst_ptrs[e]. A NULL entry in either table skips that expert.
moe_unpermute_blend: 2026-09-28: moe_unpermute_reduce_indexed then moe_batched_blend in one launch (kernel moe_unpermute_blend, kernels/gb10/common/moe_unpermute_blend.cu), with the same output bits; one 256-thread block per token. Needs hidden_size % 8 == 0.
moe_unpermute_reduce_indexed: 2026-09-25: For each token, output[t] = sum_k topk_weights[t, k] * expert_output[token_to_perm[t, k]], one block per token.
moe_w4a4_grouped_gemm_relu2: 2026-09-25: Grouped W4A4 expert up GEMM on the FP4 block-scale MMA, with relu^2 applied to the accumulator before the BF16 store. A is NVFP4 (packed E2M1 + per-16 E4M3 scales) in token order, gathered through sorted_token_ids; B is read from the per-expert NVFP4 pointer tables. blockIdx.z is the expert, and expert_offsets bounds its rows. The kernel ships only in kernels/gb10/nemotron-labs-3-puzzle-75b-a9b/nvfp4/moe_w4a4_grouped.cu.
moe_w4a16_fused_gate_up_k64_m128: 2026-09-25: moe_w4a16_fused_gate_up_k64_n128 with a 128-row M tile and a 256-thread block. max_m_tiles_m128 counts 128-row tiles: the call site halves the 64-row count (forward_prefill_routed.rs).
moe_w4a16_fused_gate_up_k64_n128: 2026-09-25: moe_w4a16_fused_gate_up_n128 for the kernels with a 64-wide K step.
moe_w4a16_fused_gate_up_n128: 2026-09-25: Gate and up grouped GEMMs in one launch: the grid spans 2 * n_out columns, the first n_out for gate and the rest for up.
moe_w4a16_grouped_gemm: 2026-09-25: Launcher for moe_w4a16_grouped_gemm, whose weights are contiguous per expert rather than in pointer tables. No production code calls it (@human-review: its grid (num_experts, 1, 1) does not match the kernel, which takes the expert from blockIdx.z).
moe_w4a16_grouped_gemm_ptrtable: 2026-09-29: Pointer-table grouped W4A16 GEMM, 128 threads; one launch covers all experts.
moe_w4a16_grouped_gemm_ptrtable_k64_n128: 2026-09-25: Same launch shape as moe_w4a16_grouped_gemm_ptrtable_n128, for the kernels with a 64-wide K step.
moe_w4a16_grouped_gemm_ptrtable_m256: 2026-09-25: moe_w4a16_grouped_gemm_ptrtable with a 256-row M tile and a 512-thread block. max_m_tiles counts 256-row tiles: MoeLayer::launch_grouped_gemm divides the 64-row count by 4. Loaded only under METRALE_MOE_GROUPED_M256=1 (moe/init.rs); the measurement behind that is on launch_grouped_gemm.
moe_w4a16_grouped_gemm_ptrtable_n128: 2026-09-29: Pointer-table grouped W4A16 GEMM, 128 threads. The name is historical: the grid follows the kernel’s published N tile, 64 or 128.
moe_w8a8_gateup_silu_e4m3: 2026-09-28: moe_w8a8_gateup_silu_e4m3_w1 / _w2: the gate and up grouped GEMMs of one work-list item (the _gu shape, MOE_E4M3_GU) followed by SiLU(gate) * up quantized to E4M3 per 128-column group, bit-identical to the two _gu launches plus silu_mul_quant_fp8. Writes out_fp8 [total_expanded, n] and out_scale [total_expanded, n / 128]. The work-list is built with MOE_E4M3_GU’s shape.
moe_w8a8_grouped_gemm: 2026-09-25: W8A8 grouped MoE GEMM on a dense grid (ceil(N/64), max_m_tiles, num_experts). a_fp8 comes from per_token_group_quant_fp8; a_scale ([total_tokens, K/128] FP32) and the per-block b_scale (scale_ptrs[e] → [N/128, K/128] FP32) are applied in an FP32 epilogue per 128-element K-group. weight_ptrs[e] points at [N, K] FP8; output is [total_expanded, N] BF16; sorted_token_ids is [total_expanded] or NULL.
moe_w8a8_grouped_gemm_e4m3: 2026-09-27: moe_w8a8_grouped_gemm_pm4 on the native e4m3 tensor-core MMA (the _gu or _dn entry, described by tile): same operands and output, bit-identical, over a work-list that moe_build_tile_worklist built on the same stream with tile.m_tile and n / tile.n_tile. grid_ctas (clamped to 1..=16384) CTAs stride over the items.
moe_w8a8_grouped_gemm_pm4: 2026-09-25: moe_w8a8_grouped_gemm over the work-list from moe_build_tile_worklist (kernel moe_w8a8_grouped_gemm_pm4, same module), which must run first on the same stream. The grid follows the rule of moe_fp8_grouped_gemm: max_tiles clamped to 1..=16384.
moe_weighted_sum_blend: 2026-09-25: The MoE combine for one token: output[j] = sum_e weights[e] * expert_out[e, j] + sigmoid(dot(input, gate_weight)) * shared_out[j]. Every block computes the gate dot product itself. A null gate_weight means an ungated shared expert, blended at weight 1.
moe_weighted_sum_blend_batch2: 2026-09-25: moe_weighted_sum_blend for two tokens; blockIdx.y is the token.
moe_weighted_sum_blend_batch3: 2026-09-25: moe_weighted_sum_blend for three tokens.
moe_weighted_sum_blend_fp8_grouped: 2026-09-25: Grouped blend, one blockIdx.y per token. expert_out rows are sorted positions; token_to_perm[token * top_k + k] maps a slot to its row. gate_weight may be NULL (ungated shared expert).
moe_weighted_sum_blend_prefill: 2026-09-25: Per token t: output[t] = sum_e expert_weights[t*top_k+e] * expert_out[t*top_k+e] + sigmoid(dot(input[t], gate_weight)) * shared_out[t]. expert_out is [num_tokens * top_k, hidden] BF16, expert_weights [num_tokens * top_k] f32, input [num_tokens, k] BF16 and gate_weight the shared-expert gate [1, k] BF16.
moe_zero_expert_add: 2026-09-25: out[t, :] += zero_accum[t] * x[t, :], the identity-expert blend.
n_tile_blocks: 2026-09-29: ceil(n / t) blocks, t being kernel’s published N tile. An error when the kernel publishes none.
native_q2_mmq_enabled: 2026-09-25: True only when METRALE_GGUF_NATIVE_Q2_MMQ is exactly 1. With it, and with the kernels resolved and group 128 on gate, up and down, the dense FFN prefill runs packed Q2 weights through q2_0_mmq_gemm instead of dequantizing them to BF16. Whether weights stay packed at all is METRALE_GGUF_NATIVE_Q2’s decision.
nvfp4_grouped_tc_shape_ok: 2026-10-02: Whether the tensor-core kernels take an n-column, k-deep projection: k in whole load groups of two 128-wide chunks, n in whole CTAs of geometry.
nvfp4_lean_repack_shape_ok: 2026-10-04: nvfp4_tc_lean_repack admits whole 16-row tiles, whole 128-K chunks, and a tile that fits the default 48 KiB of shared memory.
nvfp4_mmq_gemm: 2026-09-25: C[m, n] = A[m, k] x W[n, k] in BF16, without weight_scale_2, on the 128-wide M tile. a_fp4 is block_fp4_mmq, w_nvfp4 the nvfp4_mmq_repack output. kernel_wc is used when n is not a multiple of 128.
nvfp4_mmq_gemm_tiled: 2026-09-25: nvfp4_mmq_gemm with the M tile mmq_x chosen by the caller. 2026-09-28: With out_scale, the pipelined 128-row tile applies it in its store and the call returns true; otherwise the caller still owes nvfp4_scale_bf16. The tile choice: the kernels must be that tile’s entries. nvfp4_mmq.cu instantiates 16, 32, 64 and 128. A 128-wide tile issues MMAs for all 128 columns whatever m is, so a small m wants a small tile. grid.y = ceil(m / mmq_x), and each M tile streams the whole weight again, so m must not exceed mmq_x (debug-asserted).
nvfp4_mmq_quantize_act: 2026-09-25: Quantize BF16 activations [m, k] to block_fp4_mmq in out_y: E2M1 values with one UE4M3 scale per 16, k padded to 256, one thread per 16-value group.
nvfp4_mmq_repack: 2026-09-25: Repack a checkpoint NVFP4 weight (E2M1 [n, k/2], low nibble = even k, and E4M3 [n, k/16] scales) for the MMQ kernel. Each row holds k/64 blocks as all nibble bytes (32 per block) followed by all scale bytes (4 per block). Codes and scale bytes are copied, not requantized.
nvfp4_mmq_smem: 2026-09-25: Dynamic shared memory for M tile mmq_x, in bytes: mmq_x ids, a y tile of mmq_x * 36 ints padded to 256 ints, and an x tile of 128 * 76 ints. This is mmq_get_nbytes_shared in q4k_vendor/mmq.cuh with mmq_y = 128, MMQ_MMA_TILE_X_K_FP4 = 76 and 8 warps of 32.
nvfp4_mmq_weight_bytes: 2026-09-25: Bytes of the repacked form of an [n, k] weight; k must be a multiple of 64.
nvfp4_pipe_gemm: 2026-09-28: C[m, n] = A[m, k] x W[n, k] in BF16, without weight_scale_2, through metrale_nvfp4_gemm_pipe (same operands and output as nvfp4_mmq_gemm, bit-identical). k must be a multiple of 256. Grid of 128-channel x 192-token tiles, 256 threads. With out_scale, each output is multiplied by it as nvfp4_scale_bf16 does in place, bit for bit.
nvfp4_scale_bf16: 2026-09-25: data[i] *= scale in place over total BF16 values; the caller passes the down projection’s weight_scale_2.
nvfp4_silu_mul_quant: 2026-09-25: silu(gate * gate_scale) * (up * up_scale) quantized straight to block_fp4_mmq in out_y for the down GEMM, without writing a BF16 intermediate. k is the intermediate width, padded to 256. The kernel applies no clamp.
nvfp4_silu_mul_scaled: 2026-09-25: out[i] = silu(gate[i] * gate_scale) * (up[i] * up_scale) over total BF16 values. The kernel applies no clamp.
nvfp4_tc_lean_repack: 2026-09-27: Grouped NVFP4 down over the SiLU product. k is the intermediate width; rows of act and output are sorted positions, rows of sh_act and sh_down_out tokens. 2026-10-04: Repacks one row-major NVFP4 matrix [n, k] (packed E2M1 [n, k / 2] at packed, E4M3 block scales [n, k / 16] at scale) in place into the lean layout the _lean tensor-core expert kernels read (nvfp4_tc_lean_repack, moe_nvfp4_grouped_tc.cu): same bytes and size, outputs byte-identical to the row-major point. Sets the device u32 at bad (zeroed by the caller) on a negative or NaN scale, after which the matrix is unusable. Needs nvfp4_lean_repack_shape_ok.
paged_decode_attn_bf16: 2026-09-25: Paged BF16-KV decode attention, one block per (q_head, seq). sliding_window == 0 attends to the whole sequence; otherwise only the last sliding_window positions.
paged_decode_attn_bf16_gqa: 2026-09-25: GQA-packed BF16-KV paged decode, one block per (kv_head, seq). Same arguments, argument order and block as paged_decode_attn_bf16; only the grid’s x extent changes, to num_kv_heads, because the kernel (kernels/gb10/common/paged_decode_attn_bf16_gqa.cu) keeps a KV head’s whole query group in registers.
paged_decode_attn_bf16k_turbo2v: 2026-09-25: Paged decode attention over the BF16-K / Turbo2V cache (kernel paged_decode_attn_bf16k_turbo2v), with the same sparse-V skip as the Turbo3V kernel.
paged_decode_attn_bf16k_turbo3v: 2026-09-25: Paged decode attention over the BF16-K / Turbo3V cache (kernel paged_decode_attn_bf16k_turbo3v). The kernel skips a position’s V load and dequant when its attention weight is at most TQ_PLUS_SPARSE_V_THRESHOLD (1e-3 by default).
paged_decode_attn_bf16k_turbo4v: 2026-09-25: Paged decode attention over the BF16-K / Turbo4V cache (kernel paged_decode_attn_bf16k_turbo4v), with the same sparse-V skip as the Turbo3V kernel.
paged_decode_attn_fp8
paged_decode_attn_fp8_gqa: 2026-09-25: GQA-packed FP8-KV paged decode, one block per (kv_head, seq): the arguments and order of paged_decode_attn_fp8, grid x extent num_kv_heads, and the preconditions of paged_decode_attn_bf16_gqa.
paged_decode_attn_fp8k_turbo2v: 2026-09-25: Paged decode attention over the FP8-K / Turbo2V cache (kernel paged_decode_attn_fp8k_turbo2v).
paged_decode_attn_fp8k_turbo3v: 2026-09-25: Paged decode attention over the FP8-K / Turbo3V cache (kernel paged_decode_attn_fp8k_turbo3v), one 256-thread block per (query head, sequence). The kernel skips a position’s V load and dequant when its attention weight is at most TQ_PLUS_SPARSE_V_THRESHOLD.
paged_decode_attn_fp8k_turbo4v: 2026-09-25: Paged decode attention over the FP8-K / Turbo4V cache (kernel paged_decode_attn_fp8k_turbo4v).
paged_decode_attn_nvfp4: 2026-09-25: Paged decode attention over an NVFP4 KV cache, one block per (q_head, seq).
paged_decode_attn_reduce_fp8: 2026-09-25: Combine split-K partials into the BF16 output. The workspace is F32 [head_dim values, m, l] per split whatever the cache type, so the BF16 split-K’s reduce handle is launched through this function too (qwen3_attention/decode/splitk_dispatch.rs).
paged_decode_attn_reduce_nvfp4: 2026-09-25: Combine the split-K partials in workspace into the BF16 output.
paged_decode_attn_splitk_bf16: 2026-09-25: The BF16-KV counterpart of paged_decode_attn_splitk_nvfp4 (kernels/hopper/common/paged_decode_bf16_splitk_hopper.cu). It takes no cache_stride: the cache is [blocks, block_size, kv_heads, head_dim] and the kernel computes the page stride, as paged_decode_attn_bf16 does.
paged_decode_attn_splitk_fp8: 2026-09-25: The FP8-KV counterpart of paged_decode_attn_splitk_nvfp4.
paged_decode_attn_splitk_nvfp4: 2026-09-25: Split-K paged decode attention over an NVFP4 KV cache: each (q_head, seq) sequence is cut into num_splits parts, and each block writes its partial output and softmax state to workspace.
paged_decode_attn_turbo3k_turbo8v: 2026-09-25: Paged decode attention over the Turbo3K / Turbo8V cache (kernel paged_decode_attn_turbo3k_turbo8v).
paged_decode_attn_turbo4k_turbo3v: 2026-09-25: Paged decode attention over the Turbo4K / Turbo3V cache (kernel paged_decode_attn_turbo4k_turbo3v).
paged_decode_attn_turbo4k_turbo8v: 2026-09-25: Paged decode attention over the Turbo4K / Turbo8V cache (kernel paged_decode_attn_turbo4k_turbo8v).
parse_cublas_scope: 2026-09-25: Parse the CublasScope grammar. Returns the resolved set and the lowercased tokens that matched nothing. Pure: the environment read and the logging are in GemmDispatch::from_env.
per_token_group_quant_fp8: 2026-09-25: Per-token FP8 activation quantization over 128-element K-groups: A_fp8 [M, K] FP8 E4M3 and a_scale [M, K/128] FP32.
ple_add_highway: 2026-09-25: highway += ple_out, in FP32. The reference model (bench/qwen4_exp/ref/modeling_qwen4_exp.py) adds PLE’s output to the hidden state before that layer’s attention hyper-connection.
ple_conv: 2026-09-25: Depthwise causal conv with kernel k_size and dilation dilation, then out = gated + silu(conv). All buffers but weight (BF16) are FP32. state is [(k_size - 1) * dilation, channels] and is rolled in place, so prefill and decode use the same launch.
ple_gate: 2026-09-25: Gate the n-gram value by the highway, and write both the gated value and its norm_conv output. hidden is the FP32 highway [T, hc*H]; key and value are BF16 projection outputs. Both outputs are FP32, because the chain ends on the FP32 highway (the precision note in ple.cu). One block per token.
predequant_nvfp4_to_fp8: 2026-09-25: Dequantize an NVFP4 weight to FP8 E4M3: reads B_packed[N, K/2], B_scale[N, K/16] and scale2, writes B_fp8[N, K]. QuantizedWeight::predequant_to_fp8 calls it.
prefill_attention: 2026-09-25: Flash-attention prefill on contiguous BF16 tensors: Q and O [batch, seq_len, num_q_heads, head_dim], K and V [batch, seq_len, num_kv_heads, head_dim]. sliding_window == 0 disables the window; otherwise keys with q - k >= sliding_window are masked.
prefill_attention_64: 2026-09-25: prefill_attention with 64 query rows per block (32 on the SCALE and HIP targets).
prefill_attention_512_sink: 2026-09-25: prefill_attention for the deepseek-v4-flash attn_prefill_512 kernel (HDIM 512, 16 query rows per block) with a per-head sink logit: sinks is FP32 [num_q_heads], added to the softmax denominator only. DevicePtr::NULL means no sink.
prefill_attention_fp8kv: 2026-09-25: Contiguous prefill with BF16 Q and FP8 E4M3 K and V, which the kernel dequantizes to BF16 in shared memory; 64 query rows per block.
prefill_attention_paged: 2026-09-25: Causal flash-attention prefill of a contiguous Q chunk against K and V in the paged BF16 cache, found through block_table. Query row i is at absolute position q_offset + i; 32 query rows per block.
prefill_attention_paged_64: 2026-09-25: prefill_attention_paged with 64 query rows per block (32 on the SCALE and HIP targets).
prefill_attention_paged_512: 2026-09-25: Causal paged prefill for head_dim == 512 over a BF16 KV cache: 32 query rows and 8 warps per block, with the 101,120 bytes of dynamic shared memory attn_prefill_paged_512.cu requires.
prefill_attention_paged_batched: 2026-09-25: Batched BF16-KV paged prefill attention, 32 query rows per block.
prefill_attention_paged_batched_64: 2026-09-25: Batched BF16-KV paged prefill attention, 64 query rows per block.
prefill_attention_paged_bf16k_turbo2v_64: 2026-09-25: BF16 K with 2-bit V; the arguments of prefill_attention_paged_bf16k_turbo3v_64.
prefill_attention_paged_bf16k_turbo3v_64: 2026-09-25: BF16 K ([block_size, num_kv_heads, head_dim] per block) with 3-bit V.
prefill_attention_paged_bf16k_turbo4v_64: 2026-09-25: BF16 K with 4-bit V; the arguments of prefill_attention_paged_bf16k_turbo3v_64.
prefill_attention_paged_dflash: 2026-09-25: prefill_attention_paged launched with causal_mask_enabled = 0, the BF16-cache counterpart of prefill_attention_paged_fp8_dflash.
prefill_attention_paged_dflash_bf16_indirect: 2026-09-25: prefill_attention_paged_dflash with kv_len, q_offset and q_rope_pos read on the device from kv_len_q_offset_dev (three u32 at byte offsets 0, 4 and 8) instead of passed as scalars, so a captured CUDA graph can replay the launch with values the host writes before each replay. q_offset addresses the cache; q_rope_pos is the absolute query position the causal and sliding-window masks compare against. The kernel applies no RoPE. Kernel: attn_prefill_paged_indirect.
prefill_attention_paged_fp8: 2026-09-25: prefill_attention_paged over an FP8 KV cache.
prefill_attention_paged_fp8_64: 2026-09-25: prefill_attention_paged_fp8 with 64 query rows per block (32 on the SCALE and HIP targets).
prefill_attention_paged_fp8_batched: 2026-09-25: Batched FP8-KV paged prefill attention, 32 query rows per block.
prefill_attention_paged_fp8_batched_64: 2026-09-25: Batched FP8-KV paged prefill attention, 64 query rows per block.
prefill_attention_paged_fp8_dflash: 2026-09-25: prefill_attention_paged_fp8 launched with causal_mask_enabled = 0: the causal mask is skipped, so a query of the DFlash draft block also attends to the block’s later positions.
prefill_attention_paged_fp8k_turbo2v_64: 2026-09-25: FP8 K with 2-bit V.
prefill_attention_paged_fp8k_turbo3v_64: 2026-09-25: FP8 K with 3-bit V.
prefill_attention_paged_fp8k_turbo4v_64: 2026-09-25: FP8 K with 4-bit V.
prefill_attention_paged_nvfp4: 2026-09-25: Causal paged prefill over an NVFP4 KV cache, 32 query rows per block.
prefill_attention_paged_nvfp4_64: 2026-09-25: prefill_attention_paged_nvfp4 with 64 query rows per block (32 on the SCALE and HIP targets).
prefill_attention_paged_nvfp4_batched: 2026-09-25: Batched NVFP4-KV paged prefill attention, 32 query rows per block.
prefill_attention_paged_nvfp4_batched_64: 2026-09-25: Batched NVFP4-KV paged prefill attention, 64 query rows per block, with the arguments of prefill_attention_paged_nvfp4_batched. The caller (qwen3_attention/prefill/paged_attn_batched.rs) picks the 64-row kernels for every KV type when chunk_len >= 256.
prefill_attention_paged_turbo2_64: 2026-09-25: The turbo2 counterpart of prefill_attention_paged_turbo_64.
prefill_attention_paged_turbo3k_turbo8v_64: 2026-09-25: 3-bit K with FP8 V.
prefill_attention_paged_turbo4k_turbo3v_64: 2026-09-25: 4-bit K with 3-bit V.
prefill_attention_paged_turbo4k_turbo8v_64: 2026-09-25: 4-bit K with FP8 V.
prefill_attention_paged_turbo_64: 2026-09-25: Causal paged prefill over a TurboQuant KV cache (same format for K and V), 64 query rows per block (32 on the SCALE and HIP targets). The turbo8, turbo4 and turbo3 _64 entries share this argument list; the caller picks one through kernel and passes that cache’s block stride and data-section size.
prefill_batched_chunk_zero_admitted: 2026-09-28: Whether a batched prefill wave may carry chunk-zero streams: the first-chunk lever (prefill_batched_first_chunk_enabled) or varlen batching (prefill_varlen_enabled). The single rule for the model engine’s admission (kernel_batched_eligible), its paged metadata upload, and the attention layer’s routing; before 2026-09-28 the layer read the lever alone and refused a varlen-admitted chunk-zero wave after admission, which failed every stream in it.
prefill_batched_first_chunk_enabled: 2026-09-25: Whether chunk-zero streams may use the paged batched-prefill path: codispatch on, or METRALE_Q12_BATCHED_FIRST_CHUNK set to 1 or true. The scheduler reads codispatch alone.
prefill_codispatch_enabled: 2026-09-25: Whether cross-request co-dispatch of fresh prompts is on: --prefill-codispatch, else METRALE_PREFILL_CODISPATCH (1 or true, any case); off when neither is set. The scheduler’s levers and prefill_batched_first_chunk_enabled both read it.
prefill_varlen_enabled: 2026-09-25: Whether varlen (ragged) batched prefill is on: --prefill-varlen-batch, else METRALE_PREFILL_VARLEN (1 or true, any case); off when neither is set. The admission predicate (check_kernel_batched_eligible), the batched attention layer and the scheduler’s levers all read this one value.
q2_0_gemv: 2026-09-25: Q2_0 GEMV for M=1 decode: C[1,N] = A[1,K] @ dequant(B)^T. A and C are BF16; B is the raw block_q2_0 buffer, and each weight is dequantized as (code-1)*d inside the dot product.
q2_0_gemv_batchm: 2026-09-25: Batched Q2_0 GEMV: C[M,N] = A[M,K] @ dequant(B)^T, A and C row-major BF16. Each weight block is read once and applied to all m rows.
q2_0_gemv_vec: 2026-09-25: Q2_0 GEMV for M=1 decode: C[1,N] = A[1,K] @ dequant(B)^T. Each lane reads one uint32 of 16 codes, and the CTA stages the activation tile in shared memory. A and C are BF16; B is the raw block_q2_0 buffer; each weight is (code-1)*d.
q2_0_gemv_vec_batchm: 2026-09-25: Batched Q2_0 GEMV for any m: C[M,N] = A[M,K] @ dequant(B)^T, A and C row-major BF16. Each weight word is read once per launch and applied to every staged row.
q2_0_mmq_gemm: 2026-09-25: C[m, n] = A[m, k] x W[n, k] in BF16. a_q8 is the DS4 q8_1 activation from super::quantize_act_q8_1; w_q2_0 is the packed weight as loaded. The tile is 128 x 128 like the Q4_K GEMM, and Q2_0’s x tile row (MMQ_MMA_TILE_X_K_Q8_0) is the same 76 ints as Q4_K’s, so it reuses Q4K_MMQ_SMEM. kernel_wc is used when n is not a multiple of 128.
q2_0_mmq_gemm_packed: 2026-09-25: q2_0_mmq_gemm on a PackedQ2Weight. Returns an error unless w.group == 128, the group size of block_q2_0.
q2_0_weight_bytes: 2026-09-25: Bytes of the packed block_q2_0 form of an [n, k] weight; k must be a multiple of 128.
q4k_mmq_gemm: 2026-09-25: C[m, n] = A[m, k] x W[n, k] in BF16, with a_q8 in q8_1 and w_q4k in block_q4_K, on a 128 x 128 tile. kernel_wc is used when n is not a multiple of 128.
q4k_weight_bytes: 2026-09-25: Bytes of the Q4_K form of an [nrows, n_per_row] weight; n_per_row must be a multiple of 256.
q8_1_scratch_bytes: 2026-09-25: q8_1 activation scratch bytes for [m, k]: 4 bytes per value with k rounded up to 256, plus 1 MiB. ffn_act_q8 in gpu-runtime/src/buffers/sizes.rs uses the same expression with k the larger of the hidden and intermediate sizes.
qsa_block_pool: 2026-09-25: Pool n_new newly complete blocks starting at first_block: the mean of ratio raw keys, then RMSNorm scaled by 1 + w, then rope at the block’s first token position. Returns without launching when n_new == 0.
qsa_gather: 2026-09-25: Copy the selected tokens’ K/V rows into contiguous NHD scratch: slot i holds position sel[i].
qsa_prefill_attn: 2026-09-25: Per (row, q head), attention over exactly the selected blocks in lists plus the incomplete tail, read from the paged cache; the result overwrites that row’s context in attn_out.
qsa_qprep: 2026-09-25: One decode query: per head, RMSNorm scaled by 1 + w, then partial rope at pos; FP32 output.
qsa_qprep_rows: 2026-09-25: qsa_qprep for rows prefill rows at positions first_pos + row, reading q from rows of the indexer projection qk.
qsa_score: 2026-09-25: scores[b] = sum_h relu(q_h . k_b) / sqrt(hd) for n_blocks blocks.
qsa_score_rows: 2026-09-25: Per-row block scores for prefill rows; blocks at or past a row’s complete-block count score -inf.
qsa_score_rows_tc: 2026-09-25: Tensor-core qsa_score_rows: one block of 8 warps covers 16 rows by 64 blocks instead of one block per (row, block). q is split into BF16 high and low halves for the MMA, so the scores are not bit-identical to qsa_score_rows; they feed a top-k.
quant_gemm: 2026-09-25: GEMM (m rows) through the kernel for the weight’s format: NVFP4, FP8 or dense BF16. PackedQ2 weights return an error.
quant_gemv: 2026-09-25: GEMV through the kernel for the weight’s format: NVFP4, FP8 or dense BF16. PackedQ2 weights return an error.
quantize_act_q8_1: 2026-09-25: Quantize BF16 activations [m, k] to q8_1 in out_q8, with k padded to 256. The scale layout (DS4, D4 or D2S6) is the kernel’s.
quantize_bf16_to_fp8: 2026-09-25: Quantize a BF16 weight [N, K] to FP8 E4M3 [N, K] with one f32 scale per row [N], one CTA per row (kernel quantize_bf16_to_fp8 in kernels/gb10/common/dense_gemv_fp8w.cu). DenseWeight::quantize_to_fp8 calls it; the resulting Fp8DenseWeight feeds the row-scaled FP8 kernels such as fp8_gemm_n128_row_scaled.
quantize_bf16_to_nvfp4: 2026-09-25: Quantize a BF16 [M, K] matrix to NVFP4 with scale2 = 1.0: packed E2M1 [M, K/2] and one E4M3 scale per 16 elements [M, K/16]. One block per row; the kernel’s N parameter is the row count. Produces the W4A4 prefill activations.
quantize_weight_q4k: 2026-09-25: Quantize BF16 weights [nrows, n_per_row] to block_q4_K, one thread per 256-value block.
read_expert_ptrs_u64: 2026-09-25: One blocking device-to-host copy of an [n] u64 pointer table.
read_expert_scales_f32: 2026-09-25: One blocking device-to-host copy of an [n] f32 scale table.
requant_w_nvfp4_int8: 2026-09-25: Requantize an NVFP4 weight (packed E2M1, per-16 E4M3 block scales, per-tensor scale2) to int8 with one F32 scale per 32 elements, for int8_gemm_faith2. Reads W_packed[N, K/2] and W_e4m3[N, K/16]; writes W_i8[N, K] (signed int8) and W_scale[N, K/32] (F32). The dense FFN’s ensure_int8_weight runs it on a weight’s first int8 prefill and keeps the result.
reshape_and_cache: 2026-09-25: Copy BF16 K and V rows into the paged BF16 cache at slot_mapping (kernel reshape_and_cache_flash), one block per token. key_stride / value_stride are the row strides in elements. _cache_stride is not passed to the kernel, which derives the block stride from block_size.
reshape_and_cache_bf16k_turbo2v: 2026-09-25: Write K and V to the BF16-K / Turbo2V cache (kernel reshape_and_cache_flash_bf16k_turbo2v): K as raw BF16, V as 2-bit Lloyd-Max codes plus one FP8 scale per 16 elements, 2.5 bits per V element.
reshape_and_cache_bf16k_turbo3v: 2026-09-25: Write K and V to the BF16-K / Turbo3V cache (kernel reshape_and_cache_flash_bf16k_turbo3v): K as raw BF16, V as 3-bit Lloyd-Max codes (3/8 byte per element) plus one FP8 scale per 16 elements. v_data_section_bytes is the size of a V block’s code section, after which its scales start.
reshape_and_cache_bf16k_turbo4v: 2026-09-25: Write K and V to the BF16-K / Turbo4V cache (kernel reshape_and_cache_flash_bf16k_turbo4v): K as raw BF16, V as 4-bit Lloyd-Max codes plus one FP8 scale per 16 elements.
reshape_and_cache_flash_v_only: 2026-09-25: Copy only V into the paged BF16 cache (kernel reshape_and_cache_flash_v_only), for callers whose K is written by a fused_k_norm_rope_cache_write_* launch; a reshape_and_cache there would overwrite that K.
reshape_and_cache_fp8: 2026-09-25: Quantize BF16 K and V rows into the paged FP8 cache (kernel reshape_and_cache_flash_fp8): fp8 = sat(bf16 / scale) with the dequant scales k_scale / v_scale. k_cache / v_cache are the pool base pointers, and cache_stride is the stride between cache blocks in elements (normally block_size * num_kv_heads * head_dim).
reshape_and_cache_fp8k_turbo2v: 2026-09-25: Write K and V to the FP8-K / Turbo2V cache (kernel reshape_and_cache_flash_fp8k_turbo2v).
reshape_and_cache_fp8k_turbo3v: 2026-09-25: Write K and V to the FP8-K / Turbo3V cache (kernel reshape_and_cache_flash_fp8k_turbo3v): K as FP8 E4M3, V as 3-bit Lloyd-Max codes plus one FP8 scale per 16 elements. One 256-thread block per token.
reshape_and_cache_fp8k_turbo4v: 2026-09-25: Write K and V to the FP8-K / Turbo4V cache (kernel reshape_and_cache_flash_fp8k_turbo4v).
reshape_and_cache_nvfp4: 2026-09-25: Write each token’s K and V into the paged NVFP4 cache at its slot_mapping slot: E2M1 values with one E4M3 scale per 16. One block per token.
reshape_and_cache_turbo3k_turbo8v: 2026-09-25: Write K and V to the Turbo3K / Turbo8V cache (kernel reshape_and_cache_flash_turbo3k_turbo8v).
reshape_and_cache_turbo4k_turbo3v: 2026-09-25: Write K and V to the Turbo4K / Turbo3V cache (kernel reshape_and_cache_flash_turbo4k_turbo3v).
reshape_and_cache_turbo4k_turbo8v: 2026-09-25: Write K and V to the Turbo4K / Turbo8V cache (kernel reshape_and_cache_flash_turbo4k_turbo8v).
residual_add: 2026-09-25: In-place BF16 residual[i] += src[i] (kernel bf16_residual_add).
residual_add_rms_norm: 2026-09-25: hidden += src, then output = rms_norm(hidden) and residual = hidden, in one kernel, one block per token row.
residual_add_rms_norm_exact: 2026-09-28: residual_add_rms_norm_exact (residual_add_rms_norm_exact.cu): the arguments and launch of residual_add_rms_norm, with the bytes of bf16_residual_add followed by rms_norm_residual: the sum of squares reads the BF16-rounded sums. The circuit executor launches it across a layer boundary, with the next layer’s input-norm weight.
residual_add_rms_norm_gatef32: 2026-09-25: residual_add_rms_norm that also writes the normed row in FP32 to output_f32, the MoE router input, so routing does not see the BF16 rounding of the norm output. Callers use it when METRALE_FP32_ROUTING is active for the layer’s FFN.
rms_norm: 2026-09-25: RMS norm of num_tokens packed rows of hidden_size, one block per row. How weight scales the result is the kernel’s: rms_norm in kernels/gb10/common/rms_norm.cu multiplies by 1 + weight, and the rms_norm_vanilla kernels by weight.
rms_norm_residual: 2026-09-25: output = rms_norm(input) and residual = input in one pass, one block per token row.
rms_norm_short_row_eligible: 2026-09-25: Gate for rms_norm_warp_row: hidden_size <= 256 and even, and num_rows >= 1024. METRALE_RMS_NORM_WARP_ROW=0 turns it off; the variable is read once per process.
rms_norm_strided: 2026-09-25: RMS norm of num_groups groups of rows_per_group rows in one launch. Groups start row_stride elements apart, and rows inside a group are packed at hidden_size. The multi-sequence q/k head norms use it because each sequence’s heads sit inside that sequence’s own interleaved block.
rms_norm_warp_row: 2026-09-25: RMS norm with one warp per row and 8 rows per block, so the reduction needs no shared memory or block barrier. The prefill per-head q_norm/k_norm uses it when rms_norm_short_row_eligible holds (qwen3_attention/prefill/cache_skip.rs).
rope: 2026-09-25: Apply rotary position embeddings to Q and K in place, rows packed (num_q_heads * head_dim and num_kv_heads * head_dim elements apart). positions is a device u32[seq_len].
rope_mrope_interleaved: 2026-09-25: Interleaved multimodal RoPE (MRoPE) with three position streams (pos_t, pos_h, pos_w); rotary pair i takes its position from stream i % 3. With the same pointer for all three, the math is that of rope.
rope_mrope_interleaved_k_only: 2026-09-25: rope_mrope_interleaved for K only, for when Q was already rotated by a fused prefill kernel.
rope_proportional: 2026-09-25: Proportional RoPE (Gemma-4 full-attention layers).
rope_strided: 2026-09-25: rope for rows that are not packed: Q rows q_row_stride and K rows k_row_stride elements apart, all num_tokens rows in one launch. rope_forward_strided differs from rope_forward only in the row address, so the packed strides give the same result as rope.
rope_yarn: 2026-09-25: RoPE with frequencies from a precomputed inv_freq table ([rotary_dim/2] FP32); rope_forward_yarn ignores theta.
rope_yarn_scaled: 2026-09-25: RoPE from an inv_freq table with cosine and sine scaled by attention_factor.
scaled_add: 2026-09-25: In-place BF16 output[i] += scale * src[i] (kernel bf16_scaled_add).
set_prefill_codispatch_from_cli: 2026-09-25: Publish the command line’s --prefill-codispatch decision. Returns the value in force, which differs from enabled when the cell was already resolved; the command line then did not take effect. The serve calls this only when the flag is given, so without it the METRALE_PREFILL_CODISPATCH fallback applies.
set_prefill_varlen_from_cli: 2026-09-25: Publish the command line’s --prefill-varlen-batch decision. Returns the value in force, which differs from enabled when the cell was already resolved; the command line then did not take effect. The serve calls this only when the flag is given, so without it the METRALE_PREFILL_VARLEN fallback applies.
sigmoid_blend: 2026-09-25: In-place BF16 output[i] += sigmoid_gate * src[i], where sigmoid_gate is a host value the caller has already passed through the sigmoid. The arguments are those of kernel bf16_sigmoid_blend; nothing in the tree calls this function.
sigmoid_blend_device: 2026-09-25: output[i] += sigmoid(*gate_ptr) * src[i], where gate_ptr is one BF16 scalar in device memory (bf16_sigmoid_blend_device in kernels/gb10/common/residual_add.cu).
sigmoid_gate_mul: 2026-09-25: output[i] = input[i] * sigmoid(gate[i]) (kernel sigmoid_gate_mul).
sigmoid_gate_mul_batched: 2026-09-25: sigmoid_gate_mul over num_tokens rows in one launch: output[t, d] = input[t, d] * sigmoid(gate[t * gate_stride + d]). input and output are contiguous [num_tokens, dim].
sigmoid_gate_mul_head_broadcast: 2026-09-25: output[t, h, d] = input[t, h, d] * sigmoid(gate[t, h]): one BF16 gate value per head, [num_tokens, nq], broadcast over the head’s hd elements (kernel sigmoid_gate_mul_head_broadcast).
silu_mul: 2026-09-25: output = act(gate) * up over num_elements BF16 values, where kernel fixes the activation: callers pass moe_silu_mul (SiLU) or, for a GeLU FFN, gelu_mul. The moe_silu_mul.cu copies in gb10/deepseek-v4-flash (also compiled for longcat-flash-lite) and gb10/step3p7-flash first clamp gate to at most 10 and up to [-10, 10].
silu_mul_quant_fp8: 2026-09-25: SiLU(gate) * up quantized to FP8 E4M3 with one scale per 128-element group of each row, in one kernel (silu_mul_quant_fp8); the FP8 MoE prefill uses it before the down projection (moe/forward_prefill_fp8.rs).
silu_mul_strided: 2026-09-25: SiLU(gate) * up over row-strided BF16 operands (kernel silu_mul_strided): row r reads gate and up at r * in_stride + c for c < cols and writes output at r * out_stride + c. The fused dense-FFN gate+up projection (layers/dense_ffn_gateup_fused.rs) passes the two column halves of one [rows, 2 * cols] matrix (up = gate.offset(cols * 2)) and a contiguous output. One grid row per matrix row.
softplus_gate_mul_head_broadcast: 2026-09-25: Like sigmoid_gate_mul_head_broadcast with softplus in place of sigmoid (kernel softplus_gate_mul_head_broadcast).
ssd_scan_fits: 2026-09-25: Whether the SSD chunked scan fits in shared memory for this state_size.
ssd_scan_smem: 2026-09-25: Dynamic shared memory mamba2_ssd_scan requests for a given SSM state size. It grows linearly in state_size, so a large SSM state can exceed MAX_DYNAMIC_SMEM. The launch calls this function, so the two cannot drift.
ssm_ba_gates_hopper_enabled: 2026-09-25: The lever: the target’s [defaults] ssm_ba_gates_hopper, which METRALE_SSM_BA_GATES_HOPPER overrides (0, false, off or no turn it off, any other value on; super::target_defaults). There is no METRALE_NO_* spelling for it.
ssm_ba_gates_hopper_reject: 2026-09-25: Why the twin does not run, as one of BA_GATES_REJECTS; None means it runs. Pure, so tests cover it without a GPU. The shape guards:
ssm_h_state_f16_to_f32: 2026-09-25: FP16 -> FP32 widening of n h-state elements (ssm_h_state_f16_to_f32). Grid-stride; src and dst must not alias.
ssm_h_state_f32_to_f16: 2026-09-25: FP32 -> FP16 conversion of n h-state elements (ssm_h_state_f32_to_f16, round-to-nearest-even). The kernel is grid-stride, so the clamped grid still covers every element. src and dst must not alias: an in-place narrowing conversion is a data race.
strided_out_extent_elems: 2026-09-25: Elements from the start of the output to one past the last element a cuBLASLt GEMM writes when its m_pad output rows of n elements are ldc elements apart: (m_pad - 1) * ldc + n. Pass the padded row count, because the pad rows are written too. The strided call sites bound their buffers with it.
transpose_block_scale: 2026-09-25: Transpose FP32 block scales: [N/128, K/128] to [K/128, N/128].
transpose_fp8: 2026-09-25: Transpose an FP8 weight on the GPU: B[N, K] to B_t[K, N].
transpose_u8: 2026-09-25: Single-matrix uint8 transpose [rows, cols] -> [cols, rows] on the GPU (transpose_u8.cu, 32x32 shared-memory tiles). The quantized weight transposes in weight_map/quantized/transpose.rs call it.
use_gemv_sw: 2026-09-25: Use the single-warp kernel only when the lever is on and its handle resolved (a missing kernel is KernelHandle(0)).
w4a4_gemm: 2026-09-25: W4A4 GEMM on FP4 tensor cores (kernel w4a4_gemm). The activation comes from quantize_bf16_to_nvfp4 and is launched with scaleA2 = 1.0; the weight is an NVFP4 QuantizedWeight; the output is BF16 [M, N].
w4a4_gemm_mfast: 2026-09-25: w4a4_gemm through the w4a4_gemm_mfast kernel, with M on the fast grid axis (blockIdx.x = M block), so consecutive CTAs share a B panel.
w4a16_decode_gemv: 2026-09-25: Decode GEMV: the single-warp kernel when use_gemv_sw allows it, otherwise the base w4a16_gemv, each with its own grid.
w4a16_gemm: 2026-09-29: W4A16 GEMM, C = A @ dequant(B): A [M, K] BF16, B NVFP4 (packed E2M1, FP8 block scales, FP32 per-tensor scale), C [M, N] BF16. Every launcher here sizes grid.x from the kernel’s published N tile (n_tile_blocks).
w4a16_gemm_n128: 2026-09-29: W4A16 GEMM over a packed transposed B, N tiled by the kernel’s published tile: it launches through w4a16_gemm_n128_ldb with ldb = n.
w4a16_gemm_n128_ldb: 2026-09-25: w4a16_gemm_n128 with an explicit row stride ldb for the transposed B, so row r starts at r * ldb.
w4a16_gemm_n128_m128: 2026-09-25: W4A16 GEMM with a 128-row CTA (two 64-row chunks): each CTA loads a B tile once for both chunks.
w4a16_gemm_n128_m128_bf16: 2026-09-25: W4A16 GEMM launching w4a16_gemm_t_m128_bf16, which runs BF16 m16n8k16 MMAs, on the grid, block and transposed NVFP4 weight layout of w4a16_gemm_n128_m128. It passes 8 arguments, so it fits only the 8-parameter v1 kernel; w4a16_gemm_t_m128_bf16_v2 takes a 9th (ldb) and goes through w4a16_gemm_n128_m128_bf16_ldb (ldb = N for an unpadded twin).
w4a16_gemm_n128_m128_bf16_ldb: 2026-09-25: w4a16_gemm_n128_m128_bf16 with an explicit transposed-B row stride, for the same reason as w4a16_gemm_n128_ldb. This is the launcher for the 9-parameter w4a16_gemm_t_m128_bf16_v2.
w4a16_gemm_n128_m128_v2: 2026-09-25: W4A16 GEMM v2 (w4a16_gemm_t_m128_v2, in the minimax-m2-229b and qwen3.6-27b kernel dirs): the CTA tile of w4a16_gemm_n128_m128 (M=128, N=128, K step 32) on 256 threads instead of 128. Warps 0-3 compute the first 64 rows and warps 4-7 the second, so the two chunks’ MMAs run in parallel.
w4a16_gemm_n128_m128_v3: 2026-09-25: W4A16 GEMM v3 (w4a16_gemm_t_m128_v3, only in kernels/gb10/minimax-m2-229b/nvfp4/w4a16_gemm_v3.cu): a 128x128 CTA tile on 256 threads with a K step of 64.
w4a16_gemv: 2026-09-25: C[1, N] = A[1, K] @ dequant(B) with BF16 A and C and an NVFP4 weight.
w4a16_gemv_batch2: 2026-09-25: w4a16_gemv for 2 contiguous rows in one pass over the weight. gemv_tc::tc_fixed_m takes the launch to the tensor-core kernel when it routes it (METRALE_W4A16_TC_WIDE); otherwise kernel runs.
w4a16_gemv_batch3: 2026-09-25: w4a16_gemv_batch2 for 3 rows.
w4a16_gemv_batchm: 2026-09-25: w4a16_gemv for m contiguous rows in one pass over the weight. kernel is the CUDA-core w4a16_gemv_batch{M} tier the caller picked (W4a16BatchmTiers::kernel); gemv_tc::tc_kernel replaces it with the tensor-core w4a16_gemv_tc8/tc16 when it routes the shape.
w4a16_gemv_dual: 2026-09-25: Two W4A16 GEMVs over the same BF16 input in one launch; blockIdx.z picks weight1 or weight2. Both projections have n outputs.
w4a16_gemv_dual_batch2: 2026-09-25: w4a16_gemv_dual_batch3 for 2 rows.
w4a16_gemv_dual_batch3: 2026-09-25: Two NVFP4 GEMVs over the same 3-row input [3, K]; blockIdx.z picks weight0 or weight1, and each writes [3, N].
w4a16_gemv_dual_sw: 2026-09-25: One warp per output variant of w4a16_gemv_dual, 8 outputs per block. It shares the per-lane partial w4a16_dual_partial with the base kernel (w4a16_gemv_fused.cu). Callers pick it through use_gemv_sw with ModelLevers::gemv_sw, which is on unless METRALE_NO_GEMV_SW=1.
w4a16_gemv_grid_x
w4a16_gemv_qg: 2026-09-25: w4a16_gemv over a weight whose rows interleave Q and Gate per head; the output is stored deinterleaved, [Q of all heads | Gate of all heads], with n = num_heads * head_dim * 2.
w4a16_gemv_qg_batch2: 2026-09-25: w4a16_gemv_qg for 2 rows: A is [2, K], C is [2, N], each row deinterleaved.
w4a16_gemv_qg_batch3: 2026-09-25: w4a16_gemv_qg for 3 rows: A is [3, K], C is [3, N], each row deinterleaved.
w4a16_gemv_qkvz: 2026-09-25: w4a16_gemv over a weight whose rows interleave Q, K, V and Z per group; the output is stored deinterleaved as [Q | K | V | Z].
w4a16_gemv_silu_input: 2026-09-25: W4A16 GEMV whose activation is silu(gate_out) * up_out, computed per element from the two [k] BF16 vectors inside the kernel.
w4a16_gemv_silu_input_sw: 2026-09-25: One warp per output variant of w4a16_gemv_silu_input, 8 outputs per block, picked the same way as w4a16_gemv_dual_sw.
w4a16_gemv_sw: 2026-09-25: Single-warp-per-output W4A16 GEMV for M=1, grid ceil(N/8).
w4a16_gemv_sw_grid_x
w4a16_gemv_sw_raw: 2026-09-25: The launch of w4a16_gemv_sw for callers that hold the NVFP4 packed weight, scale and global scale as separate values rather than a QuantizedWeight (the GLM-5.3 MLP in model-arch/src/glm5next_mlp/forward.rs). It keeps the grid on w4a16_gemv_sw_grid_x.
w4a16_t_via_fp8_ldmab: 2026-09-28: w4a16_gemm_t_m128’s product on the multistage FP8 GEMM: the transposed NVFP4 weight ([K/2, N] codes, [K/16, N] E4M3 scales, scalar weight_scale_2) is dequantized to E4M3 [N, K] with that kernel’s arithmetic (fp8_predequant_nvfp4_t), then fp8_act_ldmab_gemm runs. Bit-identical to w4a16_gemm_t_m128 (kbench w4a16_bench.cu, nine shapes). k must be a multiple of 32.
w4a16_tc_rows: 2026-10-02: output [m, ldc] = input [m, lda] x W^T for the row-major NVFP4 weight [n, k] (packed E2M1, E4M3 scales of 16, weight_scale_2). Refuses a shape w4a16_tc_rows_shape_ok refuses.
w4a16_tc_rows_shape_ok: 2026-10-02: The kernel’s shape contract, without a GPU: 1..=64 rows, any positive N (the entry points are ragged: a partial last CTA loads zero weight rows and stores nothing past N), K a positive multiple of 256 (whole load groups of every entry point), and an A pitch that keeps rows 16-byte aligned and covers K; the C pitch covers N.
w8a8_act_quant: 2026-09-28: Quantize rows rows of the BF16 activation x ([rows, ldx], the first k columns) into scratch for scale’s layout.
w8a8_act_quant_silu: 2026-09-28: w8a8_act_quant of bf16(silu(gate) * up), computed in the same launch with the arithmetic of moe_silu_mul: gate and up are [rows, ld] BF16, the first k columns read.
w8a8_decode_available: 2026-09-28: The “W8A8 available for this layer and shape” hook: the kernels of w’s layout resolved, rows is in 1..=W8A8_MAX_ROWS, and the scratch holds the quantized activation. The declared-precision policy decides whether W8A8 is wanted; this says whether it can run.
w8a8_decode_proj_disabled: 2026-09-25: Kill switch: METRALE_NO_W8A8_DECODE_PROJ set to any value, empty and 0 included, turns this route off. Read once, so it is constant for the process and safe to branch on under CUDA-graph capture.
w8a8_gemv: 2026-09-28: The GEMV over the activation w8a8_act_quant left in scratch: out[m, 0..n] (row pitch ldc BF16 elements) for m < rows.
w8a8_proj: 2026-09-28: One W8A8 projection: quantize x ([rows, ldx] BF16) and multiply, writing out ([rows, ldc] BF16, the first w.n() columns).
w8a16_gemm: 2026-09-25: W8A16 GEMM for M>1: C[M,N] = A[M,K] @ dequant(B[N,K])^T, with the 2D block scale [ceil(N/128), ceil(K/128)] FP32.
w8a16_gemm_m16: 2026-09-25: Contiguous input [m, k] BF16 and output [m, n] BF16; weight / block_scale are the raw w8a16_gemv pointers ([N, K] FP8 E4M3 and [N/128, K/128] FP32).
w8a16_gemm_m16_n64: 2026-09-25: The N_TILE=64 instantiation of w8a16_gemm_m16: the same arguments and guards, ceil(N/64) CTAs instead of ceil(N/32).
w8a16_gemm_m16_strided: 2026-09-25: Strided w8a16_gemm_m16: a_row_stride / c_row_stride are the A and C row pitches in elements, for callers whose rows are not contiguous. The multi-seq decode QKV buffer (qkv_fp8_batch.rs) is [n, per_seq_qkv] with Q at offset 0, K after Q and V after K in every row, so one launch per projection writes all m rows into their slots. Same argument order as w8a16_gemv_batch16_strided.
w8a16_gemm_n128_m128: 2026-09-25: Transposed W8A16 GEMM with a 128x128 MxN tile (kernel w8a16_gemm_t_m128): two 64-row chunks, 8 warps, m16n8k16 BF16 MMA, and the block scale folded into an FP32 accumulator once per 128-K block (kernels/gb10/common/w8a16_gemm_t_m128.cu). Same arguments and layout as w8a16_gemm_t, so it takes the output of the transpose_fp8 / transpose_block_scale kernels unchanged.
w8a16_gemm_pipelined: 2026-09-25: Pipelined W8A16 GEMM (kernel w8a16_gemm_pipelined), with the same arguments and block-scale layout as w8a16_gemm. A 128x32 MxN tile with 256 threads and a cp.async prefetch (PM_M_TILE, PM_N_TILE, PM_THREADS in kernels/gb10/common/w8a16_gemm_pipelined.cu).
w8a16_gemm_pipelined_by_m: 2026-09-25: Contiguous block-scaled W8A16 GEMM, tile chosen by w8a16_pipelined_prefers_m32: the 32-tile twin on m32_kernel, else (2026-09-26) w8a16_pipelined_prefers_m64’s 64-tile twin on m64_kernel, else the 128-tile w8a16_gemm_pipelined on full_kernel. The GDN in_proj_qkvz / out_proj batched-verify arms (trait_decode_batched.rs) call it.
w8a16_gemm_pipelined_m32: 2026-09-25: Contiguous [M, K] -> [M, N] launch of the same kernel: the pitches are k and n. Same signature as super::w8a16_gemm_pipelined.
w8a16_gemm_pipelined_m64: 2026-09-26: Contiguous launch of w8a16_gemm_pipelined_m64 (pitches k and n). Same signature as super::w8a16_gemm_pipelined.
w8a16_gemm_pipelined_m32_strided: 2026-09-25: Strided launch of w8a16_gemm_pipelined_m32.
w8a16_gemm_pipelined_m64_strided: 2026-09-26: Strided launch of w8a16_gemm_pipelined_m64: the arguments and checks of w8a16_gemm_pipelined_m32_strided, grid (ceil(N/32), ceil(M/64), 1), block 256. The multi-seq attention Q/K/V and O arms take it above 32 rows under ModelLevers::fp8_attn_m32.
w8a16_gemm_t: 2026-09-25: Transposed W8A16 GEMM: C[M,N] = A[M,K] @ dequant(B_t[K,N]), with the transposed FP32 block scale block_scale_t[K/128, N/128], so weight reads are contiguous along N. Each CTA covers 64 rows and 64 columns with 128 threads (kernels/gb10/common/w8a16_gemm_t.cu, kernels/strix-hip/common/w8a16_gemm_t.cu).
w8a16_gemm_t_pipelined: 2026-09-25: Pipelined transposed W8A16 GEMM (kernel w8a16_gemm_t_pipelined), with the same arguments as w8a16_gemm_t: a 128x32 MxN tile with 256 threads (PT_M_TILE, PT_N_TILE, PT_THREADS in kernels/gb10/common/w8a16_gemm_t.cu).
w8a16_gemv: 2026-09-25: W8A16 GEMV for M=1 decode: C[1,N] = A[1,K] @ dequant(B[N,K])^T.
w8a16_gemv_batch2: 2026-09-25: Launches a batch-2 block-scaled FP8 GEMV. No kernel under kernels/ defines a w8a16_gemv_batch2 entry point.
w8a16_gemv_batch4: 2026-09-25: Block-scaled FP8 GEMV for up to 4 rows (w8a16_gemv_batchm_impl<4>); one pass over the weight serves every row. input [M, K] BF16, output [M, N] BF16; weight [N, K] FP8 E4M3 and block_scale [N/128, K/128] FP32, as for w8a16_gemv.
w8a16_gemv_batch4_strided: 2026-09-25: Strided variant of w8a16_gemv_batch4 (up to 4 rows). It reads and writes rows at caller-given pitches, so one projection of the multi-sequence decode QKV buffer [n, per_seq_qkv] is written for every row in one launch.
w8a16_gemv_batch16: 2026-09-25: The 16-row instantiation of the same template (kernel w8a16_gemv_batch16, module w8a16_gemv_batch4), with the launch geometry and per-row numerics of w8a16_gemv_batch4.
w8a16_gemv_batch16_ncol2: 2026-09-25: N_COLS=2, contiguous A [M, K] and C [M, N].
w8a16_gemv_batch16_ncol4: 2026-09-25: N_COLS=4, contiguous A and C. Grid: (ceil(N/16), 1, 1)
w8a16_gemv_batch16_ncol2_strided: 2026-09-25: N_COLS=2, explicit A and C row pitches in elements, for the multi-seq [n, per_seq_qkv] decode buffer, as w8a16_gemv_batch16_strided.
w8a16_gemv_batch16_ncol4_strided: 2026-09-25: N_COLS=4, explicit row pitches.
w8a16_gemv_batch16_strided: 2026-09-25: The 16-row variant of w8a16_gemv_batch4_strided (kernel w8a16_gemv_batch16_strided, module w8a16_gemv_batch4), with the same launch geometry.
w8a16_gemv_dual: 2026-09-25: Two FP8 E4M3 GEMVs over the same BF16 input in one launch; blockIdx.z picks projection 1 or 2. Both have n outputs. Each weight is [n, k] bytes with an FP32 [n/128, k/128] block-scale buffer.
w8a16_gemv_row_tiered: 2026-09-27: One row’s block-scaled W8A16 projection [1, k] x [n, k]^T for a single-row decode site: under RowTiers::Canonical, when m32_kernel is linked and k is whole 128-wide scale blocks, the 32-row tile (the order every other row count takes under that policy); otherwise the scalar w8a16_gemv.
w8a16_gemv_silu_input: 2026-09-25: FP8 E4M3 GEMV whose activation is silu(gate_out) * up_out, computed per element from the two [k] BF16 vectors inside the kernel. block_scale is FP32 [n/128, k/128].
w8a16_m16_smem_bytes: 2026-09-27: Dynamic shared memory of a skinny block for k: PM16_SMEM_BYTES(K), the block’s 8 weight rows k + 16 bytes apart.
w8a16_pipelined_prefers_m32: 2026-09-25: The selection rule between the 128-tile w8a16_gemm_pipelined and its 32-tile twin at a contiguous call site: the twin takes 1..=32 rows when its handle is nonzero and K is a positive multiple of 128. Above 32 rows the twin would need grid.y > 1, streaming the weight once per 32-row tile. w8a16_gemm_pipelined_by_m applies it.
w8a16_pipelined_prefers_m64: 2026-09-26: 33..=64 rows take the 64-row twin w8a16_gemm_pipelined_m64 when its handle is nonzero and K is a positive multiple of 128: one weight pass for a whole k = 2 or 3 batched verify at C = 16, where the 32-row twin would stream the weight per tile and the 128-row tile is three-quarters empty. The three kernels give every output the same sub-MMAs in the same order (native_fp8_gdn_proj_m32_microtest compares them byte for byte).
w8a16_tc_rows: 2026-09-28: The launch itself, with no policy check: input [m, a_row_stride] BF16, weight [n, k] E4M3, block_scale [n / 128, k / 128] FP32, output [m, c_row_stride] BF16. Refuses a shape w8a16_tc_rows_shape_ok refuses.
w8a16_tc_rows_launch: 2026-09-28: Launches [m, k] x [n, k]^T through w8a16_tc_rows_{16,32,64} and returns true when the lever is on, the process runs RowTiers::Canonical, the shape passes w8a16_tc_rows_shape_ok and the target ships the module; otherwise returns false and launches nothing.
w8a16_tc_rows_shape_ok: 2026-09-28: The kernel’s shape contract, without a GPU: at least one row, N a positive multiple of 128 (whole scale blocks, whole CTAs), K a positive multiple of 128, and an A pitch that keeps rows 16-byte aligned and covers K; the C pitch covers N.
wide_prefill_kernel: 2026-09-25: The head_dim > 256 prefill kernel and the BR its grid must be built for. qwen3_attention::init takes the handle from here and prefill_attention takes the BR from here, so the two cannot disagree; a grid built for the wrong BR would not fail, it would cover the wrong q-tiles.
widen_block_scale_f32: 2026-09-25: Widen an FP8 block-scale tensor to FP32 on the GPU, so the block-scale kernels read const float*. src is [total] BF16 (dtype 0), FP32 (1) or E8M0 (2); dst is [total] FP32. An E8M0 byte e becomes 2^(e - 127) (the bit pattern e << 23), e = 0 becomes 2^-127 and e = 255 becomes 0.0.

Type Aliases§

ContiguousBatchGemv: 2026-09-25: The signature of the contiguous batched GEMV launchers, such as w8a16_gemv_batch4 and w8a16_gemv_batch16, so a caller that picks a tier by row count can hold the launcher and its handle as one pair.
ContiguousM16Gemm: 2026-09-25: The shared signature of w8a16_gemm_m16 and w8a16_gemm_m16_n64, so m16_tc_kernel (dense_ffn_m16_tc.rs) can return either as one function pointer.
DenseM16Bf16Gemm: 2026-09-25: The signature both instantiations share, so a caller can pick one and hold a single function pointer.

Module ops

Module ops Copy item path

Re-exports§

Modules§

Structs§

Enums§

Constants§

Functions§

Type Aliases§

Module ops