Struct Glm5NextMlpKernels 

pub struct Glm5NextMlpKernels {

    pub gemm: KernelHandle,
    pub gemm_f32: KernelHandle,
    pub gemv: KernelHandle,
    pub gemv_f32: KernelHandle,
    pub gemv_batchm: KernelHandle,
    pub w4a16: KernelHandle,
    pub w4a16_gemv: KernelHandle,
    pub w4a16_gemv_sw: KernelHandle,
    pub w4a16_gemv_sw_moe: KernelHandle,
    pub w4a16_gemv_sw_moe_batchm: [KernelHandle; 7],
    pub moe_row_union: KernelHandle,
    pub swiglu: KernelHandle,
    pub router: KernelHandle,
    pub combine: KernelHandle,
}

Every kernel a GLM MLP site launches.

Resolved with kernel() (not try_kernel): a missing entry point is a hard error, never a silent fallback onto moe_silu_mul, whose SwiGLU does not clamp.

Fields

gemm: KernelHandle

C = A @ B^T, BF16 out — dense FFN, shared expert.

gemm_f32: KernelHandle

Same, FP32 out. 🪤 The router GEMM MUST take this one: glm5next_router_topk reads const float* logits, while gate.weight is BF16 on disk.

gemv: KernelHandle

M=1 twins of gemm / gemm_f32. gemv_f32 may be a 0 handle on a target that predates dense_gemv_bf16_fp32out; gemm() falls back to the tile arm then.

gemv_f32: KernelHandle

gemv_batchm: KernelHandle

🔴 dense_gemv_bf16_batchm — 2 ..= 8 rows in ONE weight sweep. The shared-expert and dense-FFN projections are pure weight streaming, so this is what stops a K-token verify from re-reading them K times. 0 = unavailable, falls back to the tile GEMM.

w4a16: KernelHandle

NVFP4 C = A @ B^T, tile GEMM. Kept for any M > 1 caller; the decode path must not use it — see W4A16_GEMV_MODULE.

w4a16_gemv: KernelHandle

NVFP4 C[1, N] = A[1, K] @ B[N, K]^T — the M=1 decode kernel.

🪤 Its grid is COUPLED to the kernel’s N_PER_BLOCK; use ops::gemv_sw::w4a16_gemv_grid_x, never a hand-written div_ceil.

w4a16_gemv_sw: KernelHandle

Single-warp-per-output sibling of w4a16_gemv, bit-identical to it (examples/w4a16_gemv_sw_microtest.rs): same w4a16_gemv_partial per orig-lane, 8 outputs per 256-thread block instead of 4, and no cross-warp __syncthreads() + shared-memory round trip. try_kernel — a target without it falls back to the base kernel rather than failing to load.

🪤 Its grid is ceil(N/8), NOT ceil(N/4). Dispatch through ops::w4a16_decode_gemv, which couples the two; swapping the kernel without swapping the grid writes half the outputs.

w4a16_gemv_sw_moe: KernelHandle

Grouped MoE sibling of w4a16_gemv_sw: all top_k slots in ONE launch, expert weights reached through a device pointer table indexed by the router’s on-device ids. Bit-identical per slot — same w4a16_gemv_partial, same shuffle tree.

🪤 Grid is (ceil(N/8), top_k, 1). try_kernel — a target without it falls back to the host-dispatch loop.

w4a16_gemv_sw_moe_batchm: [KernelHandle; 7]

Row-batched sibling of w4a16_gemv_sw_moe, indexed [rows - 2] for rows 2..=8: the UNION of the selected experts over a call’s rows, each swept ONCE. An expert two rows both picked costs one weight read here and two in the per-row path — 14% of the routed traffic at K=2, 22% at K=3 (measured union 8.00/13.74/18.76/23.35 at K=1..4).

🔴 Widened from 2..=4 to 2..=8 on 2026-08-31. The old stop at 4 was the compiled tier family, not a limit of the union: Self::moe_row_union resolves rows * top_k ids in ONE 64-thread block, and GLM-5.3 is 8 * 8 == 64 exactly. It matters because the batched prefill sub-chunk is 8 rows wide (ANOMALIES A65) and the routed experts were the only stage of it still paying per row.

🪤 grid.y is the UNION entry, not the slot, and its extent is rows * top_k; the entries the routing did not fill retire on u_eid < 0. Needs Self::moe_row_union.

🪤 Register cost rises with the tier — measured ptxas -v, sm_121a, no spills at any width: 43/40/48/70/72/80/80 registers at R = 2..8. R = 8 at 80 regs / 256 threads is 3 CTAs/SM against R = 4’s 5, so a WIDER tier is not free; it wins only when the union actually shrinks the expert sweeps.

moe_row_union: KernelHandle

Builds the union table the batched kernel indexes. One block, rows * top_k threads.

🪤 rows * top_k MUST be <= 64: it is a single block and threads past it never run, so an over-wide call silently drops union entries. The caller gates on it.

swiglu: KernelHandle

🪤 Clamped SwiGLU, asymmetric. Not moe_silu_mul.

router: KernelHandle

combine: KernelHandle

Implementations

impl Glm5NextMlpKernels

pub fn resolve(gpu: &dyn GpuBackend) -> Result<Self>

Trait Implementations

impl Clone for Glm5NextMlpKernels

fn clone(&self) -> Glm5NextMlpKernels

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Copy for Glm5NextMlpKernels