Skip to content

V100 (sm_70) support for the Qwen3.8-Flash-Next MoE engine - #3

Closed
noorazman wants to merge 7 commits into
Niko1221:mainfrom
noorazman:feature/v100-moe
Closed

noorazman wants to merge 7 commits into
Niko1221:mainfrom
noorazman:feature/v100-moe

Conversation

@noorazman

@noorazman noorazman commented Sep 25, 2026 •

Copy link
Copy Markdown

No description provided.

5x Tesla V100 (2x32GB PCIe, 3x16GB SXM2), all cc 7.0, driver 580.178.04,
CUDA 12.9.86, Ubuntu 22.04.5, Xeon E5-2680 v4 x2 (AVX2, no AVX-512), 125 GB RAM.
Unmodified tree builds for its sm_120 default target (BUILD_OK); without an
explicit CMAKE_CUDA_ARCHITECTURES the configure step fails on this machine
(CMake native detection quirk). Documents quirks: nvidia-smi query segfault,
low soft memlock in non-login shells.
- CMakeLists: lower the architecture floor from 80 to 70 (the only sm_80-gated
  feature in the tree is the QSA scorer's ldmatrix + tf32 MMA).
- device.cu: accept cc >= 7 instead of cc == 12 (sm_120 stays the reference).
- native_qsa_score.cu: warp-FMA fallback for __CUDA_ARCH__ < 800 - tf32-
  truncated inputs, 5-step shfl_xor butterfly per (row, head), ReLU and the
  ordered head/bias/tail-bonus additions match the MMA path's operation order.
- qsa_score_hostref_test: standalone host-reference test for the fallback
  (15 n_kv values incl. incomplete blocks and 2048-2051; worst rel err 2.6e-6
  on the V100, cells [n, max_cells) untouched).
- Docs/v100-compatibility.md: the full sm_70 audit (feature verdicts, the
  scorer contract, empirical evidence on GPU 0).

Verified: sm_70 fatbin contains only sm_70 cubins; scorer SASS has 0
ldmatrix/0 MMA; strata-device --selftest and 16/16 kernel parity selftests
pass on the V100-PCIE-32GB.
- enable_testing() unconditionally: the parity/selftest add_test() calls are
  unconditional, but ctest only saw them when the unpublished tests/ suite
  was present.
- Docs/v100-build.md: exact configure/build commands for this machine,
  Phase 4 verification evidence, ctest selection (19/20 on a clean tree;
  ple_parity needs the Q2_0 GGUF + bench/micro fixtures by design).
The unbraced 'else' after the CPU_ISSET loop bound to the INNER if (C
dangling-else), so with a successful sched_getaffinity the fallback
hardware_concurrency() list was pushed once per UNSET mask bit: on this
56-CPU machine physical_cores(true) returned 54,263 cores and the
ExpertPool ctor spawned that many 8 MB-stack workers, stalling main for
minutes (reproduced at -O0..-O3; 14k+ threads observed live). The fix is
two braces; the Windows branch was already braced. Verified: physical_cores
returns 55 and the first V100 inference completes (28.4 tok/s decode).
- v100-testing: pool-worker startup bug (dangling else), first-inference
  results, determinism, token-level llama.cpp cross-check, GPU4 16GB validation
- v100-benchmarks: throughput on both cards, peak VRAM, pool-worker and
  hugepage A/B measurements
- v100-stage1-final: phase-by-phase PASS verdict, deliverables, caveats
@noorazman noorazman closed this Sep 25, 2026
@noorazman
noorazman deleted the feature/v100-moe branch September 25, 2026 14:33
maxious added a commit to maxious/Strata_SYCL that referenced this pull request Oct 4, 2026
…a no-win; P0b Niko1221#2/Niko1221#3 and P0 Niko1221#2's XMX avenue closed

With a persistent FP16 weight copy, oneMKL's dense FP16 GEMM loses every ncols the
engine runs (ncols=1..4: 39.9/38.0/34.4 us vs native_mmvq 21.4/23.9/29.7, i.e.
0.54-0.86x) and only crosses over at the fat MTP window (ncols=6 1.33x, ncols=8
2.39x), costing 2x dense VRAM + ~12 us materialization. native_mmvq stays the
decode path; the decode levers are P0 Niko1221#2's pipe-arithmetic slimming and the NCOLS
unroll-to-loop. decode_xmx_gemm_bench + exp 25 doc; README P0b Niko1221#2/Niko1221#3 closed.
AdrianBM96 pushed a commit to AdrianBM96/Strata3060 that referenced this pull request Oct 4, 2026
- PORT_FORK.md: metodo por port (rama, cherry-pick, build sm_86, ctest, bit a bit, A/B 6+6) y
  tabla de estado. Hallazgo: el Niko1221#1 (GR_DOWN_MAX4) YA esta en nuestro arbol (lo trae el 0.1.38);
  el Niko1221#2 (bce7fbb, MTP chain) esta entrelazado con el layer-split; el Niko1221#3 (f42c58f, AVX-VNNI) es
  portable; el Niko1221#4 (CPU assist) es el ultimo y no bit-exacto.
- ESTADO_DESPLIEGUE.md: servicios, puertos, cambios adoptados con su medicion, estado del repo
  del motor y reconstruccion rapida tras un reinicio.
- ops/env-ab.sh: arnes generico de A/B de una variable de entorno (alterna off/on/off/on).
AdrianBM96 pushed a commit to AdrianBM96/Strata3060 that referenced this pull request Oct 4, 2026
…ancia en la 3060)

