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,
}

NVFP4 quantized weight: packed E2M1 data + FP8 block scales + FP32 per-tensor scale.

Fields

weight: DevicePtr

Packed E2M1 weights (2 values per byte).

weight_scale: DevicePtr

Per-group FP8 block scales.

weight_scale_2: f32

Per-tensor FP32 scale factor (extracted from GPU via D2H copy at load time).

input_scale: DevicePtr

Input activation scale (FP32 on device, for FP8 activation path).

weight_scale_2_vec: DevicePtr

Per-row FP32 scale2 on device ([N] floats). When set, the w4a16_gemv_prs kernel reads scale2 per output row instead of the scalar weight_scale_2, eliminating precision loss from per-tensor absmax on outlier rows.

Implementations

impl QuantizedWeight

pub fn null() -> Self

Null weight (all pointers NULL). Used for remote experts under EP.

pub fn has_per_row_scale2(&self) -> bool

Whether this weight has per-row scale2 (for PRS GEMV dispatch).

pub fn is_null(&self) -> bool

Whether this weight points to NULL (remote expert placeholder).

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

Concatenate two NVFP4 weights by rows: [N1, K/2] + [N2, K/2] → [N1+N2, K/2].

Both weights MUST share the same K (input dimension) and the same scalar weight_scale_2. The packed weight bytes and FP8 block scales are concatenated on-GPU via cuMemcpy.

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

Transpose weight layout from [N, K/2] to [K/2, N] for coalesced GEMM reads.

Also transposes scale from [N, K/GROUP_SIZE] to [K/GROUP_SIZE, N]. Returns a NEW QuantizedWeight with freshly allocated GPU buffers, leaving the original untouched (needed for decode kernels).

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

transpose_for_gemm with an explicit scale block size. Scale tensor is [N, K/group_size]; the packed-weight transpose is group-size-independent.

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

Transpose SEVERAL weights sharing one K and concatenate them along N into a single [K/2, N_total] twin, so three GEMMs become one.

Motivation (GB10, decode M=16): the attention k/v projections are N=1024, which against the 128-wide N tile yields 8 CTAs on 48 SMs — 40 SMs idle, 23.6 GB/s, 9.75x off the bandwidth floor. Concatenating q|k|v to N=14336 gives 112 CTAs in ONE launch. Bit-identical: every output element is the same dot product against the same column, merely relocated along N.

REQUIRES all parts to share weight_scale_2 — the GEMM applies a single scale2 to the whole launch. Callers MUST verify this (the values live on device); None is returned if the caller passes an empty list.

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

transpose_concat_for_gemm with an explicit scale block size. transpose_concat_for_gemm_gs with the output ROW STRIDE padded to align_up(n_total, align), pad columns left zero.

The transposed layout puts row r at byte offset r * stride, and the tile GEMM reads B with 16-byte cp.async, which requires a 16-byte aligned source. When n_total is not a multiple of 16 — lm_head’s N is the VOCAB SIZE, 248077 here, which is ODD — 15 of every 16 rows are misaligned and the kernel faults with CUDA_ERROR_MISALIGNED_ADDRESS. Padding the stride is what makes a transposed lm_head legal at all.

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>

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

Pre-dequant NVFP4 → FP8 E4M3 for zero-overhead prefill GEMMs.

Reads B_packed[N, K/2] + B_scale[N, K/GROUP_SIZE] + scale2 and produces B_fp8[N, K] on GPU. The resulting DevicePtr can be used with fp8_gemm_t which eliminates the per-inference dequant phase entirely.

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