Running a 35B MoE Language Model on 8 GB VRAM in Mosaic

Large Language Models (LLMs) continue to deliver dramatic leaps in reasoning, code generation, and domain knowledge, but their hardware requirements have grown just as aggressively. A state-of-the-art open-weights Mixture-of-Experts model like Qwen3.6-35B-A3B contains roughly 35 billion parameters:

• At 16-bit floating point precision (FP16 / BF16), the weights alone consume 70 GB of memory.
• Even with aggressive 8-bit quantization (FP8 E4M3), the model requires at least 35 GB of memory before allocating a single byte for the Key-Value (KV) activation cache.

For practitioners without access to multi-GPU enterprise clusters ($4\times\text{A100}$ or $8\times\text{H100}$), running these models locally has historically felt like a non-starter. Consumer laptop GPUs, such as the NVIDIA GeForce RTX 4060 Laptop GPU, are constrained to an 8 GB VRAM budget.

To solve this, I designed and implemented Mosaic—a local-first inference engine and interactive terminal CLI that runs the full 35B MoE model smoothly on an 8 GB consumer GPU. Rather than viewing model evaluation purely as a linear algebra puzzle, Mosaic approaches inference as an operating system memory hierarchy problem: combining zero-copy Linux kernel memory mapping, an intelligent three-tier expert storage hierarchy, handwritten C++ AVX2 SIMD execution, and linear recurrent attention.

1. The Inference Memory Wall & The MoE Opportunity

When an inference framework attempts to run a model that exceeds available GPU memory, the standard fallback is layer offloading (implemented in frameworks like llama.cpp or vLLM with CPU offload).

Why Naive Offloading Fails

In standard offloading, the host CPU transfers entire layer weights across the PCIe bus on every forward pass:

[ System RAM (DDR5: ~50 GB/s) ]
             │
             │ PCIe Gen 4 x8 Bus (~12 - 14 GB/s effective)
             ▼
[ GPU VRAM (GDDR6: ~280 GB/s) ]

To generate a single token through a 40-layer network, the engine must stream 35 GB of weight tensors across the PCIe bus. At an effective throughput of 14 GB/s, transferring weights alone takes:

This caps generation speed at 0.4 tokens/second—turning an interactive conversation into an excruciatingly slow slideshow.

The Sparsity & Working-Set Insight

Mixture-of-Experts architectures provide a unique structural opening. In Qwen3.6-35B-A3B:

• The model has 40 layers.
• Each layer contains 64 independent Feed-Forward Network (FFN) “experts”.
• For any single token, a learned router selects only the Top-8 experts to evaluate.

While the total checkpoint size is 35 billion parameters, the active parameter count per token is only ~3.2 billion parameters.

Crucially, natural language exhibits strong temporal and semantic routing locality. When a model generates a function implementation in C++ or explains an aerospace trajectory equation, the router does not pick random experts uniformly. Instead, consecutive tokens repeatedly route into the exact same cluster of 6 to 12 specialized experts across adjacent layers.

If we can capture that recurring working set in fast memory while avoiding redundant PCIe bandwidth transfers for cold outliers, we can break through the memory wall.

2. End-to-End System Architecture

Mosaic organizes hardware resources into three distinct tiers based on latency, throughput, and capacity:

┌─────────────────────────────────────────────────────────────┐
│ Tier 0: NVMe SSD Backing Store (~35 GB Checkpoint)         │
│ Read Throughput: ~4.5 GB/s · Latency: ~25 µs               │
└──────────────────────────────┬──────────────────────────────┘
                               │
                               │ mmap() + madvise(MADV_WILLNEED)
                               ▼
┌─────────────────────────────────────────────────────────────┐
│ Tier 1: Linux Kernel Page Cache (System RAM: ~8 - 14 GB)    │
│ Zero-copy virtual memory · No Python heap allocation        │
└──────────────┬───────────────────────────────┬──────────────┘
               │                               │
               │ Cold Experts (< 3 hits)       │ Hot Experts (≥ 3 hits)
               ▼                               ▼
