Repository navigation
Volta (sm_70) speed on strata-poly: measured GLM-5.3-Flash numbers + request for a Volta kernel path and --kv q8_0 for glm5next #88
Description
Feature request: Volta (sm_70) path for strata-poly — tensor-core kernels + --kv q8_0 for glm5next
Follow-up to #73 (the kda_step_kernel launch-bound fix, 79d51a1, which we hit and bisected). With that fix the engine runs on our rig; this issue is about speed on the same hardware, and a request for the two repo gaps that cap it.
Hardware (what you don't have — reporting per COMMUNITY_BENCHMARKS.md)
- 2x Tesla V100-PCIE 32GB (sm_70, ~897 GB/s, PCIe gen3, no NVLink) + 1x RTX 2080Ti 22GB (sm_75)
- Xeon E5-2673 v3 x2 (24c/48t, Haswell, AVX2, no AVX-512), 2 NUMA nodes; 251 GiB RAM; Linux; driver 570.211.01, CUDA 12.8
- Build: main as of 23 Oct (includes 79d51a1),
-D CMAKE_CUDA_ARCHITECTURES=70;75
Model and measured numbers (what we measured)
GLM-5.3-Flash-REAP50-Q3_K_M.gguf, 78.78 GB (half routed experts pruned, 6,192 experts left, own draft/MTP block). 3-card --layer-split auto, 128K context, --vram-reserve-mib 2048:
- experts resident: 85.7% in VRAM (rest through the CPU path)
- decode, no-graphs: 34-35 tok/s
- decode, CUDA graphs on: 37.8-44.9 tok/s
- draft/MTP active, acceptance low in our runs (~1 of 2 guessed tokens kept)
- 4-case bench against our production baseline: short 71% / mid 47% / tooljson 82% / reasoning 44%
Baseline for comparison, same rig, same day class: Qwen3.8-Flash-Next IQ3_XXS on the Strata project's own engine (v0.1.41 custom build, sm_70), 3-card split, experts 100% VRAM-resident: 75-97 tok/s decode. So poly-on-Volta runs at roughly half of what the same cards do for the Strata engine's expert path.
What we believe caps it (from reading the source; measurements not yet run per-request)
STRATA_POLY_MMAis compiled out below sm_75 (src/poly/kernels.cu~1646:__CUDA_ARCH__ < 750 -> 0), andtensor_cores()gates onptxVersion >= 75(~2150). Volta has tensor cores (m16n8k8 f16) but noldmatrix; the poly flash-attention and MMA paths therefore fall back on our cards. Upstream Strata's ownnative_mmvq.cualready carries a measured Volta table (kIlRows_70[], read off a V100-SXM2 bench), so the project has proven Volta tuning — it just isn't in the poly kernels.--kv q8_0is rejected for GLM (src/poly/engine.cpp~991: "only supported for qwen35moe models in v1"). At 128K fp16 KV the residency headroom is what keeps 14.3% of experts off-card; q8_0 for glm5next would directly attack the expert-residency gap we measured (85.7%).- The repo-seeded
--expert-profilecounts come from 5090 runs; we have not yet recalibrated on this rig (config work on our side, listed for honesty).
What we ask
- A Volta path for the poly kernels: enable the tensor-core step on sm_70 (m16n8k8 without
ldmatrix, or a Volta-specific variant), and/or route poly expert GEMV through the already-tunednative_mmvqVolta table. --kv q8_0(or the int8/int8k readers from poly: --kv int8 (and int8k) — a compute-native Q8_0 KV cache for qwen35moe and Gemma 4 #33/poly: --kv int8 and --kv int8k for Gemma 4, with the readers renamed apart from q8_0's #59) for glm5next.
We can run any A/B you want on real sm_70 hardware — we are likely the only Volta box reporting on this repo, and #87 shows the Pascal/Volta-adjacent paths are actively being worked. Happy to re-run the same 4-case bench against a same-day build with a candidate patch, and to post per-run JSON.
What we did not test
- No knob sweep yet (
--expert-swaps,--window, manual--layer-split) — config numbers above are the default/auto settings. - No HIP/SYCL builds on this rig.
- Vision: not tested beyond confirming mtmd cannot read GLM images (per GLM 5.3 Flash on V100 (sm_70): every generation fails - cudaGraphInstantiate / native MMVQ 'too many resources requested for launch' #73).
Activity
VecSzn commentedon Oct 25, 2026
Thanks for the numbers, neither are planned on my side right now, I don't have a V100 and GLM's KV cache is not laid out like Qwen3.6's, PRs are welcomed though, and it's your offer to A/B on the V100s that would let me take them.
Csavoldi commentedon Oct 26, 2026
Another 2x V100-PCIE-32GB box here (47 GB RAM, Xeon E5-2682 v4, no third card — our Qwen3.8 numbers and Docker recipe are in #95). Reading this issue pushed us to try the poly engine on Volta, so a small data point for the "it has not run on such a card" note in docs/OLDER_GPUS.md:
strata-poly runs on Volta. Built from main (7b70075) per the docs with CUDA 12.9.1,
-DCMAKE_CUDA_ARCHITECTURES=70 -DSTRATA_EXPERIMENTAL_SM60=ON, targetstrata-poly, and served Qwen3.6-35B-A3B UD-Q3_K_M split across both cards (context 32768, the setup default for this model).Baseline
strata_benchmedians, no fattn/fdec/gdn-mma (i.e. exactly the fallback paths this issue is about):Two things we hit on the way, both being investigated:
poly_attend_testdies withattention scores: cuBLAS status 1on the V100 — insideattend_batch'scublasGemmStridedBatchedEx(fp16 in / fp32 out,DEFAULT_TENSOR_OP). The engine's own prefill through the same path works fine (a 14k-token prompt went through no problem), so it looks shape-specific rather than a broken Volta path; we're narrowing down which size crosses the line.mma.sync.aligned.m8n8k4(in the exact 8-accumulator form qsa_prompt_attn.cu uses, with the fragment mapping documented there) andwmma 16x16x16both compile under CUDA 12.9, execute correctly, and show HMMA.884 in SASS on driver 580.173.02. The mapping comment in qsa_prompt_attn.cu reproduced all 256 outputs of our probe exactly.Next on our side: a
(256, 256, 64)row for the Volta fattn config table plus lettingmake_setupaccept Volta (the dormantVOLTA_MMA_AVAILABLEbranches in fattn_mma.cuh look reachable for Qwen3.6's 64-column tile), and after that a Volta arm for the GDN path modeled on the same m8n8k4 pattern. PRs to follow if the numbers hold up — happy to coordinate so we don't duplicate your A/B offer.llollaalloll commentedon Oct 28, 2026
@Csavoldi — this is the rig behind the original post (2x V100-PCIE-32GB + 2080Ti-22GB, driver 570.211, CUDA 12.8). Your data point changes our plan: poly-on-Volta clearly has headroom when the model is fully resident.
Our target, stated plainly: GLM-5.3-Flash-REAP50 within reach of our production baseline — Qwen3.8-Flash-Next IQ3_XXS on the Strata engine, 2-card split, experts 100% resident, int8 KV, 262K ctx = 132-134 tok/s decode on these same cards. Current poly numbers for GLM: 34-35 (no-graphs) / 37.8-44.9 (graphs), experts 85.7% resident on 3 cards.
Two questions before we duplicate effort:
(256,256,64)row +make_setupaccepting Volta + GDN arm): rough ETA? We'd rather A/B your patch than write a parallel one.VOLTA_MMA_AVAILABLEpath looks reachable "for Qwen3.6's 64-column tile". Does the same m8n8k4 pattern cover glm5next's attention/head shapes, or does GLM need its own table rows? (Our blocker list from the OP:STRATA_POLY_MMAgate at__CUDA_ARCH__<750,tensor_cores()ptxVersion>=75,--kv q8_0rejected for glm5next.)Coordination offer, concrete: we can run your candidate branch against our fixed 4-case bench (short / mid / tooljson / reasoning, temp 0, nonce per request, per-run JSON posted) on real sm_70, same-day build, both 2-card (2x V100 only) and 3-card configs — including a
--kv q8_0A/B if your branch grows one. We are also the only box here that already bisected a poly kernel launch bug on Volta (#73 /79d51a1e), so we can triage Volta-specific launch failures fast.One data request from our side: your 157.5 tok/s was Qwen3.6-35B-A3B (~23GB, fully resident on 2x32GB). Have you tried anything on the 60-80GB class where experts cannot all fit? That residency gap is our dominant cost and we want to know if your setup ever crossed it.