Running Qwen3.8-Flash-Next 125B on Three RTX 3090s at 80 Tokens per Second

Deploying state-of-the-art open models on consumer or workstation hardware is the holy grail for modern AI developers. When Qwen released Qwen3.8-Flash-Next, a preview of their upcoming Qwen4 architecture featuring a massive 125-billion-parameter Mixture-of-Experts (MoE) design, the benchmarks spoke for themselves: a +16.5 point increase on DeepSWE agentic coding and +22.3 on JobBench over the standard Qwen3.8-27B.

However, running a 125B parameter model locally presents a formidable hardware hurdle. In standard 16-bit float precision, the model requires ~360 GB of memory. Even a aggressively quantized 4-bit version (such as Unsloth's UD-Q4_K_XL) weighs in at 111 GB, with the routed expert tensors alone consuming 71.7 GB. On a budget workstation equipped with three NVIDIA RTX 3090 GPUs (providing a combined VRAM pool of 72 GB across non-NVLink PCIe slots), stock llama.cpp is forced to offload roughly 25% of the expert computations to system RAM.

Because CPU RAM bandwidth is roughly 10x to 20x slower than VRAM—and because system memory swaps bounce work back and forth across PCIe buses—stock execution degrades to a sluggish 23 tokens per second.

By analyzing expert activation patterns, implementing dynamic multi-precision quantization tiers, and writing ~400 lines of custom CUDA/C++ patches inside llama.cpp, it is possible to achieve over 80 tokens per second fully inside 72 GB of VRAM without sacrificing model intelligence. For developers who prefer not to maintain custom GPU kernels and hardware setups, high-throughput commercial endpoint aggregators like n1n.ai provide instant, zero-maintenance access to top-tier MoE models with minimal latency.

Benchmark Comparison: Qwen3.8-Flash-Next vs Qwen3.8-27B

To understand why fitting this 125B model onto consumer hardware is worth the engineering effort, consider the original performance metrics published by the Qwen team:

While dense architectures like Qwen3.8-27B run fast out of the box (~135 tokens/s on two RTX 3090s using vLLM), their agentic problem-solving capabilities fall noticeably short when executing complex multi-step reasoning. The 125B Flash-Next model bridges this gap, but its memory footprint demands a fundamental redesign of how llama.cpp handles expert weights.

The Bottleneck: How Stock MoE Execution Fails in VRAM-Constrained Environments

In a Mixture-of-Experts architecture, each transformer layer contains numerous small feed-forward networks called experts—512 per layer in Qwen3.8-Flash-Next. A lightweight router network dynamically evaluates the input context and selects the top 10 most relevant experts for every incoming token. Although the model contains 125B parameters, it only computes ~6B active parameters per token.

When llama.cpp attempts to load a model that exceeds available GPU VRAM, its default allocator offloads entire layer tensors to system RAM.

[ Stock llama.cpp MoE Execution ]
Token In ──► Router (GPU) ──► Selected Experts
                               ├─► Hot Experts (GPU VRAM)   --> Fast MatMul (~2.5ms)
                               └─► Cold Experts (System RAM) --> Slow CPU Transfer (~15ms)
                                                                 (PCIe Bottleneck)

In a machine with dual Xeon E5-2660 v4 CPUs running single-channel RAM, transferring weights from system RAM back and forth across PCIe lanes introduces severe synchronization penalties:

1. GPU-CPU Roundtrips: Every single token generation steps through 48+ layers. If a layer's selected experts sit in CPU RAM, the main thread blocks awaiting host-to-device transfers.
2. Page Cache Eviction: Under heavy context pressure, memory-mapped model files trigger disk I/O reads from NVMe storage, causing thread delays (state D).
3. Speculative Decoding Breakdown: When using Multi-Token Prediction (MTP) draft heads, 5 draft tokens must be verified simultaneously. Checking 6 tokens at once means activation of up to 60 distinct experts per layer instead of 10, drastically increasing the hit rate on slow CPU-side experts.

Lever 1: Expert Frequency Sorting (Hot/Cold Splitting)

The foundational insight behind this optimization comes from analyzing the importance matrix (imatrix) generated across real-world text corpora. Expert utilization in MoE models is strikingly power-law distributed:

Key Observation: Across average layers, the top 25% busiest experts process 52% of all tokens, while the top 80% handle 95% of routed requests.

Standard llama.cpp treats all 512 experts in a layer as a monolithic 3D tensor located either entirely on the GPU or entirely on the CPU. By writing an offline tensor permutation patch, each layer's expert tensor is split at load time into two separate pools:

• Hot Pool (GPU VRAM): Contains the top-performing experts ranked by historical routing frequency.
• Cold Pool (System RAM): Contains low-frequency long-tail experts.

# Conceptual representation of expert weight splitting
import torch

def split_expert_tensor(layer_experts, imatrix_scores, vram_capacity_ratio=0.75):
    # Rank expert IDs by cumulative activation frequency
    sorted_indices = torch.argsort(imatrix_scores, descending=True)
    
    num_hot = int(len(sorted_indices) * vram_capacity_ratio)
    hot_indices = sorted_indices[:num_hot]
    cold_indices = sorted_indices[num_hot:]
    
    hot_experts = layer_experts[hot_indices].to("cuda")
    cold_experts = layer_experts[cold_indices].to("cpu")
    
    return hot_experts, cold_experts

By placing the coldest 19% of experts on the CPU, cold experts serve only ~4.5% of total routed tokens (down from 19%). This simple split reduced slow-path host-to-device traffic by 4x, lifting performance from 23 tok/s to 32 tok/s. However, synchronization overhead between CPU and GPU threads remained the primary speed limiter. To eliminate CPU roundtrips completely, 100% of experts needed to live in VRAM.

Lever 2: Dynamic Expert Tiering & Mixed-Precision Allocation

To squeeze all 512 experts across all layers into 72 GB of VRAM alongside the KV cache and model embeddings, we apply variable-rate precision quantization based on expert popularity.

Rather than quantizing every expert to a uniform 3-bit format (such as IQ3_S), an offline re-packer tool (expert_tiers.py) assigns custom bit budgets:

1. Tier 1 (Busiest Experts - Top 25%): Quantized using high-precision formats (Q6_K for gate/up projections, Q8_0 for down projections).
2. Tier 2 (Moderate Experts - Middle 55%): Quantized using mid-tier formats (IQ4_XS / IQ4_NL).
3. Tier 3 (Rare Experts - Bottom 20%): Squeezed heavily using extreme low-bit formats (IQ3_XXS / MXFP4).

[ Three-Tier VRAM Allocation Model ]
Total VRAM: 72 GB
┌────────────────────────────────────────────────────────┐
│ Non-Expert Weights, KV Cache & Embeddings (~18 GB)    │
├────────────────────────────────────────────────────────┤
│ Tier 1 Experts (High Precision: Q6_K / Q8_0)  ~22 GB  │
├────────────────────────────────────────────────────────┤
│ Tier 2 Experts (Mid Precision: IQ4_XS / IQ4_NL) ~21 GB│
├────────────────────────────────────────────────────────┤
│ Tier 3 Experts (Low Precision: IQ3_XXS/MXFP4)  ~10 GB │
└────────────────────────────────────────────────────────┘

Because down-projection matrices feature 640-wide dimensions—which ultra-low 2-bit quants like IQ2_XXS cannot process due to 256-block alignment constraints—the re-packer implements an adaptive fallback ladder using non-linear quantization formats like IQ4_NL and MXFP4.

If you prefer to integrate pre-optimized endpoints directly into your production pipelines without manually managing quantization matrices and CUDA setups, platforms like n1n.ai offer scalable infrastructure for ultra-fast model inference.

Lever 3: Patching llama.cpp for Multi-Tier Execution & MTP

Stock llama.cpp execution kernels assume uniform quantization formats across all experts in a layer. Enabling three separate precision tiers per layer required modifying ggml CUDA kernels (mul_mat_id) and the execution graph (~400 lines of custom C++/CUDA code).

The key kernel modification allows a single router decision to dispatch indices into three precision shelves simultaneously:

// Patched mul_mat_id execution step in CUDA kernel
__global__ void k_mul_mat_id_tiered(
    const void * src0, const void * src1, void * dst,
    const int32_t * ids, const int min_expert_id, const int max_expert_id,
    int n_experts_per_tok
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int selected_expert = ids[idx];

    // Range-check: Process only if selected expert belongs to current precision shelf
    if (selected_expert < min_expert_id || selected_expert >= max_expert_id) {
        return; // Explicit skip to prevent out-of-bounds reading
    }

    // Compute matrix multiplication for assigned tier...
}

Critical Engineering Pitfalls Encountered & Resolved

During implementation, four subtle bugs were identified and fixed:

1. CUDA MoE Duplicate Index Crash: Initially, "not on this shelf" branches defaulted to index 0. When a token selected two cold experts, index 0 appeared twice in the same warp, violating CUDA's single-occurrence expert grouping assumption. This was resolved by creating an explicit pass-through skip mechanism (min_expert_id / max_expert_id boundary checks).
2. Parameter Slot Overwrites: Custom flags controlling skip logic were initially stored in op_params[0]. However, llama.cpp internally overwrites op_params[0] to store accumulator precision settings. The parameter flag was relocated to index op_params[7] using a magic bitmask.
3. Unsigned Integer Promotion Trap: C++ ternary checks using ids ? ids[i] : blockIdx.x implicitly promoted negative values (-1 sentinels representing skips) to unsigned 32-bit integers (4,294,967,295), causing memory access faults 4 billion rows past the CUDA buffer.
4. Fused Kernel Destination Hiding: Decode pipelines fuse gate, up-projection, and activation functions into a single CUDA kernel. The fused operator tracks its target node on activation rather than the source matrix multiplication node. The kernel patch was adjusted to trace source node properties dynamically.

Quantitative Evaluation & Performance Results

To evaluate quality loss, Kullback-Leibler Divergence (KLD) and perplexity were measured against the uncompressed 8-bit reference model (Q8_0) using 24,576 tokens of English Wikipedia text.

On standardized benchmark evaluations, the 128k Tiered Profile scored 95.5% on GSM8K (grade school math), fully matching the baseline range of the dense 27B model (95.0%–96.5%).

Throughput & Context Window Scaling

With speculative decoding enabled via a 4-token MTP draft head, inference speed scaled across context sizes as follows:

In a synthetic needle-in-a-haystack test, the 256k context profile accurately retrieved a random 10-character key (6BC5XYE8FS) hidden at 45% depth within 173,692 tokens of context, proving the retention of long-context retrieval dynamics.

Power Consumption and Operational Cost

Measurements taken directly via nvidia-smi across the three RTX 3090 GPUs yielded the following metrics:

• Active Decoding Power: ~540 Watts aggregate.
• Idle Power: ~124 Watts.
• Energy Consumption: ~7.5 to 8.4 Joules per token.
• Cost Efficiency: At standard residential power rates ($0.15/kWh), generating 1,000,000 tokens costs approximately **$0.35 to $0.40**.

Step-by-Step Implementation Guide

Follow these steps to replicate the patch build and deploy Qwen3.8-Flash-Next on local multi-GPU setups.

Step 1: Clone and Apply Custom C++ Patches

# Clone llama.cpp at target commit
git clone https://github.com/ggerganov/llama.cpp.git
cd llama.cpp
git checkout 81bc6b8

# Apply MTP speculative PR and tiering patches
git apply patches/0001-qwen4exp-mtp-pr28243.patch
git apply patches/0002-tiered-experts-hot-cold-split.patch

# Build with CUDA support
cmake -B build -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=86
cmake --build build -j$(nproc)

Step 2: Re-pack Expert Tiers Offline

# Run the tiering tool against the Q8_0 base model using importance matrix data
python3 tools/expert_tiers.py write \
  --src Qwen3.8-Flash-Next-Q8_0-00001-of-00006.gguf \
  --imatrix imatrix_unsloth.gguf \
  --budget-gib 53 \
  --gu Q6_K,IQ4_XS,IQ3_XXS \
  --dn Q8_0,IQ4_NL,MXFP4 \
  --out fn-tier53.gguf

Step 3: Launch Local Server with MTP Speculative Decoding

./build/bin/llama-server \
  -m fn-tier53-00001-of-00002.gguf \
  -ngl 99 \
  -ts 16,16,16 \
  -c 131072 \
  -b 1024 \
  -ub 256 \
  -fa on \
  --jinja \
  -md mtp-shared-iq4.gguf \
  --spec-type draft-mtp \
  --spec-draft-n-max 4

Summary and Production Recommendations

Optimizing hardware utilization through expert tiering demonstrates that 100B+ MoE architectures can run effectively on consumer hardware. However, self-hosting complex MoE setups comes with operational trade-offs:

1. Single-User Constraints: Local execution is restricted to sequential single-request throughput. Multi-tenant batching scales poorly compared to dedicated inference engines like vLLM.
2. Prefill Bottlenecks: Processing long 100k+ token prompts can require 2–3 minutes before the first token streams out.

For teams seeking production-ready, low-latency API infrastructure without hardware maintenance overhead, testing commercial endpoints via n1n.ai provides immediate scalability.

Get a free API key at n1n.ai.