Luce-Org/lucebox-hub

Source: https://github.com/Luce-Org/lucebox-hub

Local LLM inference server built for speed. Custom kernels, speculative prefill & decoding, quantized GGUF paths.
Each project is a new optimization to our engine for a specific model family and hardware target.

Projects

Each directory is a self-contained project with setup instructions and benchmark notes.

Supported models

All speedups measured vs vendored llama.cpp (-fa 1, matching KV quant).

01 · Megakernel Qwen3.5 0.8B on RTX 3090

Single-kernel CUDA inference for Qwen 3.5-0.8B on RTX 3090. All 24 layers run in one persistent dispatch.

# 1. clone + enter
git clone https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/megakernel

# 2. install (Python 3.10+, CUDA 12+, PyTorch 2.0+). Weights stream from HF on first run.
python -m venv .venv && source .venv/bin/activate   # required on Ubuntu 24+ system Python (PEP 668)
pip install --upgrade pip
pip install torch                          # install BEFORE the next step; setup.py imports torch at build time
pip install -e . --no-build-isolation      # --no-build-isolation lets the build see the torch you just installed

# 3. run the benchmark (prefill pp520 + decode tg128 vs llama.cpp BF16 + PyTorch HF)
python final_bench.py

Implementation notes: 82 blocks, 512 threads, cooperative grid sync, no CPU round trips between layers, and weights streamed from Hugging Face on first run.

Full writeup → · Benchmarks → · Blog post →

Blackwell (RTX 5090, DGX Spark / GB10): auto-detected by setup; NVFP4 decode path lands ~194 tok/s tg128 on GB10. See megakernel/README.md#blackwell-sm_120–sm_121a.

02 · DFlash DDtree Qwen3.5 & Qwen3.6 27B GGUF on RTX 3090

DFlash speculative decoding for Qwen3.5/Qwen3.6 27B GGUF targets on a single GPU. The default setup uses Qwen3.6-27B Q4_K_M plus the Lucebox Q8_0 GGUF DFlash draft.

• Up to 207 tok/s in the demo (207.6 tok/s DFlash vs 38.0 tok/s AR, 5.46×)
• 129.5 tok/s mean on the HumanEval 10-prompt bench
• 3.43× faster than autoregressive (+15% over chain speculative decoding)
• 2.8× faster than SGLang AWQ on the same hardware
• Up to 256K context in 24 GB via TurboQuant TQ3_0 KV cache (128K Q4_0 bench: 134.78 tok/s at ctx=131072)

# 1. clone with submodules (pulls the pinned Luce-Org/llama.cpp@luce-dflash fork)
git clone --recurse-submodules https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/dflash

# 2. build the C++/CUDA decoder (CUDA 12+, CMake 3.18+)
# Default compiles for Pascal/Volta/Turing/Ampere (60/61/62/70/75/86; +120 on CUDA 12.8+, +sm_121/DGX Spark on CUDA 12.9+, +sm_110/Thor on CUDA 13.0+) so the binary runs on every supported card.
# 3090-only users can add -DCMAKE_CUDA_ARCHITECTURES=86 to skip the other archs and build faster (~3 min).
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release
cmake --build build --target test_dflash -j
cmake --build build --target test_generate -j

# 3. fetch weights: ~16 GB Q4_K_M target + 1.84 GB Lucebox Q8_0 GGUF DFlash draft
hf download unsloth/Qwen3.6-27B-GGUF Qwen3.6-27B-Q4_K_M.gguf --local-dir models/
hf download Lucebox/Qwen3.6-27B-DFlash-GGUF dflash-draft-3.6-q8_0.gguf --local-dir models/draft/

# 4a. one-shot streaming generate
python3 scripts/run.py --prompt "def fibonacci(n):"

# 4b. or reproduce the paper-style bench (HumanEval + GSM8K + Math500, ~15 min)
python3 scripts/bench_llm.py

Why GGUF/Q4_K_M: on 24 GB GPUs, the target, draft, DDTree verify state, and KV cache need to fit together. The default Qwen3.6 setup uses a ~16 GB Q4_K_M target and a 1.84 GB GGUF draft.

Algorithms used:

• DFlash (z-lab, 2026): block-diffusion draft conditioned on target hidden states.
• DDTree (Ringel et al., 2026): tree-structured verify that beats chain verify at the same compute budget.

