sm_75 (2x RTX 2080 Ti): GDN 'hc0 norm' on GPU0 takes 135-213 ms/window (GPU1: 0.05 ms) — decode stuck at ~10 tok/s

Summary

On a 2× RTX 2080 Ti (sm_75) layer split, decode of Qwen3.8-Flash-Next IQ3_XXS tops out at ~8–10 tok/s (expected ~40+ from the 12 GB-card reference numbers scaled to a 99% VRAM expert-cache hit rate). STRATA_DECODE_TIMING=1 + STRATA_VERIFY_PROFILE=1 show one line item eating 92% of the verify stage: hc0 norm in GPU0's GDN layers takes 135–213 ms/window, while the same stage in GPU1's GDN layers takes 0.05–0.22 ms. Everything else in both stages sums to ~35 ms/window.

Environment

GPU: 2× NVIDIA RTX 2080 Ti 22 GB (sm_75), driver 615.71.09, PCIe 3.0 x16 (probe: 9.1–9.6 GB/s host→device), NVLink bridge present
CPU: Xeon E5-2650L v2 (Ivy Bridge: AVX, no AVX2 → experimental STRATA_ISA_FLOOR=avx build)
RAM: 157 GB DDR3-1600
OS: Arch Linux, kernel 7.2.6-arch2-1
Engine: strata 0.1.43, built from source at 7695859 (2026-10-11) with CUDA 13.4, gcc-13, -DCMAKE_CUDA_ARCHITECTURES=75 -DSTRATA_ISA_FLOOR=avx
Model: Qwen3.8-Flash-Next IQ3_XXS (native pack), MTP on, max-context 131072, KV int8 with 32768 resident cells, --layer-split auto → K=23 (GPU0 layers [0,23), GPU1 [23,48) + head + draft)

Measurements

Decode ~7.0–10.2 tok/s across story/code/Chinese prompts (drafts accepted 55–70%). Expert cache hit rate 94–99.6% (20491/24576 profiled pairs in VRAM), misses served over PCIe — so this is not expert-miss or CPU-bound. Prefill is fine: 5,065-token prompt at 640 tok/s.

Default build (hc staged path), 300-token generation:

strata decode timing: 165 windows, avg T 2.41, 1.82 tokens/window, 258.78 ms/window = verify 255.29 (GPU-reach wait 0.00 + per-layer host 0.00 [plan 0.27 actq 0.12 jobs 0.04 CPU 3.51] + stage 221.91) + commit/emit 0.69 + draft 2.20; per layer-window: CPU experts 0.20 (0.26 entries), VRAM hits 12.07, PCIe 0.01
strata decode GPU stages (ms/window): GDN layers: q8+qkv/q-idx gemv 1.03 conv 0.08 ab 0.13 z 0.70 rec 0.25 out-proj 0.68 router+plan+bell 0.33 shared+quant 0.61 waitA 0.02 VRAM hits 2.29 copy+combine 0.07 (gap) 0.02   hc0 norm 213.39   hc0 down 0.41   hc0 up 0.45   hc1 norm 0.10   hc1 down 0.40   hc1 up 0.45 | QSA layers: ...   hc0 norm 0.03 ... | total 223.92 ms/window over 177 windows
strata decode GPU stages, stage 1 (ms/window): GDN layers: ...   hc0 norm 0.13   hc0 down 0.49   hc0 up 0.53 ... | total 22.07 ms/window over 177 windows

Same build restarted with STRATA_HC_SPLIT=0 (the plain read):

strata decode timing: 163 windows, avg T 2.40, 1.84 tokens/window, 180.41 ms/window = verify 177.09 (GPU-reach wait 0.00 + per-layer host 0.00 [plan 0.22 actq 0.11 jobs 0.03 CPU 3.30] + stage 144.32) + commit/emit 0.62 + draft 2.13; ...
strata decode GPU stages (ms/window): GDN layers: ...   hc0 norm 135.57   hc0 down 0.56   hc0 up 0.45 ... | total 146.76 ms/window over 175 windows
strata decode GPU stages, stage 1 (ms/window): GDN layers: ...   hc0 norm 0.22 ... | total 21.15 ms/window over 175 windows

STRATA_HC_SPLIT=0 helps (258.8 → 180.4 ms/window, 7.0 → 10.2 tok/s) but hc0 norm on stage 0 remains the wall at 135.6 ms.