- Niko1221#1 GR_DOWN_MAX4 (ya en el arbol): A/B 3+3, off 42,05 vs on 42,55 tok/s -> +1,2%, ruido. Apagado.
- Niko1221#3 f42c58f AVX-VNNI+gather: portado (f77a640), compilado, BIT-IDENTICO verificado con
  native_expert_parity (iq_parity solo cubre GPU; iq2s_avx2_test.cpp no existe). A/B 3+3:
  on 42,80 vs off 43,25 -> dentro del ruido (sesgo termico a favor del on). Revertido; rama
  port/avx-vnni guardada. Coincide con el fork: con IQ2_XS queda dentro del ruido.
- Niko1221#2 bce7fbb (MTP chain) esta entrelazado con el layer-split (set_stage_flags, xstage):
  cherry-pick suelto no vale; extraer la parte de 1 GPU o pedirselo a Claude.
maxious added a commit to maxious/Strata_SYCL that referenced this pull request Oct 4, 2026
…is measured null, parked

The i-quants (IQ4_NL/IQ4_XS/IQ3_XXS) dequant at 75-87 GB/s, ~4-6x below Q2_0's
bandwidth-bound 472. The bottleneck is the codebook *select* over the compile-time
constexpr int8 kvalues_iq4nl (a register select-tree, not a memory gather). A
register-table hoist of d*codebook (removing per-value convert+mul) is bit-identical
(dequant_bench checksum unchanged) but null (~75 GB/s); reverted. Q2_0 stays the
prompt dequant pick. README P0 Niko1221#3 dequant half marked resist-optimization (confirms
exp 04/11).
TheHubUser4860938 added a commit to TheHubUser4860938/Strata-Mnemosyne that referenced this pull request Oct 13, 2026
…the venue's own false leak, fixed by the venue

Live run Niko1221#3 exited 2 on a leaked control - and reading the answer showed why that exit
was right AND why the criterion was wrong: the model said it did not know the machine and
then listed examples of how to check, one being 'Custom PC with an NVIDIA RTX 4090'. A
substring over the whole text called that knowledge. The venue now asks with the same
few-words/idk rule the LongMemEval harness uses (consistent protocol voice across venues)
and scores only the FINAL line - few-word answers are one line, verbose ones end with the
answer. The control that stopped the venue is the control working: it refused to pass
quietly, and the fix came from reading the record rather than widening the criterion.
FlareP1 pushed a commit to FlareP1/Strata that referenced this pull request Oct 13, 2026
… warning goes

The trace was written by the one-token pool only, which a native pack never reaches (its decode is verify windows
from the first position), so the file stayed empty and the summary counted pool calls as records. The multi-token
pool now gets the router weights the window already maps (PoolMultiFn) and writes a record per row, drafts too, and
the summary counts records. Reported by pull requests Niko1221#3 (johnmosesventura06) and Niko1221#10 (guthirry).