Implemented here:

• C++/CUDA decode engine on top of ggml (no libllama, no Python runtime, Q4_K_M target path).
• Three custom CUDA kernels for tree-aware SSM state rollback: ggml_ssm_conv_tree, ggml_gated_delta_net_tree, ggml_gated_delta_net_tree_persist.
• DDTree budget swept for RTX 3090 + Q4_K_M target: budget=22 is the sweet spot.
• TQ3_0 KV cache (TurboQuant 3.5 bpv, default) + sliding target_feat ring to fit up to 256K context in 24 GB (Q4_0 available as legacy, tops out near 128K).

Running on other GPUs (4090, 5090, DGX Spark / GB10, Jetson AGX Thor)

Supported out of the box; the build just needs the right CUDA toolkit. dflash/CMakeLists.txt already auto-adds Blackwell archs when your nvcc is new enough, so the main quickstart above works as-is on newer cards.

Verify your target:

python -c "import torch; p=torch.cuda.get_device_properties(0); print(p.name, 'sm_%d%d'%(p.major,p.minor), p.multi_processor_count,'SMs', round(p.total_memory/1e9,1),'GB')"
nvcc --version

DGX Spark / GB10 quick start:

# CUDA 12.9+ required for sm_121
nvcc --version  # must show >= 12.9
git clone --recurse-submodules https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/dflash
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release   # CMake auto-adds sm_121
cmake --build build --target test_dflash -j

Jetson AGX Thor quick start:

# CUDA 13.0+ required for sm_110 / AGX Thor.
nvcc --version
git clone --recurse-submodules https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/dflash
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release   # CMake auto-adds the Thor arch your nvcc supports
cmake --build build --target test_dflash -j

Retune per GPU:

• DDTree budget=22 tuned for 3090 + Q4_K_M + 24 GB. On the RTX 5090, budget=40 is optimal (swept). On GB10 (128 GB unified), re-sweep — larger tree = more verify throughput until memory bandwidth saturates. scripts/bench_llm.py --budget N has the sweep hooks.
• TQ3_0 KV cache + sliding target_feat ring was shaped by 24 GB (fits up to 256K context on a 3090). On GB10 (128 GB unified) / 5090 (32 GB) you can push context further or skip quantization entirely and keep F16 KV.
• Perf numbers (207 tok/s demo, 129.5 HumanEval, 2.8× vs SGLang AWQ) are RTX 3090 @ stock. RTX 5090 numbers (205 tok/s HumanEval, 4.84×) are in RESULTS.md. Ada/GB10/Thor not yet swept, PRs with RESULTS.md entries welcome.

Full writeup → · Benchmarks → · Blog post →

03 · PFlash speculative prefill on RTX 3090

Speculative prefill for long prompts. A Qwen3-0.6B BF16 drafter scores token importance, then the 27B target prefills only the retained spans. Runtime is C++/CUDA through the dflash binaries; no PyTorch is required at serving time.

• ~10.4× TTFT on 128K context: 24.8 s dflash daemon vs ~257 s llama.cpp (FA on, Q4_0 KV).
• 10.0× TTFT on 64K context: 13.5 s dflash vs 134.95 s llama.cpp.
• NIAH single-needle retrieved at every measured context (32K → 128K), keep_ratio=0.05, DFLASH_FP_ALPHA=0.85.

# 1. build dflash + BSA kernel (sm_80+ required for BSA, ~10 min cold compile)
git clone --recurse-submodules https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/dflash
cmake -B build -S . -DCMAKE_BUILD_TYPE=Release \
                    -DCMAKE_CUDA_ARCHITECTURES=86 \
                    -DDFLASH27B_ENABLE_BSA=ON
cmake --build build --target test_dflash test_flashprefill_kernels -j

# 2. fetch weights: 27B Q4_K_M target + 0.6B BF16 drafter (GGUF) + DFlash spec-decode draft
hf download unsloth/Qwen3.6-27B-GGUF Qwen3.6-27B-Q4_K_M.gguf --local-dir models/
hf download unsloth/Qwen3-0.6B-GGUF Qwen3-0.6B-BF16.gguf --local-dir models/
hf download Lucebox/Qwen3.6-27B-DFlash-GGUF dflash-draft-3.6-q8_0.gguf --local-dir models/draft/

