[Bug]: Qwen 3.8 Flash Next is too slow

[MichaelFomenko]

Commit or version

git clone --branch dev --single-branch https://github.com/JustVugg/colibri.git

Environment

Debian 13, AMD 7950X, 64GB DDR5, 2TB Gen 4 Samsung Pro 990.

Reproduction steps

COLI_FORCE_CPU=1 Q38_MAXT=190000 Q38_NATIVE_FP8=1 Q38_NATIVE_BF16=1 Q38_PREFILL_BATCH=1 PROF=1 CTX=190000 URING=1 PILOT=1 SPEC_PIN=1 TOPP=1.0 COLI_TOOL_SALVAGE=1 COLI_DEBUG=1 COLI_TIMERS=1 COLI_API_KEY=local-secret ./coli web --model /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f/LLM/Qwen3.8-Flash-Next-FP8 --ram 43 --temp 1.0 --topp 1.0 --host 127.0.0.1 --port 8000 --model-id Qwen3.8-Flash-Next-FP8

Expected behavior

I watched this Video here: https://www.youtube.com/watch?v=_TCRH725hAM and was wondering why he is getting 20 tok/s with llama.cpp by loading only the experts to the ram but the integrated lookup table is not loaded to the ram, it is acceded directly from ssd.

Actual behavior and logs

Qwen 3.8 Flash Next have 0.3 tok/s and Kimi K3 have 0.5 tok/s how is this possible? Qwen have only 6B Active parameter and it barely using the SSD to get there Experts. Using only 200 to 500 MB/s instead of the maximal available 3000 MB/s.

[JustVugg]

Thanks for the detailed report, and for the exact command line. That made this quick to pin down, and the numbers tie together in a way that is worth showing, because the token rate and the disk rate are not two problems here. They are the same one.

The arithmetic

One Qwen3.8 decode token routes ten of 512 experts in each of 48 layers, 4.7 MiB each. That is 2.2 GiB of expert weights per token with nothing cached, and roughly half that at a typical hit rate. At 0.3 tok/s you are moving about 1.1 GiB every 3.3 seconds, which is 330 MB/s. That is exactly the 200 to 500 MB/s you are seeing on the device.

So the engine is not failing to use the disk. It is using all of the disk it is being given, and the question is why the disk is delivering a fraction of what it can.

For reference, docs/qwen38.md records a measured campaign on comparable hardware (i9-14900K, 61 GiB, Samsung 990 EVO on ext4, cap 32): 1.09 to 1.18 tok/s, with iobench at 2.26 GB/s cold, and 96 of the 140 second median request spent in synchronous expert disk service. Your 0.3 is low against that 1.1, and the ratio is close to the ratio of the two disk rates.

Three things I found, in order of how much they are costing you

1. Q38_MAXT=190000 is reserving 9.78 GiB before the server is even ready.

Context state is 54 KiB per token, and it is allocated for the whole Q38_MAXT ceiling up front, not as the prompt grows. At 190,000 tokens that is 9.78 GiB. At the 8,192 default it is 432 MiB. Unless you are actually sending 190k-token prompts, that is 9.4 GiB taken from the expert cache, which is the only thing standing between you and the disk. This is the single biggest lever you have right now and it costs nothing to try.

Start from what the planner says for your box:

coli plan --model <dir> --ram 43 --ctx 8192 --gpu none

2. Most of the environment variables on that command line are not read by this engine.

URING, PILOT, SPEC_PIN, TOPP and PROF are knobs of colibri.c or of other engines. CTX is never read by getenv in qwen38.c, the ceiling comes from Q38_MAXT, which you did set. And COLI_FORCE_CPU does not appear anywhere in the tree, so it is doing nothing at all.

The ones that are real on this engine and worth keeping: Q38_MAXT, Q38_NATIVE_FP8, Q38_NATIVE_BF16, Q38_PREFILL_BATCH, COLI_TIMERS. COLI_TIMERS=1 in particular will print the disk service time directly, which turns the next question into a measurement instead of a guess.

3. URING=1 doing nothing is not only a typo on your part. It is a real gap on ours.

colibri.c has an io_uring path. qwen38.c has none: expert reads go through st_read_range_raw_cap, which is a synchronous pread, one request at a time. That is what "synchronous expert disk service" in the doc above means. An NVMe drive at queue depth one, on reads this size, gives a few hundred MB/s and needs several requests in flight to reach its rated figure. We measured roughly an eight times scaling from queue depth three to eight on a different engine's expert reads, so this is not a small item, and I am tracking it separately from your configuration.