┌───────────────────────────────┐ ┌───────────────────────────┐
│ Tier 1B: C++ AVX2 CPU Engine  │ │ Tier 2: GPU VRAM Cache    │
│ Evaluated in RAM via SIMD FMA │ │ ~3.2 GiB LRU on GPU VRAM  │
│ Zero PCIe bus traffic!        │ │ Full Tensor Core speed    │
└──────────────┬────────────────┘ └─────────────┬─────────────┘
               │                                │
               └───────────────┬────────────────┘
                               ▼
            [ Layer Activation Accumulation ]

1. Tier 0 (NVMe SSD): Stores the 40 per-layer model shards (layers-00.safetensors through layers-39.safetensors, ~0.84 GB each).
2. Tier 1 (Kernel Page Cache): Mmaps all shard files once into the virtual address space. Pages fault into system RAM on demand without duplicate Python heap allocations.
3. Tier 2 (GPU VRAM LRU Cache): Allocates a dedicated ~3.2 GiB arena on the GPU specifically for active experts. Frequently invoked experts remain in high-speed GDDR6 memory.
4. Tier 1B (CPU AVX2 SIMD Engine): When an expert is “cold” (rarely called), transferring its weights over PCIe would waste more bandwidth than computing it in place. Mosaic evaluates cold experts directly on the CPU using handwritten AVX2 assembly intrinsics while the GPU computes parallel self-attention.

3. Shard Catalog & Blockwise FP8 Quantization

Standard weight formats store floating point matrices in uniform 16-bit tensors. To minimize disk I/O and cache footprint, Mosaic formats weights in blockwise FP8 (E4M3) with paired block scale factors.

The Math of Blockwise E4M3 Quantization

A standard FP8 E4M3 byte contains 1 sign bit, 4 exponent bits, and 3 mantissa bits:

 7   6   5   4   3   2   1   0
┌───┬───────────────┬───────────┐
│ S │ E   E   E   E │ M   M   M │
└───┴───────────────┴───────────┘

Because deep neural network weight distributions develop high dynamic ranges and outliers across large projection matrices, uniform quantization causes severe degradation. Mosaic divides weight matrices into $128 \times 128$ blocks. Each block has an independent 16-bit Brain Floating Point (bfloat16) inverse scaling factor $s$:

Where $Q_{i, j} \in \text{FP8}(\text{E4M3})$ is the quantized integer byte, and $W_{i, j}$ is the reconstructed floating-point weight. This preserves sub-normal dynamic range while squeezing checkpoint storage by 50% relative to FP16.

Zero-Copy Memory Mapping

Traditional machine learning loaders copy weight buffers into Python objects, immediately consuming $2\times$ the memory footprint. Mosaic bypasses Python memory entirely:

import mmap
import os

class ShardCatalog:
    def __init__(self, shard_paths: list[str]):
        self.handles = []
        self.mappings = []
        for path in shard_paths:
            fd = os.open(path, os.O_RDONLY)
            mapping = mmap.mmap(fd, 0, access=mmap.ACCESS_READ)
            # Advise the kernel to expect sequential window access
            mapping.madvise(mmap.MADV_WILLNEED)
            self.handles.append(fd)
            self.mappings.append(mapping)

The underlying memory pages are owned directly by the Linux kernel. If the system experiences memory pressure, the kernel can drop clean clean pages and reload them from NVMe in sub-millisecond timeframes without triggering swap death.

4. Tiered Dynamic Expert Execution

At each layer $l$, the input activation $x_t \in \mathbb{R}^d$ passes through the gating network:

The router selects the Top-$k$ expert indices ($k = 8$) and computes their normalized softmax weights:

Mosaic tracks routing frequency across a rolling context window. When evaluating expert $i$:

# Tiered routing decision logic
class TieredMoERunner:
    def execute_expert(self, layer_idx: int, expert_idx: int, x: torch.Tensor) -> torch.Tensor:
        hit_count = self.tracker.get_hits(layer_idx, expert_idx)
        
        if hit_count >= 3:
            # Hot Expert: fetch from GPU VRAM LRU cache (or promote)
            gpu_expert = self.vram_cache.get_or_load(layer_idx, expert_idx)
            return gpu_expert.forward(x)
        else:
            # Cold Expert: execute directly on CPU via C++ AVX2 kernel
            # Zero PCIe bus transfer overhead!
            self.tracker.record_hit(layer_idx, expert_idx)
            weight_ptr = self.catalog.get_expert_ptr(layer_idx, expert_idx)
            return cpu_executor.forward_fp8(x.data_ptr(), weight_ptr)

