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,
    pub moe_sort_by_expert: KernelHandle,
    pub moe_grouped_gemm: KernelHandle,
    pub combine_indexed: KernelHandle,
}

2026-09-25: The kernels a GLM MLP site launches.

resolve fails when a gpu.kernel entry point is missing (the GEMM/GEMV bases, w4a16, w4a16_gemv and the three glm5next_ffn.cu kernels). The try_kernel ones are KernelHandle(0) when absent, and the forward takes another path for each.

Fields

gemm: KernelHandle

2026-09-25: BF16 C = A @ B^T tile GEMM: dense FFN and shared expert.

gemm_f32: KernelHandle

2026-09-25: The same GEMM with FP32 output, used for the router, because glm5next_router_topk reads FP32 logits.

gemv: KernelHandle

2026-09-25: M=1 GEMVs for gemm / gemm_f32. gemv_f32 is optional; when it is 0, the router’s M=1 calls run the tile GEMM (ops::dense_mm_bf16).

gemv_f32: KernelHandle

gemv_batchm: KernelHandle

2026-09-25: dense_gemv_bf16_batchm: 2 to DENSE_GEMV_BATCHM_MAX_M (16) rows in one weight sweep, for the dense FFN and shared expert. When 0, those rows run the tile GEMM.

w4a16: KernelHandle

2026-09-25: NVFP4 w4a16_gemm tile GEMM. Resolved, but not launched by this module.

w4a16_gemv: KernelHandle

2026-09-25: NVFP4 C[1, N] = A[1, K] @ B[N, K]^T, the per-expert decode GEMV.

Its grid is tied to the kernel’s N_PER_BLOCK; use ops::w4a16_gemv_grid_x.

w4a16_gemv_sw: KernelHandle

2026-09-25: Single-warp-per-output variant of w4a16_gemv, checked bit-identical to it by examples/w4a16_gemv_sw_microtest.rs. When 0, the base kernel runs.

Its grid is ceil(N/8) (N_PER_BLOCK_SW), not the base kernel’s ceil(N/4); launch it through ops::w4a16_gemv_sw_raw or ops::w4a16_decode_gemv, which pick the grid.

w4a16_gemv_sw_moe: KernelHandle

2026-09-25: All top_k slots of one row in one launch, grid (ceil(N/8), top_k, 1). Weights come from the global-id pointer tables indexed by the router’s on-device ids. When 0, the forward reads the ids back to the host and launches per local expert.

w4a16_gemv_sw_moe_batchm: [KernelHandle; 7]

2026-09-25: Row-batched w4a16_gemv_sw_moe, indexed [rows - 2] for rows 2..=8 (w4a16_gemv_sw_moe_batchm_m2 .. _m8). Each expert in the union of the rows’ selections is swept once for all the rows that picked it.

grid.y is the union entry, not the slot, with extent rows * top_k; unfilled entries return on u_eid < 0. Needs Self::moe_row_union.

moe_row_union: KernelHandle

2026-09-25: Builds the union table the batched kernel indexes: one block of rows * top_k threads. forward_moe uses it only when rows * top_k is at most MOE_ROW_UNION_MAX_IDS (64).

swiglu: KernelHandle

2026-09-25: glm5next_swiglu_clamp, the asymmetric clamped SwiGLU. moe_silu_mul does not clamp.

router: KernelHandle

combine: KernelHandle

moe_sort_by_expert: KernelHandle

2026-09-25: moe_sort_by_expert (moe_permute.cu): counting sort of the [rows, top_k] ids into expert-contiguous order, writing sorted_token_ids, expert_offsets and token_to_perm. When 0, the grouped prefill path is off.

moe_grouped_gemm: KernelHandle

2026-09-25: The tensor-core grouped W4A16 GEMM (mma.sync.aligned.m16n8k16) for prefill, at the tile forward_prefill_gemm::gemm_tile picks. When 0, the grouped prefill path is off.

combine_indexed: KernelHandle

2026-09-25: Self::combine reading the routed rows in expert-sorted order through token_to_perm, with the same accumulation order and single rounding.

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