# 3. run the daemon: compress (drafter scoring) + generate (target spec decode)
DFLASH_FP_USE_BSA=1 DFLASH_FP_ALPHA=0.85 \
./build/test_dflash models/Qwen3.6-27B-Q4_K_M.gguf models/draft/dflash-draft-3.6-q8_0.gguf --daemon
# stdin protocol: `compress <ids.bin> <keep_x1000> <drafter.gguf>` →
#                 stream of compressed token ids, then `generate <…>` →
#                 stream of generated tokens.

Daemon stdin commands: compress runs the drafter with FlashPrefill block-sparse attention and returns the compressed token-id stream; generate runs the target on that stream with normal speculative decode + DDTree. park / unpark / free drafter swap weights in and out of VRAM so target + drafter coexist on a 24 GB card.

Runtime tunables (full list in dflash/src/flashprefill.h):

DFLASH_FP_USE_BSA=1     # dispatch sparse FA forward through BSA (sm_80+)
DFLASH_FP_ALPHA=0.85    # block-selection threshold; higher = stricter = fewer K-blocks per Q-row
DFLASH_FP_PROFILE=1     # log mean / score / select / forward stage timings

What’s ours, what isn’t. Algorithms are from Cross-Family Speculative Prefill (Liu et al., ICLR 2026) for the scoring + selection layer and FlashPrefill (Fan et al., 2026) for the drafter sparse-attention forward. What we built:

• C++/CUDA daemon-resident speculative prefill in front of a quantized GGUF target — no PyTorch, no Triton, no per-request subprocess.
• BSA wired without libtorch via a 3-header ATen/c10 stub set under dflash/deps/bsa_stubs/.
• Custom Qwen3-0.6B forward (qwen3_0p6b_*) so the drafter runs through the same ggml allocator as the 27B target.
• 4 CUDA kernels (flashprefill_kernels.cu) for the FlashPrefill mean_K / score / select / sparse_fwd algorithm.

Full writeup → · Daemon-side build / tunables → · Blog post →

AMD Strix Halo (HIP backend)

Same DFlash + PFlash stack on an AMD iGPU. PR #119 ports the Phase 2 rocWMMA flashprefill kernels to HIP. End-to-end on a single Ryzen AI MAX+ 395 box (Radeon 8060S iGPU, gfx1151, 128 GiB LPDDR5X-8000 unified): 37.0 tok/s DFlash decode on Qwen3.5-27B Q4_K_M, 27.6 s TTFT at 16K context with NIAH retrieval intact. That is 3.08× decode and 2.24× prefill over llama.cpp HIP AR on the same iGPU. End-to-end wall clock at a realistic 16K prompt + 1K generation workload: 2.66× faster than vanilla llama.cpp.

git clone --recurse-submodules https://github.com/Luce-Org/lucebox-hub && cd lucebox-hub/dflash

# Build for gfx1151 (Strix Halo). Swap the arch for gfx1100 / gfx1201.
cmake -B build -S . \
  -DCMAKE_BUILD_TYPE=Release \
  -DDFLASH27B_GPU_BACKEND=hip \
  -DDFLASH27B_HIP_ARCHITECTURES=gfx1151 \
  -DDFLASH27B_HIP_SM80_EQUIV=ON
cmake --build build --target test_dflash -j

DFLASH27B_HIP_SM80_EQUIV=ON enables the rocWMMA Phase 2 flashprefill kernels (the path that delivers the prefill speedup). OFF falls back to ggml’s flash_attn_ext (slower but no rocwmma headers needed).

Per-arch DDTree tuning: gfx1151 (Strix Halo iGPU, bandwidth-bound on LPDDR5X) peaks at --ddtree-budget=22. gfx1100 (7900 XTX, GDDR6) prefers budget=8 per the PR #156 cross-arch perf plan. Run scripts/bench_he.py --ddtree-budget N to verify on your card.

Drafter recipe for max decode: target = Qwen3.5-27B Q4_K_M, drafter = same gen quantized to Q8_0 via dflash/scripts/quantize_draft_q8.py. The matching Q8_0 GGUF on the unsloth Qwen3.6 target needs DFLASH27B_DRAFT_SWA=2048 for sliding-window correctness.