By filtering out one-off “cold” experts from the GPU transfer queue, PCIe bus utilization drops by over 70%. The GPU spends its time performing high-throughput matrix multiplications on stable, resident working sets.

5. Handwritten C++ AVX2 SIMD Kernel

Evaluating cold experts on the CPU could easily become the next bottleneck if implemented naively. Standard implementations unpack FP8 bytes using a 256-entry lookup table (LUT). However, when multiple CPU worker threads read a lookup table concurrently, L1 data cache eviction and translation lookaside buffer (TLB) stalls degrade performance.

Mosaic implements an arithmetic dequantization and Fused Multiply-Add (FMA) kernel in C++17 (src/mosaic/kernels/cpu_executor.cpp) using direct Intel AVX2 SIMD intrinsics:

#include <immintrin.h>
#include <omp.h>
#include <cstdint>

// Unpacks 8 packed FP8 E4M3 bytes directly into a 256-bit AVX float vector
// Arithmetic decoding without memory lookup tables
inline __m256 dequantize_fp8_avx2(const uint8_t* src, float scale) {
    // Load 8 unsigned bytes and zero-extend to 32-bit integer vector
    __m128i raw8 = _mm_loadl_epi64(reinterpret_cast<const __m128i*>(src));
    __m256i raw32 = _mm256_cvtepu8_epi32(raw8);

    // Extract sign, exponent, and mantissa components
    __m256i sign = _mm256_and_si256(raw32, _mm256_set1_epi32(0x80));
    __m256i exp  = _mm256_and_si256(raw32, _mm256_set1_epi32(0x78));
    __m256i mant = _mm256_and_si256(raw32, _mm256_set1_epi32(0x07));

    // Construct IEEE-754 single-precision float (1 sign, 8 exponent, 23 mantissa)
    __m256i ieee_sign = _mm256_slli_epi32(sign, 24);
    __m256i ieee_exp  = _mm256_slli_epi32(_mm256_add_epi32(_mm256_srli_epi32(exp, 3), _mm256_set1_epi32(127 - 7)), 23);
    __m256i ieee_mant = _mm256_slli_epi32(mant, 20);

    __m256i ieee_bits = _mm256_or_si256(ieee_sign, _mm256_or_si256(ieee_exp, ieee_mant));
    __m256 float_val = _mm256_castsi256_ps(ieee_bits);

    // Multiply by per-block bfloat16 inverse scale factor
    return _mm256_mul_ps(float_val, _mm256_set1_ps(scale));
}

// 8-way vectorized GEMV (Matrix-Vector multiplication) inner loop
void gemv_fp8_avx2(
    const float* __restrict__ activation,
    const uint8_t* __restrict__ weight_matrix,
    const float* __restrict__ block_scales,
    float* __restrict__ output,
    int in_features,
    int out_features
) {
    #pragma omp parallel for schedule(static)
    for (int row = 0; row < out_features; ++row) {
        __m256 acc0 = _mm256_setzero_ps();
        __m256 acc1 = _mm256_setzero_ps();

        const uint8_t* row_weights = weight_matrix + (row * in_features);

        for (int col = 0; col < in_features; col += 16) {
            float scale = block_scales[(row / 128) * (in_features / 128) + (col / 128)];

            __m256 w0 = dequantize_fp8_avx2(row_weights + col, scale);
            __m256 w1 = dequantize_fp8_avx2(row_weights + col + 8, scale);

            __m256 x0 = _mm256_loadu_ps(activation + col);
            __m256 x1 = _mm256_loadu_ps(activation + col + 8);

            // Fused Multiply-Add: acc = acc + (w * x)
            acc0 = _mm256_fmadd_ps(w0, x0, acc0);
            acc1 = _mm256_fmadd_ps(w1, x1, acc1);
        }

        // Horizontal sum across 8-wide SIMD vector
        __m256 sum = _mm256_add_ps(acc0, acc1);
        __m128 lo = _mm256_castps256_ps128(sum);
        __m128 hi = _mm256_extractf128_ps(sum, 1);
        __m128 sum128 = _mm_add_ps(lo, hi);
        sum128 = _mm_hadd_ps(sum128, sum128);
        sum128 = _mm_hadd_ps(sum128, sum128);

        output[row] = _mm_cvtss_f32(sum128);
    }
}

On an 8-core AMD Ryzen 7 processor, this vectorized C++ kernel evaluates a cold expert projection in ~0.08 milliseconds, keeping CPU computation fully overlapped with asynchronous GPU attention kernels.

6. Hybrid Attention: Recurrent Delta Rule & Constant Memory Scaling

Memory challenges in large models are not limited to weights. As conversational context lengthens, standard Multi-Head Attention (MHA) consumes memory quadratically with sequence length $L$:

At a context window of 32,768 tokens, an uncompressed FP16 KV cache for a 35B model consumes ~5.8 GB of VRAM. On an 8 GB GPU with 3.2 GB reserved for weights, the KV cache alone would trigger an out-of-memory crash.

Mosaic integrates a Hybrid Attention Architecture:

Layer Architecture Distribution:
┌─────────────────────────────────────────────────────────────┐
│ 85% of Layers: Gated Linear Recurrent Attention (Delta Rule) │
│ State Buffer: O(1) Constant Size · Compute: O(1) per token  │
└─────────────────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────────────────┐
│ 15% of Layers: Full Multi-Head Self-Attention (Paged KV)     │
│ Anchored at strategic intervals for global associative recall│
└─────────────────────────────────────────────────────────────┘

The Linear Recurrent Delta Rule

For the majority of layers, Mosaic uses a recurrent delta update rule over a fixed-dimension memory matrix $S_t \in \mathbb{R}^{d \times d}$:

Because $S_t$ has fixed dimensions ($d \times d$), its memory footprint never grows whether the conversation has 10 tokens or 32,000 tokens. The computational cost per generated token is strictly $O(1)$.

For the remaining 15% of anchor layers, full attention is computed using a paged block allocator (similar to vLLM’s Virtual Memory Paging) to maintain surgical factual retrieval over long documents without memory fragmentation.

7. Real-World Benchmarks & Performance Metrics

We benchmarked Mosaic on a consumer laptop:

• GPU: NVIDIA GeForce RTX 4060 Laptop GPU (8 GB GDDR6 VRAM, 140W TGP).
• CPU: AMD Ryzen 7 7840HS (8 cores / 16 threads, AVX2 / AVX-512 support).
• RAM: 16 GB DDR5-5600.
• Storage: 1 TB PCIe 4.0 NVMe SSD (~4,800 MB/s sequential read).

Measured Generation Throughput

Memory Footprint Breakdown

RTX 4060 (8.0 GiB Physical VRAM Limit)
┌───────────────────────────────────────────────────────────┐
│ GPU Expert LRU Cache (3.2 GiB)                            │
├─────────────────────────┬─────────────────────────────────┤
│ Paged KV Cache (0.6 GiB)│ Free Headroom (3.8 GiB)         │
│ & Model Embeddings      │ Display compositor & OS buffers │
└─────────────────────────┴─────────────────────────────────┘

The system stabilizes at ~4.2 GiB total VRAM usage, leaving almost 4 GB of headroom for the desktop compositor, browser windows, and operating system buffers.

8. What This Means for Local AI

Mosaic highlights an important shift in AI systems engineering. Running state-of-the-art models locally does not always require purchasing enterprise datacenter GPUs. When models are treated through the lens of classical systems programming:

1. Memory hierarchies are malleable: With fast NVMe storage, virtual memory page tables, and kernel prefetching, disk can serve as an effective extension of RAM.
2. Sparsity enables locality: Mixture-of-Experts models allow software to exploit working-set cache locality in ways dense models cannot.
3. CPUs are underutilized matrix engines: SIMD instruction sets (AVX2 / AVX-512) and OpenMP parallelism make modern multicore CPUs surprisingly capable at evaluating sparse, non-critical sub-networks in parallel with GPU tensor cores.

By aligning neural network execution with the realities of consumer hardware, Mosaic makes private, local 35B model inference a reality on a standard developer laptop.