FireEcho

FireEcho Engine

High-performance single-GPU inference kernel for 30B+ MoE models

Created by Luis E. Davila Flores

What is FireEcho?

FireEcho is a from-scratch inference engine that runs Qwen3-Omni-30B (30.5 billion parameters, 128-expert MoE) on a single RTX 5090 at 45+ tokens/second using only 20 GB VRAM.

It achieves this through custom Triton kernels that fuse dequantization inside the matmul loop — no separate dequantization step, no global memory writes, no NVIDIA proprietary libraries.

Key Results

Benchmark Results (RTX 5090, 200 tokens/prompt, 8 diverse prompts)

Peak VRAM: 21.5 GB | Baseline forward: ~23.1ms/tok | Full-stack forward: ~17.5ms/tok

Note: L1-L5 show slight overhead vs L0 due to additional dispatch logic. CUDA Graph (L6) eliminates all Python overhead and captures the full 48-layer forward as a single graph replay. The compression layers (FE-XC/INT2) reduce memory bandwidth which compounds with speculative decoding — with a trained EAGLE-3 head at 70% acceptance, projected throughput is ~457 tok/s.

Speed Optimization History

Starting from a naive Python loop over 128 MoE experts (0.4 tok/s), each optimization layer compounds:

The Goliath Kernel: Why It's Fast

Standard quantized inference dequantizes weights to BF16 in global memory, then runs a matmul. This doubles memory traffic.

Goliath FP4 dequantizes inside the Triton matmul tile loop — in registers, with zero global memory writes:

# Simplified Goliath FP4 inner loop
for k_block in range(0, K, BLOCK_K):
    w_packed = tl.load(weight_ptr + offsets)     # Load FP4 packed bytes
    w_lo = (w_packed & 0xF) * scale              # Dequant low nibble in-register
    w_hi = (w_packed >> 4) * scale               # Dequant high nibble in-register
    acc += tl.dot(a_tile, w_tile)                # Tensor core matmul

Packed MoE eliminates the Python expert loop entirely. All 128 experts are packed into one contiguous [128, K//2, N] buffer. A single Triton kernel launch reads expert IDs from a GPU tensor and indexes into the buffer — zero .item() calls, zero CPU-GPU synchronization.

Quantization Formats

Unique Features

Hebbian Memory

Biologically-inspired fast weights that learn during inference (no backpropagation). Implements competitive learning, STDP traces, intrinsic plasticity, PMI correction, and GHA decorrelation. Papers: Lansner BCPNN, Triesch 2005, Sanger's GHA.

Atlas Gatekeeper (FE-AGK)

Runtime expert management:

• Ban & Pick: Profiles expert impact, bans bottom 25% per layer (8->6 effective experts)
• MoDES: Skips entire MoE computation for uncertain tokens (saves ~50% compute on many layers)

FE-XC / INT2 Cold Expert Demotion

Automatically compresses rarely-used experts to 2-bit:

• FE-XC: Codebook 2-bit (2x8 AQLM-style, near-FP16 quality, 5.3x faster kernel)
• INT2: Scalar 2-bit (simple but lower quality)
• Age-adaptive: hot->FP4, cold->FE-XC, coldest->INT2

FlashDecode

Custom Triton attention kernel for M=1 decoding:

• Online softmax (no separate softmax pass)
• Reads only valid KV positions (no padding waste)
• GQA support (4 KV heads -> 32 query heads)
• 15.8ms per token (48 layers, 4096-token context)

EAGLE-3 Speculative Decoding (infrastructure ready)

Draft-then-verify acceleration. Draft head predicts K=5 tokens, target model verifies all 6 in one forward pass. Infrastructure complete, draft head training in progress.

Quick Start

from fireecho_kernel import FireEchoEngine

# Load (streams layer-by-layer, 110s, 20 GB VRAM)
engine = FireEchoEngine.from_pretrained("path/to/Qwen3-Omni-30B")

# Enable optimizations
engine.enable_flat_decode(kv_dtype='fp8')    # FP8 KV cache
engine.enable_cuda_graph_decode()            # CUDA Graph

# Generate
input_ids = engine.tokenizer.encode("Hello, world!", return_tensors='pt').cuda()
output = engine.generate(input_ids, max_new_tokens=200, temperature=0.7)
print(engine.tokenizer.decode(output[0], skip_special_tokens=True))

Requirements

• GPU: RTX 4090 (24 GB) minimum, RTX 5090 (32 GB) recommended
• CUDA: 12.4+
• Python: 3.10-3.12
• PyTorch: 2.4.0+
• Triton: 3.0+
• OS: Linux x86_64

Hardware Independence

FireEcho uses zero NVIDIA proprietary libraries:

• No cuQuantizer, CUTLASS, TensorRT, cuBLAS (except via torch.matmul for attention)
• All custom kernels are pure Triton (compiles to NVIDIA CUDA, AMD ROCm, Intel XPU)
• Runs anywhere Triton runs

Architecture

FireEcho Engine
├── fireecho_kernel.py      # Main engine (9000+ lines)
│   ├── FireEchoEngine      # Load, generate, speculate
│   ├── MoEFFN              # Packed MoE with fused dispatch
│   ├── HebbianMemory       # Fast weights (learn at inference)
│   ├── FlashDecode         # Triton M=1 GQA attention
│   └── CUDA Graph          # Graph-captured decode
├── goliath_kernel.py       # Quantized GEMM kernels (3000+ lines)
│   ├── GoliathFP4          # FP4 fused dequant-matmul
│   ├── GoliathFP8          # FP8 fused dequant-matmul
│   ├── GoliathINT2         # INT2 scalar quantization
│   ├── GoliathFEXC         # FE-XC codebook 2-bit
│   └── Packed MoE          # Contiguous expert buffers
├── triton_hebbian.py       # Fused Hebbian kernels
├── femx_storage.py         # Block floating point storage
└── persistent_memory.py    # AGI-like persistent memory

License

CC BY-NC 4.0 — Free for research and non-commercial use with attribution.

For commercial licensing: @Joysulem on X/Twitter

Citation

@software{fireecho2026,
  author = {Davila Flores, Luis E.},
  title = {FireEcho Engine: High-Performance Single-GPU Inference for 30B+ MoE Models},
  year = {2026},
  url = {https://github.com/Joysulem/FireEcho}
}