The startup warning that the GPU hit path is NOT CORRECT predates the native kernels: the GPU and the CPU pool
compute an expert from different activation forms, so a run with the cache can part from a cache-off one at a near
tie, while the distributions agree with llama.cpp's.
svgop added a commit to svgop/Strata that referenced this pull request Oct 17, 2026
…iority order (T1.2a attention Niko1221#1, expert-fusion Niko1221#2, HC-fusion Niko1221#3, doorbell Niko1221#4, smoothness Niko1221#5); 3 research agents in flight
svgop added a commit to svgop/Strata that referenced this pull request Oct 17, 2026
…ged: the gated draft path launches step j only after the host read step j-1 (mtp.cpp) - each step pays a host launch bubble while the GPU idles; the ungated arm enqueues all step graphs back-to-back (the min_p<=0 twin's shape) and the readout still trims at the first prob under min_p, so drafts/dprob are bit-identical; price = executing the max_steps-n steps past the trim (~50us each). stair50/stair35 configs (STRATA_SPEC_DEPTH=1 counter; min-p 0.5 vs 0.35) staged for the depth-stair diagnostic
svgop added a commit to svgop/Strata that referenced this pull request Oct 17, 2026
…; only 300-class lane = verify-window megakernel restructure (Cohere 1.58x proof); FREE BOOT Niko1221#3 STRATA_DUMMY_NODES prices per-node tax; four tail levers R2-R5 priced; four more negatives banked
ProTekk added a commit to ProTekk/Strata that referenced this pull request Oct 17, 2026
…21#3)

The code-byte run is read as floor-aligned 32-bit words + one uniform funnel
shift per word instead of per-byte LDG.E.U8. A 16-byte floor window would need
a runtime word offset and dynamic register indexing would spill it to local
memory; the 66/130-B row strides keep every run at 0 or 2 mod 4, so a 4-byte
floor needs no word offset - off is 0 or 2, the shift is 0 or 16 bits.

SASS (shipped objects, turbo4 mode 5): LDG.E.U8 64 -> 0, instructions
2456 -> 1848 (now below q4_0 mode 4's 1936), REG 96 STACK 0 LOCAL 0; every
mode-5 tier spill-free (96-114 regs), all U8 loads gone.

The window overreads <=7 B past the block: the arena pools now take +16 B
tails (layer.cpp); the parity harness already mallocs +64. The elastic-VMM
corner (an array an exact multiple of the granularity) is documented at both
sites - elastic is opt-in and off in every bench/live config.

Receipts: qsa_prompt_attn_parity 12/12 PASS, all NEWHASHes bitwise identical
(pure instruction-level change); ms/chunk turbo4 15.21 -> 12.77 (-16%),
turbo8 20.87 -> 15.27 (-27%), t3 15.03 -> 13.67, t2 13.53 -> 13.11,
t1 12.72 -> 12.43; q4_0/int8/fp16 controls unchanged. turbo4's kernel gap to
q4_0: 6.15 -> 3.71 ms/chunk.
tirex999 pushed a commit to tirex999/DACAN that referenced this pull request Oct 18, 2026
The pause note is gone: the sweep continues from t = 0.9, seed 9, Niko1221#3.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
tirex999 pushed a commit to tirex999/DACAN that referenced this pull request Oct 18, 2026
…18.10 21:4x MSK)

Cell 0.9 / seed 9: Niko1221#3 after the restart matched #1 and Niko1221#2 byte for byte; the warm-up before it (first long answer after the start) differed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
lbbit pushed a commit to lbbit/Strata that referenced this pull request Oct 28, 2026
用户反馈:web Monitor 看不到并发数/资源占用,且多任务时只显示 queued(实为引擎
未开 --batch 的单并发)。两处根因:
1. dev 服务未加 --batch → 引擎单并发,第2个请求只能排队; api_monitor 默认关 → Monitor 页无数据。
2. 真正的 Monitor 页是 index.html + app.js(单页应用),此前我把 KV pool 加到了未使用的
   monitor.html/monitor.js(已回退),从未在真实页面显示。

本提交在真实 Monitor 页加(符合既有卡片/muted 文本风格,不破坏布局):
- index.html: Model state 卡加一行 #conc-detail;Context fill 卡在 Experts in VRAM 后
  加 "KV pool" bar-row + 进度条。
- app.js renderMonitor: 从 /metrics 的 live 渲染并发(N-way parallel · running · waiting —
  逐 slot 状态/速度),从 engine 渲染 KV 池占用(used/total pages + %)。

实测(dev 版 build-hip/strata,双卡,--kv-pool --batch 4,端口 12346):
- 空闲: "4-way parallel · 0 running · 0 waiting — #0 idle, Niko1221#1 idle, Niko1221#2 idle, Niko1221#3 idle"
- 2 个长请求并发: "4-way parallel · 2 running · 0 waiting — #0 decoding 8.1t/s, Niko1221#2 decoding 10.2t/s"
- KV pool: "2,060 / 2,407,603 pages (0%)"
- 两个长请求均正常完成;截图确认样式与既有行一致。

附带(运行配置,非提交):dev 服务需 --batch N 才有并发、api_monitor:true 才有 Monitor 数据。

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
RandoTeam referenced this pull request in RandoTeam/4gbVramStrata Nov 1, 2026
[Docs] 4 GB: the real start flow, what ran, and the settings that change speed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant