Running 753B MoE Models on Consumer GPUs: Hand-Written SASS Kernels for 2-Bit Experts

Hook

A single RTX 5090 running a 753-billion parameter model shouldn't be possible—yet someone reverse-engineered NVIDIA's entire Blackwell instruction set and wrote 5,600 lines of assembly to make it happen.

Context

Frontier mixture-of-experts models like DeepSeek-V4 (753B total parameters, 51B active) and GLM-5.2 (621B total, 54B active) represent the current state-of-the-art in open-weight language models, but they're effectively out of reach for researchers without datacenter access. Even with aggressive FP8 quantization, these models require 300GB+ of VRAM—well beyond consumer hardware. The official DeepSeek-V4 FP4 checkpoint from NVIDIA drops memory requirements significantly, but vLLM doesn't support loading it on consumer Blackwell cards, and even if it did, a single RTX 5090 (32GB) or dual-card setup (64GB) still falls short.

The problem runs deeper than just model size. MoE architectures route each token to a small subset of expert networks, creating a natural caching opportunity—but vLLM treats all parameters equally, loading what fits and failing when it doesn't. Meanwhile, NVIDIA provides no public toolchain for writing custom CUDA kernels at the SASS assembly level for consumer Blackwell (SM120), making it impossible to exploit undocumented tensor core instructions for aggressive quantization. vLLM-Moet exists because its author wanted to run these frontier models on hardware they actually own, and the only path forward was to build the entire stack from scratch: reverse-engineer the instruction set, write assembly kernels by hand, and rearchitect how vLLM thinks about model residency.

Technical Insight

The core architecture inverts the traditional GPU memory model by treating MoE experts as a three-tier cache hierarchy: 2-bit compressed experts in pinned host RAM (base tier), frequently-used experts in GPU DRAM (hot cache), and an FP4 "delta" pool for precision recovery on critical routing decisions. The 2-bit quantization uses a sign-symmetric codebook {-4, -1, 1, 4} discovered through empirical analysis—asymmetric codebooks cause catastrophic bias accumulation across layers, with 99% negative bias measured on GLM-5.2's expert weights. By forcing equal L2 error across positive and negative quantization bins, the patch achieves better acceptance rates (2.73 vs 2.68) than the baseline FP4 checkpoint despite using half the bits.

The SASS kernels implement in-register decompression using PRMT-LUT instructions feeding block-scaled QMMA.SF tensor cores. Here's the conceptual data flow for a 2-bit expert matmul:

// Conceptual kernel flow (actual implementation is 800+ lines of SASS)
__global__ void moe_w2_mm_kernel(
    half* output,           // [tokens, hidden_dim]
    half* activations,      // [tokens, hidden_dim] 
    uint8_t* w2_experts,    // [num_experts, hidden_dim, hidden_dim/4] packed
    int32_t* routing_table, // [tokens] -> expert_id
    float* scales           // [num_experts, hidden_dim] block scales
) {
    // Each token processes its routed expert
    int token_id = blockIdx.x;
    int expert_id = routing_table[token_id];
    
    // Load 2-bit weights: 4 values per byte
    uint8_t packed = w2_experts[expert_offset(expert_id, thread_coord)];
    
    // PRMT instruction extracts and sign-extends 2-bit -> 4-bit
    // Maps {0,1,2,3} -> {-4,-1,1,4} through LUT
    uint32_t unpacked = __prmt_lut(packed, SIGN_SYMMETRIC_LUT);
    
    // Block-scale and feed to QMMA.SF (undocumented FP4 tensor core)
    float scale = scales[expert_id * scale_stride + block_id];
    half fragment[8];
    __qmma_sf_m16n8k32(fragment, activations, unpacked, scale);
    
    // Accumulate to output
    atomicAdd(&output[token_id * hidden_dim + out_col], fragment[i]);
}

The actual kernel runs at 4 CTAs per SM with carefully orchestrated register pressure—Blackwell's 65,536 registers per SM enable keeping entire weight blocks in registers during decompression, avoiding shared memory roundtrips.

Precision recovery happens through a confidence-gated replay mechanism. During forward pass, the routing layer tracks confidence scores for each token's expert selection. When confidence drops below a threshold (indicating the model is uncertain), the kernel bumps an in-graph counter but zeros out that expert's contribution. After the graph completes, the CPU-side runner detects non-zero counters, issues a batched H2D transfer for all missed experts at FP4 precision (51.6 GB/s measured bandwidth), promotes them into the delta cache, and replays the entire CUDA graph. This is the only way to maintain bit-identical outputs under vLLM's graph execution model—you can't partially re-execute a captured graph.