Observations

The asymmetry is extreme: identical op, ~0.1 ms on stage 1's GDN layers (GPU1), 135–213 ms on stage 0's (GPU0). Same GPU model, same driver.
GPU0's other GDN stages look normal (q8+qkv gemv 1.03, z 0.70, out-proj 0.68, VRAM hits 2.29 ms), so this doesn't look like a general "GPU0 is slow" problem.
GPU0 utilization reads 90–96% during decode at only ~105–125 W (of the 300 W cap) at full SM clock — lots of small work or spinning, not a compute-saturated kernel.
--spec 6 instead of 4 does not help (window time, not window count, is the bottleneck).
Possibly the hc0 norm timing slot on stage 0 includes a cross-device wait (stage boundary / doorbell?) rather than pure kernel time — but then 135 ms of per-window wait on a 2-stage split would still be the anomaly.

Questions

Is this a known sm_75 (or 2-stage layer split) issue? Any knob beyond STRATA_HC_SPLIT I should try?
If it helps, I can rerun with extra instrumentation (e.g. STRATA_DECODE_TIMING, layer dumps, a forced --layer-split 26, or single-GPU) — happy to provide numbers or test a patch.

Thanks for the engine — first time this box has ever run a 125 B model at all, and the answers look right. Even at 10 tok/s it's usable; fixing this would put it in daily-driver territory.

Follow-up with a round of eliminations, plus what looks like the key discrepancy against other Turing reporters.

Same-hardware references I found after filing: #1349 (2× RTX 2080 Ti 22 GB, IQ3_XXS, v0.1.40.1) reports 74-80 tok/s decode (89 with their three-switch stack), and #2197 (single 2080 Ti, v0.1.43 7695859 — same commit as mine) reports 62.8 tok/s baseline. I get ~10 t/s on both single-GPU and split configs, so something on my box is off by ~7x, independent of the split.

Eliminated on my box (each measured with the same 300-token story/code prompts):

change	decode
baseline (v0.1.43, CUDA 13.4 build, split K=23)	7.0-10.2 t/s
rebuilt with CUDA 12.8 instead of 13.4	9.6-9.9 t/s (no change)
single GPU (`gpu: [0]`, no layer split)	9.3-11.3 t/s (no change)
split + `--pipeline-windows 2 --layer-split 24` + `STRATA_SPLIT_OWN=0 STRATA_STAGE_TRIM=1 STRATA_PREFILL_HELP=1 STRATA_BF16_TC=1` (the #1349 stack)	8.3-11.6 t/s (no change)
CPU governor schedutil → performance (1200→2100 MHz)	9.1/8.1 t/s (no change)
expert cache hit rate in all runs	94-99.6% (not expert-miss bound)

What differs from the fast reporters: my CPU is a Xeon E5-2650L v2 (Ivy Bridge, 1.7-2.1 GHz, AVX1 only, no AVX2) so I run the experimental STRATA_ISA_FLOOR=avx build; #1349 has a Ryzen 5700X (AVX2), #2197 an i7-6700 (AVX2). Also DDR3-1600 here vs DDR4 there. During decode the engine burns 2-6 CPU cores continuously, so there is real per-window host work; on this CPU it may simply be the critical path.

Question: does the ISA-floor build change anything in the decode host path besides ggml-cpu's expert rows (dispatch, router lookahead, sampling)? And is there a knob to make per-window host work visible per-subsystem (the decode timing line already shows per-layer host 0.00, which doesn't match the observed CPU burn — where should I look)?

Happy to bisect or test patches. If the answer is "the 2012 CPU is just the wall", that's a fair answer too — I'd then document it as the floor for this hardware class.

Built the #2197 PR stack locally (all nine PRs merged cleanly onto 7695859: #2069 #2070 #2072 #2093 #2071 #2091 #2092 #2146 #2145), rebuilt with CUDA 12.8 / STRATA_ISA_FLOOR=avx, enabled the full switch set (STRATA_BF16_TC=1 STRATA_SELECT_SIMT=2 STRATA_PA_V2T=1 STRATA_PF_FUSED=1 STRATA_PF_FUSED_SM75=1, plus the #1349 stack already in place).

Results on my box (same prompts as before):

	before PRs	with PR stack
decode (300-512 tok story/code/中文)	8.3-11.6 t/s	9.4-12.2 t/s
prefill, 5065-token prompt	640.8 t/s	708.2 t/s (+10.5%)

So the PRs work and help (+10-20%), but the ~7x baseline gap vs #1349/#2197's machines remains. The hc0 norm anomaly is untouched by them, which fits: #2070 optimizes a kernel that costs those machines ~65 µs/read, while mine shows 135 ms/window in that slot.

This keeps pointing at something in my environment (the AVX-floor build or the Ivy Bridge host path) rather than the sm_75 kernels themselves. Any pointer on how to make the host side of a verify window visible (or a known slow path in the ISA-floor build) would be appreciated — happy to instrument and report back.

Found it. gdb sampling of the engine during a decode shows the main thread spending its time in:

#0  strata::kernels::iq4nl_dequant_row(unsigned char const*, float*)
#1  strata::kernels::PleTable::read_row(unsigned int, float*)
#2  strata::kernels::PleTable::gather_batch(...)

The PLE table's rows are read from the GGUF shard per window, and this box's storage is the slow part: startup measured --ple-io direct at 222 rows/s (~1 MB/s of random 4 KiB reads) — a mergerfs pool on spinning disks — and the mmap fallback still page-faults to it. 16 rows × T tokens per window × ~4.5 ms per faulted read ≈ the 135-213 ms "hc0 norm" slot (it's the first stamped node after the window's gather — the GPU waits there for the host).

That also explains why the fast reporters (#1349, #2197) never saw it: NVMe random reads at ~50-80 µs make the same gather ~5-8 ms/window, invisible.

Fix on my side: --ple-io ram (loading now; the box has 157 GB RAM so the 28.8 GB table fits). I'll report the before/after numbers when it's up. Two suggestions for the engine: (1) the startup probe could warn louder than a single log line when ple-io is disk-bound — on this box it costs ~85% of decode speed; (2) the "hc0 norm" profile slot absorbing the wait made this look like a GPU kernel problem for a long time — a stamp before the gather would separate "host wait" from "hc read".

Confirmed fixed. With --ple-io ram (the whole 28.8 GB PLE table in RAM) on the same box, same build, same config otherwise:

	before (PLE from slow disk)	after (`--ple-io ram`)
decode, fresh prompts (greedy EN)	9.3-11.6 t/s	80.6-86.8 t/s
decode, fresh prompts (t=0.6/0.7, EN+中文)	8.1-9.3 t/s	71.9-76.6 t/s
prefill, 5065-token prompt	640.8 t/s	1070.3 t/s

That's ~8x on decode and +67% on prefill, now matching (slightly beating) the other Turing reports. The GPU was never the problem — the per-window PLE gather from a slow disk was. Thanks for the engine; and for the record the layer-split, CUDA 13.4, and the ISA-floor build were all innocent. 🙂

(Feel free to close; the two engine-side suggestions from my previous comment stand — a louder warning when the ple-io probe measures disk-class latency, and a stamp before the gather so the wait doesn't land in "hc0 norm".)

@ifdog

Thanks @ifdog, this is a very clear report. Short answer: not a known issue, and we cannot reproduce it with the cards we have, so here is what we see and what would pin it down.

What the numbers say:

On our two-card box (2x RTX A4000, sm_86, layer split auto, the same IQ3_XXS pack, STRATA_DECODE_TIMING=1 STRATA_VERIFY_PROFILE=1) stage 0's GDN hc0 norm is 0.31 ms per window and stage 1's is 0.08 ms; the stage totals are 11.6 and 12.1 ms and decode runs at about 72 tok/s. So the stage-0 norm slot is not slow by design.
The slot is a 16 us kernel on T blocks. 135-213 ms on one card, with the other card's identical kernel at 0.05-0.2 ms and the card at 90-96% utilisation but only 105-125 W, means the kernel is waiting, not computing: either for memory it reads, or for a flag. Only the GDN layers show it on stage 0 (the QSA layers' hc0 norm is 0.03), which also fits waiting rather than a slow kernel.
STRATA_HC_SPLIT=0 moving it from 213 to 135 ms (the plain read does the same norm) says it is the norm's own first access and not the split path.

Our best guesses, in order:

Something the norm kernel reads on GPU0 sits in memory the GPU has to fault in or reach over PCIe on every window (system-allocated memory with the open kernel module's HMM, or a mapped-host table), where GPU1's copy is in VRAM. The Arch kernel 7.2 with driver 615 open modules is the newest combination we know of.
The NVLink bridge: stage 0 reads something the other card wrote (the previous window's hand-off) through a peer mapping that is degraded or falls back to a slow path.
Power or clock state of GPU0 (a low-power link state that only wakes on use).

What would tell them apart, any of these helps:

STRATA_DECODE_TIMING=1 STRATA_VERIFY_PROFILE=1 again with the cards swapped (CUDA_VISIBLE_DEVICES=1,0): does the slow slot follow the card or stay on stage 0?
Single card: the same model with --gpus 0, then --gpus 1, and a smaller --max-context, so everything fits in one stage. If a single stage on GPU0 is slow, it is the card or the driver, not the split.
--layer-split 26 instead of auto: does the stall change with the split point?
nvidia-smi topo -m, nvidia-smi nvlink -s, and dmesg | grep -i -E 'nvrm|xid' after a run.
nvidia-smi -q -d PERFORMANCE,POWER during decode for both cards (clock throttle reasons, power state), and whether nvidia-smi -pm 1 (persistence mode) changes anything.
If you can, nsys profile --trace=cuda,nvtx -d 20 on a decode: the 135 ms bar will show whether it is a kernel, a memory fault or a wait on an event.

We would rather fix the cause than add a knob, so the swap and single-card runs are the most useful first step. Thank you for the AVX-floor build on a no-AVX2 Xeon too; it is good to see it run a 125 B model.

Thanks — our messages crossed: I found the cause right after posting those, and it was none of the three candidates (so no need for the swap/single-card/nsys runs). Your read of the symptom was exactly right, though: the kernel was waiting, not computing — waiting on the host.

Root cause: the per-window PLE gather read from a spinning-disk pool. gdb sampling of the engine during decode:

#0  strata::kernels::iq4nl_dequant_row(unsigned char const*, float*)
#1  strata::kernels::PleTable::read_row(unsigned int, float*)
#2  strata::kernels::PleTable::gather_batch(...)

The startup probe on this box measures --ple-io direct at 222 rows/s (~1 MB/s of random 4 KiB reads) — a mergerfs pool on HDDs; the mmap fallback page-faults to the same disks. 16 rows × T tokens per window × ~4.5 ms ≈ the 135-213 ms I was seeing in "hc0 norm" (the first stamped node after the gather, so the GPU's wait lands there). QSA layers showed it less because their slot ordering differs; the NVLink was innocent (2×25.781 GB/s, NV2, no throttle reasons, links healthy — nvidia-smi topo -m / nvlink -s output below if you want it for the probe work).

Fixed with --ple-io ram (28.8 GB table in RAM):

	before	after
decode, fresh prompts	9.3-11.6 t/s	80.6-86.8 t/s
prefill, 5065 tok	640.8 t/s	1070.3 t/s

That matches your two-card A4000 number (72 t/s) — the GPU side here is fine.

nvidia-smi topo -m / nvlink -s:

        GPU0    GPU1    CPU Affinity    NUMA Affinity
GPU0     X      NV2     0-9,20-29       0
GPU1    NV2      X      10-19,30-39     1
NVLink 2 links per card, 25.781 GB/s each; no SW power cap, no HW slowdown

Two fix-the-cause suggestions (rather than knobs), now that the mechanism is clear:

Attribute the wait correctly. The gather happens on the host before the window; the first GPU stamp after it is hc0 norm, which is why this looked like a GPU kernel problem for hours. A stamp pair around the gather (or starting the verify timing after it) would have pointed at the PLE immediately.
Treat disk-class PLE latency as a start-level condition, not a footnote. The probe already knows (222 rows/s) and silently falls back to mmap — which is the same disks. Since the fallback choice costs ~85% of decode speed on such a box, auto-selecting ram when the probe measures disk-class latency and the table fits in RAM (it says the size already) would have saved me an afternoon; a line in the start summary ("PLE is disk-bound: decode will be ~8x slower until --ple-io ram") would do it too.

Separately, #2331 is still open: UD-Q4_K_XL never reaches "session is up" here — the engine exits silently at the same point in four configs (71 GiB budget → dies in FileExpertSource: page-locking it in steps (#1250); no budget → dies right after the GPU line; split and single-GPU alike; no OOM, no core). Happy to run whatever you want on that one — the files and pack are in place.

Cannot reproduce on 2x RTX 2080 Ti with engine 0.1.46.3: `hc0 norm` is 0.35 ms/window on GPU0

Same card class as the report, newer engine, and the stage is healthy here. I could not reproduce the 135-213 ms/window hc0 norm on GPU0.

Setup: 2x RTX 2080 Ti 22 GB (sm_75), NVLink NV2 present (2 links, 25.78 GB/s), PCIe 3.0 x8 host link (6.5 GB/s host->device), driver 595.99.02, CUDA 12.8, Ubuntu 24.04.5, Strata 9de04cd4 (v0.2.1.4), engine 0.1.46.3 compiled from source for sm_75 (-DCMAKE_CUDA_ARCHITECTURES=75), i7-10700K (AVX2, no AVX-512), 94.1 GiB RAM. Qwen3.8-Flash-Next GSQ-RCO IQ3_XXS native pack, MTP on, 262144 context, kv int8 with --kv-resident 32768, --layer-split 26, --batch 3, STRATA_SM75_PERF=1.

With STRATA_DECODE_TIMING=1 + STRATA_VERIFY_PROFILE=1:

strata decode GPU stages (ms/window): GDN layers: ... hc0 norm 0.35   hc0 down 0.46   hc0 up 0.49   hc1 norm 0.12 ...
strata decode GPU stages, stage 1 (ms/window): GDN layers: ... hc0 norm 0.09   hc0 down 0.36   hc0 up 0.40   hc1 norm 0.09 ...
total 14.54 ms/window over 87 windows

hc0 norm on GPU0 is 0.35 ms/window (stage 1: 0.09 ms) - about 400x less than the 135-213 ms reported, and the same order as the reporter's healthy GPU1 (0.05-0.22 ms). The whole GDN stage on GPU0 is ~4 ms/window and the window totals ~14.5 ms, which matches a normal decode rate: 74 tok/s single stream, and 73-84 tok/s across the bench runs.

The engine also says what it does on this card now:

strata hc: CUDA0: the hyper-connection read runs as staged (the norm per token and stream, the down projection's
activations staged ahead by cp.async); checked bit for bit against the plain read on this card

The difference between the two setups that I can see: the report's CPU is a Xeon E5-2650L v2 (no AVX2 at all, an experimental STRATA_ISA_FLOOR=avx build) with DDR3-1600 and 157 GiB, and it ran engine 0.1.43. This box has AVX2, DDR4, and 0.1.46.3.

If the report is still reproducible on 0.1.46.3 with the same STRATA_VERIFY_PROFILE=1 line, the hc0 norm timing would show it directly; if it now reads like the numbers above, the issue looks fixed or was specific to that build. I can also re-run the same measurement with an AVX-only build if that is what is needed to tell the two apart.

Thank you for the full write-up. Closing: the cause you found, PLE rows read from a slow disk (--ple-io direct measured 222 rows/s), is a setting and not an engine fault, and --ple-io ram took decode from 9-11 to 80-87 tok/s on your box. The engine prints the measured rows per second at start and warns when an explicit --ple-io direct measures slow.

Navigation Menu

Repository navigation

sm_75 (2x RTX 2080 Ti): GDN 'hc0 norm' on GPU0 takes 135-213 ms/window (GPU1: 0.05 ms) — decode stuck at ~10 tok/s #2239

Summary

Environment

Measurements

Observations

Questions

Cannot reproduce on 2x RTX 2080 Ti with engine 0.1.46.3: `hc0 norm` is 0.35 ms/window on GPU0

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

sm_75 (2x RTX 2080 Ti): GDN 'hc0 norm' on GPU0 takes 135-213 ms/window (GPU1: 0.05 ms) — decode stuck at ~10 tok/s #2239

Description

Summary

Environment

Measurements

Observations

Questions

Activity

ifdog commented on Oct 13, 2026

ifdog commented on Oct 13, 2026

ifdog commented on Oct 13, 2026

ifdog commented on Oct 13, 2026

Niko1221 commented on Oct 13, 2026

ifdog commented on Oct 14, 2026

yusheng227507 commented on Oct 21, 2026

Cannot reproduce on 2x RTX 2080 Ti with engine 0.1.46.3: hc0 norm is 0.35 ms/window on GPU0

Niko1221 commented on Oct 29, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions

Cannot reproduce on 2x RTX 2080 Ti with engine 0.1.46.3: `hc0 norm` is 0.35 ms/window on GPU0