The measurement that decides where the remaining gap is

Your model sits under /media/user/6f1e9b86-.../LLM/, which is the udisks automount path for removable media. If that drive is in a USB enclosure, or on a filesystem going through FUSE, 200 to 500 MB/s is what you get and no engine change would move it.

One minute separates the two cases:

make -C c iobench
./c/iobench /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f/LLM/Qwen3.8-Flash-Next-FP8

If that reports around 2 GB/s, the bottleneck is on our side and point 3 is your answer. If it reports around 400 MB/s, it is the mount, and the fix is moving the checkpoint to the internal drive on a native filesystem. Please post what it says along with findmnt -T <that path>.

On the comparison in the video

The behaviour described there, keeping the lookup table on the SSD rather than in RAM, is already what this engine does for Qwen3.8: the 51B-parameter PLE table is never materialized, it is sixteen 160-byte row reads per token. What is not comparable is the rest. This checkpoint is 185.5 GB with 120.8 GB of routed experts, and every decode token pulls 2.2 GiB of them. Until those reads are queued rather than serialized, and until enough of them are cached, the ceiling is set by the disk, and it is measured in single-digit tokens per second rather than twenty.

Kimi K3 at 0.5 tok/s is the same story with a larger checkpoint, so I would expect both to move together once the disk question is settled.

[JustVugg]

Correction to point 3 of my previous comment, and it matters because it points at the wrong thing.

I wrote that qwen38 reads experts one request at a time. That is wrong. The engine already issues them concurrently:

• Q38_EXPERT_PARALLEL_READS defaults to 1. The misses for a layer are collected into a batch and read under #pragma omp parallel for num_threads(workers), with one job per missing expert.
• Q38_EXPERT_PREFETCH defaults to 1 and issues posix_fadvise(WILLNEED) over the ranges ahead of the read.

A decode token routes ten experts per layer, so the depth in flight is up to ten, not one. URING=1 still does nothing on this engine, but it is not the gap I claimed it was: the queueing is done with threads instead of io_uring.

That makes your mount the prime suspect rather than a secondary one. Ten concurrent reads of a few MiB each should not land at 330 MB/s on a 990 Pro on a native filesystem, and /media/user/<uuid> is the udisks automount path for removable media.

Also, my iobench invocation was wrong: it takes a file, not a directory, and its defaults (8 threads, 19 MB blocks) do not mirror this engine. Use one of the checkpoint's own shards, with a block size and thread count close to what the engine does:

