Struct GemmDispatch 

pub struct GemmDispatch {
    pub fp8_blockscaled_prefill: bool,
    pub cublas_gemm: bool,
    pub cublas_fp8: bool,
    pub cutlass_gemm: bool,
    pub cutlass_nvfp4_gemm: bool,
    pub cutlass_nvfp4_qkvz: bool,
    pub cutlass_nvfp4_attn_q: bool,
    pub cutlass_nvfp4_attn_kv: bool,
    pub cutlass_nvfp4_attn_o: bool,
    pub cutlass_nvfp4_ssm_out: bool,
    pub w4a16_variant: u8,
}

Which GEMM implementation each projection takes.

Plain Copy data, resolved from the environment at model construction.

Fields

fp8_blockscaled_prefill: bool

Block-scaled FP8 prefill (per-128-block weight scales + per-token activation scales). The DEFAULT for block-scaled FP8 checkpoints since 2026-06-17: it matches vLLM’s per-block precision and avoids the single-scale path, whose collapse of per-block dynamic range pushed long-context tool-arg decode into the FP8 argmax-flip regime (B1 drift gauge ~1400 → ~100 once block-scaled prefill is on). Opt out with ATLAS_FP8_SINGLE_SCALE=1 — diagnostic/fallback only.

cublas_gemm: bool

cuBLASLt BF16 GEMM. The hand-written mma.sync projection GEMMs reach only ~30% of the cuBLAS bf16 ceiling on GB10.

cublas_fp8: bool

Native-FP8 cuBLASLt GEMM.

cutlass_gemm: bool

CUTLASS BF16 GEMM, scoped to dense projections using the same FP8→BF16 cached dequant as cuBLASLt.

cutlass_nvfp4_gemm: bool

Native CUTLASS NVFP4 GEMM: quantizes activations to CUTLASS NVFP4 and consumes transposed Atlas NVFP4 weights after repacking scales into the CUTLASS SM120 layout. Implies every per-projection NVFP4 flag below.

cutlass_nvfp4_qkvz: bool

cutlass_nvfp4_attn_q: bool

cutlass_nvfp4_attn_kv: bool

cutlass_nvfp4_attn_o: bool

cutlass_nvfp4_ssm_out: bool

w4a16_variant: u8

ATLAS_W4A16_VARIANT — 1/2/3 pin a kernel variant, 0 = auto (v2). A dispatch decision like every other field here, so it belongs on the struct the forward pass already carries rather than in a OnceLock that would pin the first model’s choice.

Implementations

impl GemmDispatch

pub fn from_env() -> Self

Resolve from the environment. Called once, when the model is built.

pub fn defaults() -> Self

Everything off, block-scaled FP8 prefill on — the shape a build with no ATLAS_* set in the environment resolves to. Tests construct a context with this instead of mutating the process environment.

pub fn cutlass_nvfp4_attn_qkv(&self, label: &str) -> bool

NVFP4 attention Q/K/V enabled for the named projection.

Trait Implementations

impl Clone for GemmDispatch

fn clone(&self) -> GemmDispatch

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for GemmDispatch

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Default for GemmDispatch

fn default() -> Self

Returns the “default value” for a type.

impl PartialEq for GemmDispatch

fn eq(&self, other: &GemmDispatch) -> bool

Tests for self and other values to be equal, and is used by ==.

fn ne(&self, other: &Rhs) -> bool

Tests for !=. The default implementation is almost always sufficient, and should not be overridden without very good reason.

impl Copy for GemmDispatch

impl Eq for GemmDispatch

impl StructuralPartialEq for GemmDispatch