Appeared that 4bit kv-cache attention generation speed is almost twice slower on 3090 than fp8 - tried both turboquant_4bit_nc and int4_per_token_head. Taking a look into --enable-experts-parallel mode with group-size 128 for experts to get more space for fp8 KV cache.

update: i have tested with gs=128 for expert modules that work only with --enable-expert-paralleland the available context-size with fp8 kv cache with--gpu-memory-utilization 0.945 is:

• current checkpoint (tp over expert modules as well): 145K
• gs128 checkpoint (ep): 175K, but in my case token generation slower by 10%

also appeared that autoround RTN quantization does not require patched vllm for fp8 kv and takes around 10 minutes, does not require complex hardware, preserves MTP layer. receipe for gs 128 experts:

auto_round /home/ubuntu/models/Hy3-tencent \
  --model_free \
  --scheme W4A16 \
  --group_size 128 \
  --layer_config "{ '.*shared_mlp':{'group_size':64}, 'experts':{'group_size':128}, '.*eh_proj':{'bits':16,'data_type':'float'} }" \
  --ignore_layers "model.layers.0" \
  --output_dir /home/ubuntu/models/autoround/tencent-Hy3-RTN-EP

to preserve speculative decoder quality ignore layer 80 as well.
for me MTP works slower and uses more VRAM, so less space for kv cache.

Would you release this AutoRound version with gs=128? Using turboquant here degrades the speed too much (70% slower than FP8), so it would be wonderful to have this version with more space for FP8 kv cache

done: avtc/Hy3-autoround-RTN-4bit-tp8-ep

Thank you so much! I am going to test it right now

Tencent released their own GPTQ variant. What do you say about it? https://huggingface.co/AngelSlim/Hy3-GPTQ-Int4
Is it beneficial to use this one ?

In my hardware, with the Tencent model I would get 96k maximum of kv cache. With avtc's model I get 163k. So it's a considerable diference. I wouldnt be able to tell the difference in accuracy between this too models but I think it wouldn't be very different.

I just ran the AA_LCR benchmark with the avtc model and it got 74% which is on par with the BF16. I am trying to run other benchmarks to assess the inteligence but the AA_LCR score was very good. ** I need to make a correction. When I say avtc model I mean the AutoRound version with FP8 kvcache, not this GPTQ version. I am running the SWE-bench right now.

** I need to make a correction. When I say avtc model I mean the AutoRound version with FP8 kvcache, not this GPTQ version.

I have chery-picked commit to allow using KVarN near lossless 3bit cache, it seems it does not hurt token generation speed, but have not tested yet with large context length (on 10-30% context length I see parity with fp8 in speed).
the branch to build vllm from (readme updated): https://github.com/avtc/vllm/tree/fix/load-quantized-hy3-mimo-v2.5-kvarn-0.26

Excelent!! Do you have a command/script example to run this 3bit cache?

Yes, i have updated readme - so need to build from another branch and run with --kv-cache-dtype kvarn_k4v2_g128

Perfect! Thanks!!

It works like a charm. Does it work with the AutoRound model or the AngelSlim GPTQ model?

work with the AutoRound model - yes
AngelSlim GPTQ - should work - if prev.branch worked for you with it.

@avtc sorry for going offtopic, what are your plans on further models? what about deepseek 4 flash final release or mimo 2.5?

@maglat ,
deepseek v4 flash does not work in vllm on RTX 3090, I have researched that, and there is no easy way to make it work.

mimo v2.5 - i have tried to create RTN quant - and it output gibberish, then I have converted https://huggingface.co/spectator2026/MiMo-V2.5-AWQ-int4 to GPTQ weights format and was able to load it with 279K context window, and output was +- fine. But every second turn it output-ed reasoning in bare output. And for tasks that hy3 RTN close in 2min with 60t/s - Mimo runs 10mins with 90t/s, consuming very large amount of context window. And on trying to make a simple one-shot airplane prompt - it, before presenting me result - realized that there were issues - and rewrote the file - and goes in circles idk how many times never stopping. So this converted quant does not look like good to be published, (or maybe it is a model normal behavior). The branch (in readme) able to load the MimoV2.5 AWQ quant, and GPTQ format as well. I do not plan to make calibrated GPTQ quant for it yet, but maybe will change my mind and will try in future.