The cache itself exploits MoE routing concentration: benchmarks show 51% expert coverage serves 91% of routing decisions on GLM-5.2. The eviction policy is frequency-based with background promotion (H2D transfers during idle cycles) and CUDA graph-safe updates (no synchronization inside captured graphs). For experts that miss both the 2-bit and FP4 caches, the system batches all misses in a forward pass, transfers them together, and does a single graph replay—amortizing the ~3ms replay overhead across multiple misses.

The project also ships a complete reverse-engineered Blackwell ISA database: 1,994 instruction forms with 128-bit encoding templates, validated across 47,000 instructions extracted from compiled CUDA binaries. This enabled discovering undocumented QMMA.SF type codes for block-scaled FP4 (e2m1) and E3M4 formats that NVIDIA uses internally but doesn't expose through CUDA APIs. The assembler/disassembler toolchain generates the same binary encoding as nvdisasm but supports inline assembly integration with vLLM's C++ extension layer.

One non-obvious optimization: the AFRAG (fragment-major activations) kernel variant reorganizes activation layout so each tensor core load instruction fetches one complete 128-bit fragment instead of gathering across strided addresses. This addresses a load-issue bottleneck—the kernel wasn't DRAM-bound, it was spending cycles waiting for load instructions to issue. Switching to fragment-major layout yielded 1.3× GEMM throughput and +12% end-to-end prefill improvement with bit-identical outputs, demonstrating that memory layout matters as much as compression ratio at this optimization level.

Gotcha

Every line of SASS code in this project is hard-coded for SM120 (consumer Blackwell). If you're running vLLM on H100s, A100s, or even datacenter Blackwell (which may have a different SM version), these kernels won't execute—you'll hit invalid instruction errors at module load time. This isn't a theoretical limitation: the entire value proposition assumes you own RTX 5090s or PRO 6000 cards and want to run models that shouldn't fit. Production ML infrastructure runs on Hopper clusters where this code is dead weight.

The quality evaluation is concerningly thin for a quantization project. The paper reports acceptance rate deltas and shows example outputs, but there are no MMLU scores, no MT-bench results, no perplexity numbers on standard benchmarks. One example shows the model answering "What is the capital of Poland?" with "Krakow" at 2-bit precision, requiring FP4 recovery to fix to "Warsaw"—that's a factual error on a trivial question. The author notes "no systematic quality change observed" when allowing cache misses to zero out expert contributions (miss-tolerance mode), but follows with "quantitative eval pending." For production use, you need hard numbers on how often these errors occur and whether they're acceptable for your workload.

The host-resident mode requires ~200GB of free pinned RAM for GLM-5.2. Pinned memory allocation can fail silently on systems with insufficient physical RAM or competing processes, and the batched H2D transfers add 3ms+ latency spikes on cache misses. The miss-tolerance knob trades correctness for speed—you're explicitly accepting that some expert computations will output zeros when you enable it. The patch also targets vLLM 0.24.0 (January 2025), and the author had to fix multiple broken SM120 support issues in upstream vLLM to make it work. This signals that neither NVIDIA nor the vLLM maintainers are prioritizing consumer Blackwell, so expect this patch to bitrot quickly as vLLM's main branch evolves. Maintaining a 5,600-line assembly-heavy patch across vLLM's rapid release cycle is not a hobby project—it's a part-time job.

Verdict

Use if: You own RTX 5090s or Blackwell PRO cards and need to run DeepSeek-V4, GLM-5.2, or similar frontier MoE models for research, demos, or exploratory fine-tuning where FP8 precision doesn't fit in your VRAM budget. You're comfortable reading assembly, debugging CUDA graph replay issues, and accepting that quality evaluation is incomplete. You have 200GB+ of free RAM and are running single-stream or low-batch inference where the cache eviction policy can shine. You're also prepared to rebase this patch yourself when vLLM updates break compatibility—this is a research artifact, not a maintained product. Skip if: You're running production serving workloads on datacenter GPUs (H100/H200) where these kernels don't execute, you need proven quantization quality with benchmark numbers before deployment, you're serving high-concurrency workloads (100+ streams) where cache thrashing and replay latency will crater throughput, or you want portable code you can modify and iterate on without learning Blackwell SASS. For most ML engineers, wait for official vLLM FP4 support or run smaller models with standard quantization—this tool is a impressive technical flex that solves a problem 99% of practitioners don't have.