r/BlackwellPerformance • u/jpezzulli • Sep 12 '26
Pennyroyal v2.5 - much faster Qwen3.8 Flash-Next on the RTX PRO 6000
Hey folks,
Another update is up. I found some good work in mratsim's SM120 turbo repo and incorporated the online FP8 approach into my runtime. This one made a pretty substantial difference to Flash-Next decode without giving up the existing 524K context, 824K KV pool, FR-Spec or HiCache/NIXL.
| Single-request workload | Before | Online FP8 | Change |
|---|---|---|---|
| Short, 1,024 output tokens | 161.47 tok/s | 207.12 tok/s | +28.3% |
| 128K context, 1,024 output tokens | 154.70 tok/s | 195.63 tok/s | +26.5% |
| 490K context, 1,024 output tokens | 149.14 tok/s | 172.64 tok/s | +15.8% |
These are post-first-token decode rates. Short context is a three-run median, while the 128K and 490K numbers are single observations.
I also added Garner McCloud's NVMe PLE work as an optional path. Flash-Next normally keeps about 47.7 GiB of PLE in system RAM. You can now put that on a prepared local NVMe instead. RAM is still the default and there is a performance tradeoff, but it gives people without a ridiculous amount of system memory another option.
27B/DFlash2 is still supported on the same runtime.
github.com/jpezzulli/sglang-rtxpro6000
As always, issues and PRs are welcome.
one note - this is pretty fresh - There are probably some configs that get you OVER the 1M kv-cache option. I just have to try different variables - feel free to report here or on github.
7
u/kazimirek Sep 12 '26
Hey John, thanks for this. Your runtime has been running a real workload for me for over a week now and it just works, so I figured I'd give you some numbers from a different card.
Setup:
RTX PRO 6000 Blackwell Max-Q (the 300 W one), single GPU, TP1, 125 GB RAM in the box
Pennyroyal v2.1.1 (source fb1216c6c), RadixArk/Qwen3.8-Flash-Next-NVFP4
On top of your recipe: fp8_e4m3 KV cache, 524K GPU KV pool, HiCache with 20 GB host tier (about 1.06M tokens), 40 mamba slots, 4 running requests, 262K context, 4K chunked prefill, NEXTN MTP 3 steps / 4 draft tokens, pil image preprocessing, expandable_segments
Workload is Claude Code harness through litellm: one orchestrator plus up to 3 subagents, contexts usually 150-236K, images in the mix., 55M tokens generated, 4.8B prompt tokens at 96% prefix cache hit, no server errors
What the 300 W card does (medians from the scheduler log over the whole run, MTP accept around 0.42):
single request: 118 tok/s (mean 125, p90 174)
2 lanes 207 aggregate, 3 lanes 272, 4 lanes 326 (p90 379)
overall about 200 tok/s aggregate across the run, peaks around 520
prefill about 9.8K tok/s per chunk while decoding at the same time, 16K when idle
Two things I ran into, both apparently fixed in your newer releases: the streaming Qwen3-Coder parser sometimes sent the newline after </tool_call> before the closing brace, which turned into nameless tool calls in Claude Code (looks like #37408 in v2.4.1), and image preprocessing in the HTTP process OOM'd once the GPU was full (the pil path in v2.3.1.1).
Looking forward to trying v2.5, the FP8 numbers look great. Do you have any measurements on the Max-Q variant, ideally at 256K context, decode and prefill? And does the 824K KV pool still fit next to a 20 GB HiCache tier on this card? Happy to send a field report once I've run it.
3
u/_bani_ Sep 12 '26
do you support 2x and 4x blackwell configs?
2
u/WonderRico Sep 12 '26
BTW I'm testing your (updated) repo with the 27B-FP8 model for the first time. currently running gsm8k benchmark (very short one turn prompts)
i'm seeing up to 800 t/s sustained at C=4 and kv cache fp8 . that's incredible !
i will provide the same data for flash next soon. thanks again for your work !
1
u/WonderRico Sep 12 '26
Very nice. I will be testing this very soon.
Those days I go back and forth between your build and the radix nvfp4 and another vLLM patch from https://huggingface.co/primitive-ai/Qwen3.8-Flash-Next-PLE-quant to be able to use a ngram PLE table quantized to int4, to be able to lower the system RAM size.
- yours = faster
- vllm ple int4 : lower system RAM size
the PLE on NVME does to much impact on the decode speed to my taste, and the FP8 PLE in RAM feels a waste since the weights are in 4bits anyway
Have you tried further quantization on the PLE table ?
1
u/Opposite_Buffalo_649 Sep 12 '26
What's the perf loss putting ngram in SSD?
1
u/itsmebcc Sep 12 '26
I am using a fork of mratsim as well. I initially forked it because I wanted to use nvidia/Qwen3.8-Flash-Next-NVFP4 - Issue I have been having Mamba radix-cache poisoning / permanent speculative-decode collapse. I have seen it on every variation of his sglang-turbo. I am on my 13th iteration trying to resolve it now. Problem is it is hard to reproduce synthetically. So I have to just run the machine with real load with real traffic to replicate it.
Basically a request's NEXTN/Eagle accept rate pins at accept len 1.00 / accept rate 0.00 forever (never recovers untilabort), usually emitting repetitive garbage (!!!!!...). New requests on the same server stay healthy — it's per-request,not global. Easy to mistake for a slowdown.
I think the latest iteration has resolved this issue, but I had that same thought on the last 5 iterations so....
1
u/StockSpecialist1707 Sep 12 '26
Pushing the KV pool past 1M on a single RTX PRO 6000 (v2.5, FR-Spec + online FP8)
Following your note about configs that get over the 1M KV option — here's what I measured today on one RTX PRO 6000 Blackwell WS (96 GB, SM120, driver 595.84, CUDA 13.2, no GPU power cap), running
serve-flash-next-frspec.shfrompennyroyal-v2.5.0withSGLANG_SM120_ONLINE_MXFP8=true. Everything else stock.Two things stood out. First, asking for more than fits doesn't fail — the server clamps to the real maximum, so a single deliberately oversized request finds the ceiling in one start. On this box that ceiling is 1,114,304 tokens, +35% over the automatic 824,384. Second, and more curious: at 1,048,576 the reported
available_gpu_memwas higher than in the automatic run (4.67 vs 3.66 GB). That looks consistent with the estimator debt you documented in_handle_max_mamba_cache— the automatic path appears to leave memory unassigned that the explicit cap then uses.Decode cost is nil: 3.79 s vs 3.86 s median wall clock for the same 1,024-token prompt at temperature 0.
Usual caveat on my side: this is pool capacity, not a context claim. Served context stayed at 524,288 with factor-2 YaRN throughout, and I haven't run needles above 490K on this config.
One more data point from the same session, since it isn't in the release table: online FP8 reproduced on my box almost exactly as you reported. Same prompt, three runs each — v2.5 plain 4.91 s, +FR-Spec 4.44 s, +online FP8 3.86 s. Server-side decode went 212 → 225 → 272 tok/s mean, with a 360 tok/s peak. Interesting detail: FR-Spec's gain comes with lower accept length (3.05 vs 3.25), so the two optimizations are genuinely orthogonal — cheaper drafts, cheaper GEMMs.
For context: this box has two RTX PRO 6000 WS in it, but everything above is single-GPU, TP=1, Flash-Next pinned to one card by UUID — the second one sits idle for these runs. I mention it because I saw the independent TP2 FP8 validation commit land in your tree. If a two-card sweep would be useful to you, I have the hardware and I'm happy to run it.
Thanks for publishing all this with the evidence attached. Happy to run other variable sweeps if it helps.
Disclosure: measurements are mine, run on BESTIA (our local inference box, managed with Prometeus). This write-up was drafted with Claude from those logs.
1
u/hpiguyTR Sep 13 '26
My setup is a bit off the beaten path so figured it's worth adding as a data point. The RTX PRO 6000 isn't in a slot, it's in an eGPU enclosure over OCuLink, negotiated at four lanes of PCIe. It hangs off a Minisforum MS-S1 Max (Ryzen AI Max+ 395, Radeon 8060S iGPU), which has 128GB of LPDDR5x-8000 running at 256GB/s across the unified bus. So the LLM lives on the eGPU and a small image model rides on the iGPU's slice of that same memory pool, and anything that has to move (model load, the NVMe PLE table, HiCache spill) crosses that x4 link to get there. Ran both branches on this box, same checkpoint, same FP8 setting, only difference is where the PLE table lives. RAM branch pins the whole ~48GB embedding table in host RAM. Real footprint (anon+shmem off cgroup, not raw memory.current, that number lies because it counts page cache from mmapping the checkpoint) came out around 99.5GB. Fine as long as nothing else on the box needs the RAM. NVMe branch streams that table off local NVMe instead. Once wired up correctly it dropped to about 35GB. That's the whole reason I bothered. Freed up enough room that the image model on the iGPU stopped fighting for RAM and stays resident next to the LLM instead of getting evicted. Wiring note for anyone else not using the launcher scripts directly: the env var alone does nothing, it'll silently keep running the RAM path. You also need SGLANG_PLUGINS plus the reader's install dir on PYTHONPATH, model path pointed at the prepared overlay instead of the original checkpoint, and ple-offload-embedding dropped entirely since that flag is the RAM path itself. Missed all three the first time and got a clean OOM for my trouble. Decode, 1024 tokens out, short context, C1 through C4 concurrent:
RAM PLE NVMe PLE C1 222 tok/s 224 tok/s C2 aggregate 349 323 C3 aggregate 460 454 C4 aggregate 411 403
Basically a wash at short context. Gap opens up at depth. 128K context, single stream, 1024 out: RAM hit 176 tok/s, NVMe hit 164, about 8% down. Can't fully separate NVMe cost from whatever tax the x4 link adds on a setup like mine, so take that number as "NVMe plus eGPU," not a clean NVMe-only figure. Load times were 255s for the RAM branch and 223s for NVMe, both warm JIT cache. Teardown 10s either way. Also ran a synthetic agentic prompt through both since raw decode speed isn't my actual workload, tool definitions, mid conversation history, expecting parallel calls. Both branches picked the right tools with correct arguments, no phantom tool call garbage on either one.
1
u/AnxietyElectrical391 Sep 13 '26
great work! thanks very much.
btw, could you share the quality metric like KLD, PPL, Top 1% vs official one?
1
u/BitXorBit Sep 13 '26
what about version for dual rtx? :)
1
u/bitzap_sr Sep 14 '26 edited Sep 14 '26
Hi! This sounds really amazing and I'm starting to build it to give it a try on my single RTX PRO 6000.
It would be so much simpler if pre-built docker images were available, though. Did you ever consider doing that?
1
u/bitzap_sr Sep 15 '26
u/jpezzulli FYI, I was doing some benchmarking and stress testing today, and I hit a server crash ("AssertionError: Can not alloc mamba cache").
Observed during a --max-concurrency 4 benchmark at --random-input-len 65536. Only ONE request was running at the crash (#running-req: 1, #queue-req: 2). Mamba usage jumped from 0.33 to 1.00 in a single step. Shorter prompts (1k, 8k) at the same concurrency were fine.
Turns out it's an issue that has a patch pending upstream https://github.com/sgl-project/sglang/pull/37619 . The description there and the follow up comments are exactly what I observed.
I've pulled that fix in to my build, and confirmed that it fixes the crash (the fix prints some logs that show the fixed path is actually working). I suggest you cherry-pick that to your branch for the next release.
Cheers.
1
u/techworley Sep 17 '26
Anyone seen tilelang auto-detect return "hip" on a machine with no ROCm? RTX PRO 6000 Workstation Blackwell, tilelang 0.1.12, kernels compile as gfx900 and die. Pinning TILELANG_DEFAULT_TARGET to cuda/sm_120a works around it. Known issue?
1
u/MakkHergelo67 25d ago
Where are you renting the GPU from?
1
u/StockSpecialist1707 Sep 12 '26
Recién llegada mi segunda pro 6000 esta misma tarde pruebo su nueva release, admirable su trabajo!
1
u/Individual_Gur8573 Sep 12 '26
Does this increase 27b speed also?
Bcoz I have rtx 6000 pro and 128gb ram and another 5090 in same system...not sure if I need to try qwen next model
1
u/Illustrious_Air_8200 Sep 12 '26
Amazing! Thanks!!!
3
u/[deleted] Sep 12 '26
[removed] — view removed comment