Blog post → · PR #119 → · PR #156 cross-arch perf plan →

Why this exists

Local AI should be a default, not a privilege: private data, no per-token bill, no vendor lock-in. The hardware to run capable models already sits on desks. The software to run those chips well doesn’t.

General-purpose frameworks dominated the last decade because hand-tuning kernels per chip was too expensive to justify. One stack, decent on everything, great on nothing. Most of the silicon’s capability stays on the floor.

AI-assisted development flips that calculus. Rewrites that took a quarter now fit in a release cycle. Lucebox is where we publish them, one chip and one model family at a time. Apache 2.0 source, full writeup, reproducible benchmarks.

Requirements

All experiments in this repo are built, tuned, and benchmarked on NVIDIA RTX 3090 (2020), the reference target. Supported GPU families:

• Ampere (sm_86, RTX 3090 / A-series): reference, CUDA 12+.
• Ada (sm_89, RTX 40xx): should work, unverified, CUDA 12+.
• Blackwell consumer (sm_120, RTX 50xx incl. 5090): supported, CUDA 12.8+.
• DGX Spark / GB10 (sm_121, compute capability 12.1): supported, CUDA 12.9+.
• Jetson AGX Thor (sm_110): supported, CUDA 13+.
• Turing (sm_75, RTX 2080): supported, CUDA 12+.

PyTorch 2.0+. dflash/ needs CMake 3.18+ and --recurse-submodules for the pinned Luce-Org/llama.cpp@luce-dflash fork (three tree-mode ggml ops); multi-arch build is automatic (see Running on other GPUs).

Megakernel porting note. megakernel/setup.py auto-detects the GPU arch and SM count at build time via torch.cuda.get_device_capability(). The decode grid is persistent (one block per SM) and is clamped to the resident-block ceiling at runtime, so no manual tuning is needed. On SM < 80 (Turing), the kernel uses FP16 instead of BF16 via a compile-time TARGET_SM flag; on SM ≥ 80 (Ampere+), BF16 is used. Just pip install -e . --no-build-isolation and the right code path is selected automatically.

Optional, find your GPU’s sweet spot: sudo nvidia-smi -pl 220 (megakernel hits best tok/J at 220 W on 3090; re-sweep for other cards).

Repository layout

lucebox-hub/
├── megakernel/    · fused forward pass for Qwen 3.5-0.8B
├── dflash/        · DFlash speculative decoding port for Qwen 3.5/3.6-27B on RTX 3090
├── pflash/        · speculative-prefill harness in front of dflash (12.5× TTFT at 128K)
└── assets/        · banners, cards, diagrams

Roadmap

  Q1 2026    ▮▮▮▮▮▮▮▮▮▮    RTX 3090 kernels & optimizations
  Q2 2026    ▮▮▮▮▮▯▯▯▯▯    Ryzen AI MAX+ 395 optimizations
  Q2 2026    ▮▮▯▯▯▯▯▯▯▯    Heterogeneous CPU + GPU latency optimizations
  Q2 2026    ▮▯▯▯▯▯▯▯▯▯    Lucebox OS for local AI machines
  Q3 2026    ▯▯▯▯▯▯▯▯▯▯    Lucebox official launch

Citation

@software{lucebox_2026,
  title  = {Lucebox: Open LLM Inference, Rewritten by Hand for One Specific Chip at a Time},
  author = {Lucebox},
  url    = {https://github.com/Luce-Org/lucebox-hub},
  year   = {2026}
}

Per-project citations live in each subproject’s README.

Inspired by

• Hazy Research: megakernel idea and the intelligence-per-watt methodology.
• z-lab/DFlash (Wang et al., 2026): block-diffusion speculative decoding algorithm. We use their published Qwen3.5/Qwen3.6-27B-DFlash draft weights as-is.
• DDTree (Ringel & Romano, 2026): tree-structured verify that DFlash 27B uses for its 3.5× speedup over chain spec decoding. liranringel/ddtree.
• AlpinDale/qwen_megakernel, Infatoshi/MegaQwen: prior art on fused Qwen kernels.

Community

• Discord: discord.gg/yHfswqZmJQ
• Website: lucebox.com
• Issues: github.com/Luce-Org/lucebox-hub/issues
• Blog: lucebox.com/blog

_{Apache 2.0 · Lucebox.com}