Struct QuantizedWeight 

pub struct QuantizedWeight {
    pub weight: DevicePtr,
    pub weight_scale: DevicePtr,
    pub weight_scale_2: f32,
    pub input_scale: DevicePtr,
    pub weight_scale_2_vec: DevicePtr,
    pub act: Nvfp4Act,
}

2026-09-25: 4-bit quantized weight: packed E2M1 data, per-group scales and a host-side per-tensor scale.

Fields

weight: DevicePtr

2026-09-25: Packed E2M1 weights, 2 per byte.

weight_scale: DevicePtr

2026-09-25: Per-group scales: FP8 for NVFP4, E8M0 for native MXFP4.

weight_scale_2: f32

2026-09-25: Per-tensor scale, held on the host.

input_scale: DevicePtr

2026-09-25: Input activation scale on the device, or NULL.

weight_scale_2_vec: DevicePtr

2026-09-25: Per-row scale2 on the device, or NULL. The loaders in this module always set NULL; the transposes copy it through.

act: Nvfp4Act

2026-09-28: The activation format the weight-quantization policy set for this weight (metrale_config::WeightQuantPolicy::nvfp4_act). Unstamped unless a loader asks the policy; derived copies (transposes, concatenations, TP shards) carry it through.

Implementations

impl QuantizedWeight

pub fn transpose_for_gemm( &self, gpu: &dyn GpuBackend, n: usize, k: usize, ) -> Result<QuantizedWeight>

2026-09-25: Transpose the packed weight [N, K/2] to [K/2, N] and the scales [N, K/16] to [K/16, N], into new buffers.

pub fn transpose_for_gemm_gs( &self, gpu: &dyn GpuBackend, n: usize, k: usize, group_size: usize, ) -> Result<QuantizedWeight>

2026-09-25: transpose_for_gemm with an explicit scale block size. The scale tensor is [N, K/group_size]; the packed-weight transpose does not depend on it.

pub fn transpose_concat_for_gemm( gpu: &dyn GpuBackend, parts: &[(&QuantizedWeight, usize)], k: usize, ) -> Result<QuantizedWeight>

2026-09-25: Transpose several weights (weight, n) sharing one K and concatenate them along N into one [K/2, N_total] twin, so one GEMM replaces several.

The result carries the first part’s weight_scale_2, input_scale and weight_scale_2_vec, so the caller must ensure every part shares them; nothing here checks. An empty parts is an error.

pub fn transpose_concat_for_gemm_padded( gpu: &dyn GpuBackend, parts: &[(&QuantizedWeight, usize)], k: usize, group_size: usize, align: usize, ) -> Result<(QuantizedWeight, usize)>

2026-09-25: transpose_concat_for_gemm_gs with the output row stride padded to align_up(n_total, align) and the pad columns zeroed.

Row r of the result starts at byte r * stride, and the tile GEMM reads B with 16-byte cp.async, which needs 16-byte aligned rows; an N such as an odd vocab size is not. Returns (weight, stride); pass the stride to w4a16_gemm_n128_ldb.

pub fn transpose_concat_for_gemm_gs( gpu: &dyn GpuBackend, parts: &[(&QuantizedWeight, usize)], k: usize, group_size: usize, ) -> Result<QuantizedWeight>

2026-09-25: transpose_concat_for_gemm with an explicit scale block size.

impl QuantizedWeight

pub fn null() -> Self

2026-09-25: All pointers NULL and weight_scale_2 0.0.

pub fn has_per_row_scale2(&self) -> bool

2026-09-25: True when weight_scale_2_vec is not NULL.

pub fn is_null(&self) -> bool

2026-09-25: True when weight is NULL.

pub fn concat_rows( &self, other: &QuantizedWeight, n1: usize, n2: usize, k: usize, gpu: &dyn GpuBackend, ) -> Result<QuantizedWeight>

2026-09-25: Concatenate two NVFP4 weights by rows into new buffers: [N1, K/2] + [N2, K/2] → [N1+N2, K/2], with group size 16.

Both must have the same K, which the caller guarantees. Their weight_scale_2 must be bit-identical, or it returns an error, because the result carries self’s scale for every row. The result has NULL input_scale and weight_scale_2_vec.

pub fn predequant_to_fp8( &self, gpu: &dyn GpuBackend, predequant_kernel: KernelHandle, n: usize, k: usize, stream: u64, ) -> Result<DevicePtr>

2026-09-25: Dequantize this NVFP4 weight [n, k] into a new FP8 E4M3 buffer [n, k] with weight_scale_2 applied, for prefill GEMMs that read FP8. It synchronizes stream before returning.

Trait Implementations

impl Clone for QuantizedWeight

fn clone(&self) -> QuantizedWeight

Returns a duplicate of the value.

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

Performs copy-assignment from source.

impl Debug for QuantizedWeight

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

Formats the value using the given formatter.

impl From<QuantizedWeight> for QuantWeight

fn from(w: QuantizedWeight) -> Self

Converts to this type from the input type.

impl Copy for QuantizedWeight