Modern Quantization: FP8, AWQ, and Marlin

High-Fidelity Low-Precision Representations

Quantization compresses floating point weights into smaller bit widths, cutting memory bandwidth pressure in half or by three-quarters.

GGUF v3 Binary Block Layouts (Rust #[repr(C)] Invariants)

In ggml-common.h and pure-Rust GGUF v3 loaders, K-quants partition weight matrices into QK_K = 256 weight super-blocks, whereas symmetric Q8_0 uses QK8_0 = 32 weight blocks:

#[repr(C)]
pub struct BlockQ4K {          // 2 + 2 + 12 + 128 = 144 bytes / 256 weights (4.50 bpw)
    pub d: half::f16,          // Super-block scale for quantized scales (2 bytes)
    pub dmin: half::f16,       // Super-block scale for quantized mins (2 bytes)
    pub scales: [u8; 12],      // Packed 6-bit scales & mins across 8x32 sub-blocks (12 bytes)
    pub qs: [u8; 128],         // 256 x 4-bit unsigned nibbles (128 bytes)
}

#[repr(C)]
pub struct BlockQ6K {          // 128 + 64 + 16 + 2 = 210 bytes / 256 weights (6.5625 bpw)
    pub ql: [u8; 128],         // Lower 4 bits of 256 quants (128 bytes)
    pub qh: [u8; 64],          // Upper 2 bits of 256 quants (64 bytes)
    pub scales: [i8; 16],      // 8-bit signed scales across 16x16 sub-blocks (16 bytes)
    pub d: half::f16,          // Super-block FP16 delta (2 bytes)
}

#[repr(C)]
pub struct BlockQ8_0 {         // 2 + 32 = 34 bytes / 32 weights (8.50 bpw; 272B per 256w)
    pub d: half::f16,          // Block FP16 scale factor (2 bytes)
    pub qs: [i8; 32],          // 32 x 8-bit signed integer weights (32 bytes)
}

Cross-Vendor GPU Dequantization: wgpu v30 & WGSL Subgroup Reductions

By requesting wgpu::Features::SUBGROUPS on the host and declaring enable subgroups; in WGSL, compute shaders accumulate dot products across hardware warp/wave lanes directly in registers without shared-memory barriers (achieving 215.05 μs Activation-Sparse INT4 GEMV and 311.10 μs Dense Q4 GEMV at 313.9 GB/s effective bandwidth on an RTX 5060 in Infinity's HARNESS Engine):

enable subgroups;

@group(0) @binding(0) var<storage, read> packed_q4: array<u32>;
@group(0) @binding(1) var<storage, read> x_vec: array<vec4<f32>>;
@group(0) @binding(2) var<storage, read_write> out_y: array<f32>;

@compute @workgroup_size(64)
fn q4_gemv_subgroup(
    @builtin(workgroup_id) wg: vec3<u32>,
    @builtin(local_invocation_index) tid: u32,
    @builtin(subgroup_invocation_id) lane_id: u32,
) {
    let packed = packed_q4[wg.x * 64u + tid];
    let w_lo = vec4<f32>(
        f32(packed & 0xFu), f32((packed >> 4u) & 0xFu),
        f32((packed >> 8u) & 0xFu), f32((packed >> 12u) & 0xFu)
    );
    let partial = dot(w_lo, x_vec[tid]);
    let warp_sum = subgroupAdd(partial);
    let pair_xor = subgroupShuffleXor(partial, 1u);
    if (lane_id == 0u && tid == 0u) {
        out_y[wg.x] = warp_sum + 0.0 * pair_xor;
    }
}

FlashAttention-3 & Online Softmax Tile Rescaling

Standard attention materializes an O(N2) score matrix in HBM. FlashAttention-3 tiles Q, K, V in on-chip SRAM, overlapping Hopper/Blackwell Tensor Memory Accelerator (TMA) loads with asynchronous WGMMA Tensor Core math while updating exact softmax statistics in a single pass via the online softmax rescaling identity:

m_new = max(m_old, rowmax(S_block))
l_new = exp(m_old - m_new) * l_old + sum_j exp(S_block[j] - m_new)
O_new = exp(m_old - m_new) * O_old + sum_j exp(S_block[j] - m_new) * V_block[j]
// Final normalized output after sweeping all K/V tiles: O_final = O_last / l_last

Quantization Formats Compared