make -C c iobench
ls -S /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f/LLM/Qwen3.8-Flash-Next-FP8/*.safetensors | head -1
./c/iobench <that file> 4.7 64 10 1

That is 64 random reads of 4.7 MB across 10 threads with O_DIRECT, which is close to one layer's worth of expert misses. Alongside it please post:

findmnt -T /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f
lsblk -o NAME,TRAN,ROTA,MODEL -d

TRAN is the one I care about most: nvme or usb answers this on its own.

The Q38_MAXT=190000 point stands unchanged, and it is still worth fixing first: 9.78 GiB reserved before READY against 432 MiB at the default is the cheapest 9 GiB of expert cache you will ever get back.

[MichaelFomenko]

./c/iobench /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f/LLM/Qwen3.8-Flash-Next-FP8/model-00130-of-00131.safetensors 4.7 64 10 1
O_DIRECT x10 threads: 64 reads x 4.703 MB = 0.3 GB in 0.06s -> 5.03 GB/s (1.0 effective ms/block)
O_DIRECT x10 threads: 64 reads x 4.703 MB = 0.3 GB in 0.06s -> 5.66 GB/s (0.9 effective ms/block)
O_DIRECT x10 threads: 64 reads x 4.703 MB = 0.3 GB in 0.06s -> 4.97 GB/s (1.0 effective ms/block)
O_DIRECT x10 threads: 64 reads x 4.703 MB = 0.3 GB in 0.05s -> 5.79 GB/s (0.9 effective ms/block)
O_DIRECT x10 threads: 64 reads x 4.703 MB = 0.3 GB in 0.05s -> 5.97 GB/s (0.8 effective ms/block)


findmnt -T /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f

TARGET                      SOURCE   FSTYPE OPTIONS
/media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f /dev/nvme0n1p1 ext4   rw,nosuid,nodev,relatime,errors=remo

lsblk -o NAME,TRAN,ROTA,MODEL -d

NAME    TRAN   ROTA MODEL
nvme0n1 nvme      0 Samsung SSD 990 PRO with Heatsink 2TB
nvme1n1 nvme      0 KINGSTON SA2000M8500G

lsblk -a
NAME        MAJ:MIN RM   SIZE RO TYPE MOUNTPOINTS
nvme0n1     259:0    0   1,8T  0 disk 
└─nvme0n1p1 259:1    0   1,8T  0 part /media/user/6f1e9b86-52a9-4ef6-a5fd-8965ef1db31f
nvme1n1     259:2    0 465,8G  0 disk 
├─nvme1n1p1 259:3    0   512M  0 part /boot/efi
└─nvme1n1p2 259:4    0 464,3G  0 part /

[JustVugg]

That settles two of the three candidates, and it moves the blame to us.

The mount is fine. TRAN=nvme, ROTA=0, ext4 with ordinary options. The /media/user/<uuid> path is just the udisks automount point of an internal drive, not an enclosure. I was wrong to make that the prime suspect.

The device is fine too, and then some. 5.0 to 6.0 GB/s on exactly the pattern the engine issues: 64 random reads of 4.7 MB across 10 threads with O_DIRECT, 0.8 to 1.0 effective ms per block.

So the comparison with the recorded campaign inverts:

	docs/qwen38.md reference	your box
iobench	2.26 GB/s	5.0 to 6.0 GB/s
result	1.09 to 1.18 tok/s	0.3 tok/s

A disk 2.5 times faster producing a result 3.6 times slower. The closing line of my first comment, that the ceiling here is the disk, does not hold on your machine. The disk offers 5 to 6 GB/s and the engine is taking 330 MB/s out of it, which is 15 to 18 times below what the drive just demonstrated.

Where I think it is going

q38_expert_get_batch in c/qwen38_core.h is all or nothing. The concurrent read path I described in my second comment is real, but it is refused outright, falling back to serial reads, if any one of these holds:

if(!m->expert_parallel_reads || count<2 || count>Q38_MAX_TOPK)     return 0;
if(cache->cap < count || !q38_prepare_expert_scale_bank(m,layer))  return 0;
for(previous...) if(experts[previous]==expert)                     return 0;
if(!q38_native_fp8_expert_tensors(m,layer,expert,weight))          return 0;

Serial means queue depth one, and queue depth one on reads this size is a few hundred MB/s. That is the signature you are seeing, against a drive that just did 5 to 6 GB/s at depth ten.

This also changes my point 1, and not in your favour. Q38_MAXT=190000 is not merely 9.78 GiB less expert cache, which would degrade smoothly. cache->cap is the slot count per layer, and if the remaining budget puts it below topk, then cache->cap < count is true on every layer and the parallel path is switched off entirely. That is a cliff, not a slope, and it would explain a better drive landing below the reference.

The number that decides it is already on your screen

You do not need to run anything new. You have COLI_TIMERS=1 on your command line, so the engine already printed this at shutdown:

[qwen38 expert I/O] weight-ranges=... scale-ranges=... coalesced-gate-up=...
                    prefetched=... parallel-batches=... resident-scales=... MiB

Please post that line, plus the cap the launcher started the engine with (the startup banner shows it). If parallel-batches=0, every expert read was serial and the cause is settled. If it is large, the batch is running and the problem is inside the read itself, which is a different investigation.

If you want to try something while that is in flight, drop Q38_MAXT to the default and keep the rest identical. That is the one change that could move cap back above topk.

On our side regardless of the answer

Six separate conditions can silently turn off the concurrent reads and none of them says a word. That is on us, and it is why this thread took three days to reach a log line that was already being printed. Whatever your numbers show, that path is getting a one-time warning that names the reason it refused.

[JustVugg]

1.12.1 changes this. The qwen38 trunk is int8 with integer dot products and the FP8 expert kernel is vectorized (#1668): on our 8-thread server decode went from 0.61 to 1.42 tok/s and a 512-token prefill from 495 to 149 s, perplexity within 0.5%. The prefill knobs are #1686. On a CUDA card the VRAM expert tier and the int8 trunk in VRAM apply since 1.12.0. Closing; if 1.12.1 is still slow on your box, open a new report with the COLI_TIMERS=1 output.