solar-open2 - promising model size, but for me it hallucinated very much from first prompts, - the quant for it already available from the nota-ai, context size that could fit on 8x3090 = 1157K (max 1M).

Right now it seems that Hy3 is the best model for 8x RTX3090, right?

I am using mostly Qwen3.6-27B full precision, for verification, code-research tasks seems not bad, and rigorous. Hy3 looks much smarter for me, but less following instructions. So depending on task. Idk if there are better models exists for this setup, monitoring reddit+youtube periodically ).

The new Deepseek 4 Flash will be insane! https://x.com/deepseek_ai/status/2083084415157022911?s=20 What they have posted today looks like the new to go model. Sadly without vision :/

@DoktorMincs
Regarding evals on hy3 - the model has 3 thinking levels: none, low, high. When you have run benchmarks (including SWE-bench) - which thinking level was active? and was --default-chat-template-kwargs '{"interleaved_thinking": true}' \ this active as well?

Yes. The flag was: --default-chat-template-kwargs '{"interleaved_thinking": true, "thinking": true, "reasoning_effort": "high"}' . The SWE-bench results that I got was just not reliable - I had a lot of trouble trying to run them, so thats why I didnt post the results. But the AA_LCR result seems all right.

Here is a recipe how to run Inkling-Small awq int4 ( https://huggingface.co/cyankiwi/Inkling-Small-AWQ-INT4 ) on 8x3090.

Follow readme first, then:

# switch the editable install onto the mixed-precision branch
git remote add avtc https://github.com/avtc/vllm.git
git fetch avtc feature/inkling-w4a16-mixed-precision-ampere
git checkout feature/inkling-w4a16-mixed-precision-ampere

example start command:

export TORCH_CUDA_ARCH_LIST="8.6"
export CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7
export RAY_memory_monitor_refresh_ms=0
export NCCL_CUMEM_ENABLE=0
export VLLM_ENABLE_CUDAGRAPH_GC=1

export VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS=0
export VLLM_ALLOW_LONG_MAX_MODEL_LEN=1
export VLLM_MARLIN_USE_ATOMIC_ADD=1
export VLLM_FLOAT32_MATMUL_PRECISION=high
export OMP_NUM_THREADS=1

vllm serve /home/ubuntu/models/Inkling-Small-AWQ-INT4-cyankiwi \
    -tp 8 \
    -dp 1 \
    --served-model-name ink-small \
    --port 8000 \
    --host 0.0.0.0 \
    --uvicorn-log-level info \
    --trust-remote-code \
    --gpu-memory-utilization 0.87 \
    --seed 1234 \
    --enable-sleep-mode \
    --skip-mm-profiling \
    -O3 \
    --no-use-tqdm-on-load \
    --tokenizer-mode inkling \
    --tool-call-parser inkling \
    --reasoning-parser inkling \
    --enable-auto-tool-choice \
    --performance-mode balanced \
    --enable-chunked-prefill \
    --max-num-seqs 4 \
    --max-model-len auto \
    --default-chat-template-kwargs '{"interleaved_thinking": true, "preserve_thinking": true}' \
    --kernel-config.enable_flashinfer_autotune=False \
    --enable-prefix-caching \
    --max-num-batched-tokens 2048

Notes:
fp16 kv-cache gives 414912 context size, vram usage 23.85GB per card. (lower --gpu-memory-utilization in case of cuda oom in the process)
fp8 or KVarN are not supported.

@avtc
What is your take on Inkling Small? Is it better than Hy3?

@DoktorMincs - i have tried only AWQ quant for now, did not tested autoround RTN yet.
The AWQ quant at first seemed to me not bad with one-shot short outputs. But for agentic tasks it doesn't work well for me, much worse than Hy3, requiring much more hand-guiding from me and still producing bad results. While Hy3 without hand-guiding done everything well.

I was able to run the original DeepSeek V4 Flash 0731 in vllm with the RTX3090 (I actually used 10x 3090s to get to 1M context) with help of Fable 5. I asked it to create a guide on how to make it run:

Running DeepSeek-V4-Flash on RTX 3090 GPUs — Complete Guide

A step-by-step guide to serving DeepSeek-V4-Flash-0731 on consumer RTX 3090 cards with vLLM.

Everything here was validated in production on 2026-08-08 on the exact hardware described below.

1. Why this is hard (read this first)

DeepSeek-V4-Flash is built for datacenter GPUs. Three hardware features it expects simply do not exist on the RTX 3090 (chip architecture "Ampere", compute capability SM86):

Because of this, stock vLLM refuses to run the model on a 3090. The gate is explicit in the source (sparse_mla.py: capability must be 9.x or 10.x). No command-line flag gets around missing silicon.

The solution is a community fork, Lvllmds4-x (by guqiong96, the author of Lvllm), which rewrites the entire DS4 kernel stack with portable fallbacks: Triton kernels that upcast FP8→BF16 before the matrix multiply, Marlin kernels for the MXFP4 experts, and tilelang JIT kernels for the mHC layers. It also bundles lk_moe, a NUMA-aware CPU/GPU MoE engine.

2. What you need

Hardware (what we used)

• Machine 1: 8× RTX 3090 (24 GB each), 56-core Xeon, 120 GB RAM
• Machine 2: 2× RTX 3090 (+ other cards, excluded), connected to machine 1 by a fast link (we use 100G RoCE; 10G Ethernet also works, slower boot)
• ~160 GB free disk on each machine (the model must exist on both)

Can I run it on 8× 3090 (one machine)? We tried hard, and the answer is yes, but with some limitations. With the original DeepSeek model weights shard to 20.3 GiB/GPU with TP=8, and the KV cache needed for even a 128k context (1.4 GiB/GPU) plus activation peaks does not fit in the remaining ~2 GiB. Every boot ends in CUDA/Triton OOM during warmup. CPU-offloading experts works but destroys decode speed. 10 GPUs (two nodes) is the practical minimum for full context. With 10 GPUs the weights drop to ~16 GiB/GPU and everything fits — we serve the full native 1M-token context. Maybe you might be able to run it with quantized versions that will give you more memory headroom to load the weights with acceptable context length.

Software

• Linux with NVIDIA driver working (nvidia-smi shows all GPUs)
• CUDA toolkit with a real nvcc (we use 13.3) — the fork JIT-compiles kernels at startup and will crash without nvcc (see §7, trap 3)
• Python 3.12 — the SAME micro-version on both machines (e.g. 3.12.13 everywhere; Ray refuses to form a cluster across 3.12.3 vs 3.12.13 — §7 trap 6)
• The model checkpoint: deepseek-ai/DeepSeek-V4-Flash-0731 (48 safetensors shards, 156 GB)

3. Install the fork (both machines, identical)

Create a dedicated virtualenv — do not install the fork into an existing vLLM environment, it replaces vLLM entirely:

mkdir -p /root/lvllmds4x
python3.12 -m venv /root/lvllmds4x/venv
V=/root/lvllmds4x/venv/bin

# 1. torch — PIN THE VERSION. Both machines must match exactly.
$V/pip install "torch==2.11.0+cu130" --index-url https://download.pytorch.org/whl/cu130

# 2. the fork wheel (download from the Lvllmds4-x GitHub releases page)
$V/pip install lvllmds4_x-2.3.9-cp312-cp312-manylinux_2_34_x86_64.whl

# 3. ray, for the two-node cluster — same version everywhere
$V/pip install ray==2.56.1

# sanity check — should print matching versions on BOTH machines:
$V/python -c "import sys,torch,vllm,ray; print(sys.version.split()[0], torch.__version__, vllm.__version__, ray.__version__)"

The wheel pulls in lk_moe, flashinfer, tilelang, and Triton automatically. It reports itself as vllm 2.3.9.

Copy the model to the same path on both machines and verify byte counts match (du -sb on each side, or use rsync with checksums).

4. Two-node topology

64 attention heads constrain tensor parallelism to TP ∈ {1, 2, 4, 8, 16...}.
With 10 GPUs the only workable geometry is:

TP=2 × PP=5    (2-way tensor parallel × 5 pipeline stages = 10 workers)
ranks 0-7  → machine 1 (8 GPUs, pipeline stages 0-3)
ranks 8-9  → machine 2 (2 GPUs, pipeline stage 4)

Consequence: DSpark speculative decoding (the model's built-in draft) is incompatible with pipeline parallelism — the fork raises NotImplementedError if you try. You give up the spec-decode speedup in exchange for fitting the model at all. Measured impact: ~37 tok/s instead of a projected ~45-50.

5. Launch scripts

Machine 1 — head + server (lvl_cluster.sh)

The key parts, annotated. Every env var here is load-bearing — each one is a crash we actually hit (§7):

#!/bin/bash
set -e
V=/root/lvllmds4x/venv/bin
MODEL=/root/models/deepseek-ai/DeepSeek-V4-Flash-0731

# CUDA_HOME must contain bin/nvcc (tilelang JIT) — trap 3
export CUDA_HOME=/usr/local/cuda-13.3
export PATH=$CUDA_HOME/bin:$PATH
export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH
export CUDA_DEVICE_ORDER=PCI_BUS_ID
export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True

# lk_moe knobs — traps 1 and 2 (deadlocks!)
export LVLLM_MOE_NUMA_ENABLED=1
export LK_THREADS=12                # per WORKER; keep <= cpu_count / local_workers
export OMP_NUM_THREADS=12
export LK_THREAD_BINDING=NUMA_NODE  # NOT CPU_CORE inside a container — trap 1
export LK_POWER_SAVING=0            # 1 deadlocks the init barrier — trap 2
export LVLLM_GPU_RESIDENT_MOE_LAYERS="0-42"   # all 43 MoE layers stay in VRAM
export LVLLM_GPU_PREFETCH_WINDOW=1
export LVLLM_GPU_PREFILL_MIN_BATCH_SIZE=128
export FLASHINFER_DISABLE_VERSION_CHECK=1

# networking for the cluster (adapt interface names/IPs to your setup)
export NCCL_SOCKET_IFNAME=mlx0
export GLOO_SOCKET_IFNAME=mlx0
export TP_SOCKET_IFNAME=mlx0
export RAY_ADDRESS=10.0.0.10:6379
# (add NCCL_IB_* here if you have RDMA/RoCE; NCCL_IB_DISABLE=1 for plain TCP)

# start the Ray head, then WAIT until the second machine joins (10 GPUs total)
$V/ray start --head --node-ip-address 10.0.0.10 --port 6379
echo "Now run the worker script on machine 2, waiting for 10 GPUs..."
until [ "$($V/python -c "import ray; ray.init(address='auto', logging_level=40); print(int(ray.cluster_resources().get('GPU',0)))")" = "10" ]; do sleep 10; done

exec $V/vllm serve "$MODEL" \
  --host 0.0.0.0 --port 3000 \
  --served-model-name deepseek-v4-flash \
  --trust-remote-code \
  --tokenizer-mode deepseek_v4 \
  --tensor-parallel-size 2 \
  --pipeline-parallel-size 5 \
  --distributed-executor-backend ray \
  --dtype bfloat16 \
  --kv-cache-dtype fp8_ds_mla \
  --gpu-memory-utilization 0.90 \
  --max-model-len 1048576 \
  --max-num-seqs 2 \
  --enable-chunked-prefill \
  --max-num-batched-tokens 2048 \
  --enable-prefix-caching \
  --compilation_config.cudagraph_mode FULL_DECODE_ONLY \
  --enable-auto-tool-choice \
  --tool-call-parser deepseek_v4 \
  --reasoning-parser deepseek_v4 \
  --disable-custom-all-reduce

Flag notes:

• --kv-cache-dtype fp8_ds_mla — the DSA compressed KV layout: 584 bytes per token per layer (~25 KB/token total). This is why a 1M context is even possible: 1M tokens ≈ 25 GB of KV spread across 10 GPUs.
• --max-model-len 1048576 — full native context. Requires ~2.4 GiB of KV headroom per GPU, which exists at 16 GiB weights/GPU. (On the failed 8-GPU attempt even 131072 didn't fit.)
• --max-num-batched-tokens 2048 — chunked-prefill chunk size. Bigger = faster prefill but bigger activation peaks. 2048 is safe; long prompts prefill at ~2,600 tok/s regardless (a 120k prompt takes ~45 s).
• --compilation_config.cudagraph_mode FULL_DECODE_ONLY — CUDA graphs for decode only; avoids capture-time OOM on 24 GB cards.
• --disable-custom-all-reduce — required with expandable_segments on multi-GPU boxes (IPC handle incompatibility).

Machine 2 — worker (ray2.sh)

#!/bin/bash
set -e
V=/root/lvllmds4x/venv/bin

# select ONLY the 3090s if the box has mixed GPUs (ours: 3090 at PCI 0 and 2)
export CUDA_DEVICE_ORDER=PCI_BUS_ID
export CUDA_VISIBLE_DEVICES=0,2

export CUDA_HOME=/usr/local/cuda-13.3
export PATH=$CUDA_HOME/bin:$PATH
export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH
export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True

# same lk_moe env as machine 1 (workers inherit at spawn, set defensively)
export LVLLM_MOE_NUMA_ENABLED=1
export LK_THREADS=12
export OMP_NUM_THREADS=12
export LK_THREAD_BINDING=NUMA_NODE
export LK_POWER_SAVING=0
export LVLLM_GPU_RESIDENT_MOE_LAYERS="0-42"
export FLASHINFER_DISABLE_VERSION_CHECK=1
export NCCL_SOCKET_IFNAME=mlx0
export GLOO_SOCKET_IFNAME=mlx0

$V/ray start --address 10.0.0.10:6379 --node-ip-address 10.0.0.11

Boot order

1. Machine 1: bash lvl_cluster.sh (starts head, waits)
2. Machine 2: bash ray2.sh (joins; machine 1 script proceeds automatically)
3. Wait ~10 minutes: 48 shards load → lk_moe repacks the MXFP4 experts (CPU-heavy, the log goes quiet — this is normal, see trap 5) → tilelang JIT compile → KV allocation → warmup → port 3000 opens.

Success looks like:

GPU KV cache size: 1,141,979 tokens
Maximum concurrency for 1,048,576 tokens per request: 1.09x
Application startup complete.

6. Verify

curl -s localhost:3000/v1/models        # shows max_model_len: 1048576

curl -s localhost:3000/v1/chat/completions -H 'Content-Type: application/json' -d '{
  "model": "deepseek-v4-flash",
  "messages": [{"role": "user", "content": "Count from 1 to 30."}],
  "max_tokens": 200, "temperature": 0.3,
  "chat_template_kwargs": {"enable_thinking": false}
}'

Measured results on this setup:

7. The trap list — every way this failed before it worked

Each of these cost us a failed boot. In order of appearance:

1. lk_moe deadlock #1 — LK_THREAD_BINDING=CPU_CORE inside LXC/containers.
   lk_moe computes CPU core IDs arithmetically and pins pool threads withsched_setaffinity. If your container's cpuset has holes (ours: 96 CPUs with gaps), pinning fails for some threads (sched_setaffinity failed: Invalid argument spam in the log), the pool never assembles, and the MOE_MXFP4 constructor waits on its init barrier forever. The main thread spins at 100% in sched_yield; the log stops right after Backend_NUMA init. Fix: LK_THREAD_BINDING=NUMA_NODE. Diagnosis tool: py-spy dump --pid <worker> — stuck in _process_mxfp4.

2. lk_moe deadlock #2 — LK_POWER_SAVING=1. Puts pool threads to sleep on futexes between jobs; under concurrent constructor init the wakeup gets lost and the barrier hangs the same way. Fix: LK_POWER_SAVING=0.

3. Missing nvcc — [Errno 2] No such file or directory: '.../bin/nvcc'. The mHC layers use tilelang, which JIT-compiles CUDA at startup. A runtime-only CUDA install (no toolkit) crashes the boot late, after weights are already loaded. Fix: point CUDA_HOME at a full toolkit.

4. Oversubscribed thread pools. LK_THREADS is per worker. 8 local workers × 56 threads = 424 threads fighting over 96 CPUs = spinlock convoy that looks like a hang. Fix: LK_THREADS ≈ cpu_count / local_workers.

5. "The log went quiet, is it stuck?" After shard loading, lk_moe repacks every expert tensor on CPU — several minutes of silence with high CPU and 0% GPU. Check worker CPU% before assuming a hang: busy = fine, and Initialized lk_moe with N experts for layer ... lines confirm progress.

6. Ray version paranoia (two-node only). Ray requires exact match of Python micro-version (3.12.13 ≠ 3.12.3 → join refused), Ray version, and you want matching torch too (pip will happily resolve a newer torch on the second machine — pin it).

7. vLLM memory profiler vs JIT compilers. The one-time tilelang/Triton compile spike during profile_run gets accounted as permanent overhead, which can make the auto-sized KV cache negative ("No available memory for cache blocks" even though nvidia-smi shows gigabytes free). Workaround if you hit it: set --kv-cache-memory-bytes explicitly (bytes per GPU); this skips profiling entirely.

8. Killing vLLM leaves GPU-holding orphans. Workers rename themselves (VLLM::Worker_TP0, ray::RayWorkerProc.run), so pkill -f vllm misses them and the next boot hits "GPUs busy". Kill in this order:

   pkill -9 -f '[b]in/vllm serve'
   ps -eo pid,args | grep -E '[R]ayWorkerProc.run|[V]LLM::' | awk '{print $1}' | xargs -r kill -9
   

8. Alternatives if you have different hardware

• RTX 4090 (SM89): has FP8 tensor cores. The fork's Ada path is faster; there is also a separate vllm-deepseek-v4-sm89 repo gated to exactly 8.9.
• A100 (SM80): same missing-FP8 story as the 3090. Documented working path: SGLang 0.5.16 + a monkeypatch converting FP8→BF16 offline (HF discussion #49 on the model page).
• Fewer GPUs / less VRAM: llama.cpp with the GGUF conversion works and supports RPC across machines (~22 tok/s in our earlier setup), or use lk_moe's CPU-offload (LVLLM_GPU_RESIDENT_MOE_LAYERS with fewer layers) at a large decode-speed cost.

Validated 2026-08-08: DeepSeek-V4-Flash-0731 · Lvllmds4-x v2.3.9 ·
10× RTX 3090 · TP2×PP5 · 1M context · 37 tok/s decode.

Special thanks to guqiong96 (https://github.com/guqiong96/Lvllm), yhfgyyf (https://github.com/yhfgyyf/vllm-deepseek-v4-sm89) and nudt-eddie ((https://huggingface.co/deepseek-ai/DeepSeek-V4-Flash-0731/discussions/49)

Regarding Hy3 and not only, have investigated how to improve performance, and pushed to branch https://github.com/avtc/vllm/tree/fix/load-quantized-hy3-mimo-v2.5-kvarn-0.26 two small changes:

export VLLM_KV_OFFLOAD_DISABLE_PIN=1
export VLLM_CUSTOM_AR_MAX_SIZE_MB=20
...
    --kv-offloading-backend native \
    --kv-offloading-size 36

VLLM_KV_OFFLOAD_DISABLE_PIN=1 - unblocks usage of kv-offloading (LRU KV cache moved to RAM instead of drop)
kv-offloading-size 36 - this is in GB of RAM
VLLM_CUSTOM_AR_MAX_SIZE_MB=20 - this removes fallback from custom all reduce to nccl for prefill chunks > 1K. In case of having p2p and using custom AR gives about 10-25-50% prefill speedup.

These flags should work for other models as well (when using this branch).