Skip to content

Verify window: PDL on sm_90+, graph branches, batched Q4_0 KV append and PLE post-ops (bit-identical) - #904

Closed
Hardin22 wants to merge 848 commits into
Niko1221:mainfrom
Hardin22:pr/dense-pdl
Closed

Hardin22 wants to merge 848 commits into
Niko1221:mainfrom
Hardin22:pr/dense-pdl

Conversation

@Hardin22

@Hardin22 Hardin22 commented Oct 5, 2026 •

Copy link
Copy Markdown
Contributor

This is the part of the fork's "fewer, fuller launches" work that main doesn't have yet. It comes as two commits, and each piece has its own switch. Every piece keeps the arithmetic of the path it replaces, so the output is unchanged; the checks are below.

  1. Programmatic dependent launch on sm_90+ (STRATA_DF_PDL). A decode window is a long chain of small kernels, and on a fast card many of them finish before the next one has fetched its weights. With PDL, the quantize kernel, the multi-column MMVQ and the multi-row BF16 GEMV load their weights before griddepcontrol.wait.
    • Where it applies: the attribute is only set while capturing, and only when every predecessor of the next node is a kernel (never after a memcpy, memset, host or event node). STRATA_DF_PDL=2 restricts it to single-predecessor edges.
    • When it is on: pdl_supported() is false on HIP and on CUDART < 12.3. It also needs cc_major >= 9 and a binary built for sm_90+ (PTX JIT from older code would lack the wait).
    • Other changes:
      • copy_rows_strided replaces a memcpy2D node.
      • The three PDL kernels drop __restrict__ (llama.cpp's PDL fix), which can change speed but not values.
  2. Graph branches and batched forms:
    • STRATA_DF_BRANCH: the GDN a/b/z, query and indexer-query sides and the step/pos inputs on a side branch. Opt-in since f89d31f: on Linux (open modules, GSP) it stalled with NVML queries between requests (Pipelined windows + async adaptive tier together (stacked on #859 and #876) #905).
    • STRATA_DF_QB_Q4, STRATA_DF_KVAPP: the Q4_0 KV's window batched (main still appends per token with --kv q4_0).
    • STRATA_DF_PLE: the PLE post-ops in one pass.

Already in main, so left out: native_swiglu_quantize_q8_1 / sigmoid_scale_rows, fd95405's attention change, the shared-expert branch, and STRATA_PLE_BATCH's projections. I also dropped a one-row-per-warp MMVQ variant, because it did not beat main's ROWS=2.

Measured

Swift 1.5 IQ3_XXS at 160K (q4_0 KV), RTX 4060 Ti + RTX 5080 layer split, with the resident RAM mode (#848), pipelined windows and the async tier (#859, #876). Greedy, five prompts × two rounds, decode tok/s.

run 1 run 2
every switch on (default) 123.5 119.7
every switch off (STRATA_DF_PDL=0 STRATA_DF_BRANCH=0 STRATA_DF_QB_Q4=0 STRATA_DF_KVAPP=0 STRATA_DF_PLE=0) 118.1 117.5

About +3%. PDL only acts on the 5080 (sm_120); the 4060 Ti (sm_89) has none.

Checks

  • Exact pass (STRATA_IQ_MT_MIN=1 --pcie-frac 0 --adapt-every 0, serial loop, 4 prompts × 500 tokens): all switches on and all off give the same text.
    • Against a build without these commits, 2 of 4 differ. The larger binary leaves 2 fewer expert slots on each card (5740 vs 5742, 4762 vs 4764), so a few experts move to the CPU, which rounds them differently.
  • Unit tests pass on both cards (RTX 5080, RTX 4060 Ti):
    • pdl_parity: 0 of 200 replays differ; programmatic edges only go kernel → kernel; also with STRATA_DF_PDL=2.
    • New: verify_batch_parity (every batched form against its per-token path, bitwise).
    • Existing: mmvq_multi_parity, kv_q4_parity, bf16_gemv_parity.

Not tested

  • HIP and CUDA 12.3–12.9 builds were not compiled here.
    • HIP takes the plain launch.
    • griddepcontrol is only compiled under !__HIPCC__ && __CUDA_ARCH__ >= 900.

Niko1221 and others added 30 commits October 2, 2026 09:35
… a failed tuning says why (Niko1221#447)

tools/calibrate.py read any "gpu" list as a layer split and added --layer-split auto, so a one-card config with a
helper card for the expert tier ("gpu": [0] + --expert-cache-device1) was refused by the engine before it read a
weight, a "0,2" string never got its split, and split_skip_if_fits was ignored. The tuner now builds the arguments
with serve.server.engine_args, the same as a normal start.

When the tuning fails, setup prints the engine's own last error line from the log (only lines written during the
tuning), and says the failure again where it is not lost: before --calibrate starts the model, and in the
"All set." summary of an install that tuned.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…etup recommends the big card alone (Niko1221#448)

- engine (layer splits only): when a stage's cache cannot fund the chunk CUDA0 alone would choose, a WARNING
  names it (CUDA index, name, cache slots, the slots the chunk borrows) and points at the helper cache
  (--expert-cache-device1, docs/SECOND_GPU.md).  The chunk choice itself is unchanged: fits() is split into
  fits_one() with the same tests; one GPU never reaches the new log.
- setup: a pair whose second card is under 12 GB beside a first card of 12 GB or more is still offered, but
  the first card alone is recommended (the default with --yes and Enter; offer_together asks "n" by default),
  with a note.  A named --gpus pair is kept.  Equal pairs (2x 8/16/24 GB) are unchanged.
- The automatic drop of the small card from an auto split is not done: stages load their weights before the
  boundary search and the exact prompt-buffer bytes need the sessions made after it.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…the totals (Niko1221#457)

The server already read drafts_offered / drafts_accepted from the engine's DONE line but only returned them in a
response's timings (draft_n / draft_n_accepted). Each /metrics request entry now carries both (null when the
engine did not report them) and the totals keep their running sums since the server started, so a dashboard such
as llm-visuals can show the acceptance. Additive: no existing field changes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…he --family/--model to use (Niko1221#444)

With Unsloth's UD-IQ3_XXS or UD-Q2_K_XL in the folder, setup either said only "has no IQ3_XXS model file"
(--family unsloth --model IQ3_XXS) or took the UD-IQ3_XXS file for the GSQ-RCO IQ3_XXS by the size in its name.
Now:
- a chosen shard whose name is a quantization setup has no size for (UD-IQ3_XXS, UD-Q2_K_XL, a K-quant) stops
  with "<file> is UD-IQ3_XXS, a GGUF Strata cannot run" and what Strata runs: ISTA-DASLab's GSQ-RCO files (and
  Swift 1.5's and the Coder's) and Unsloth's UD-Q4_K_XL only;
- a folder without the chosen file but with other GGUFs lists the ones Strata cannot use and the --family/--model
  of the ones it can (an empty folder keeps check_shards' "missing");
- a size the family does not have names the families that do (--family qwen --model IQ3_XXS), and with
  --gguf-dir also what Strata runs.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e unpruned sizes for that (Niko1221#438)

README and docs/MODELS.md said only "Weaker outside coding". Niko1221#438 measured Chinese answers wrong or looping on the
Coder where English was fine; both now say plainly that this includes Chinese and other CJK text and point to the
sizes that keep every expert (Q2_0, IQ2_XS, IQ3_S).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#461), the low-RAM resident mode on Windows 32 GB PCs (Niko1221#467), fixes Niko1221#460 Niko1221#459 Niko1221#446 Niko1221#447 Niko1221#457 Niko1221#448 Niko1221#444; setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The fill from the GGUF (--stream-experts) read each expert's three slices,
copied the blob and waited for the copy before the next read: ~12,000-18,500
experts with neither the SSD nor PCIe kept busy. Now the slots are admitted
in profile order first (same placement), reads go in file order by up to 8
threads into page-locked batches of 64, and a batch's copies overlap the next
batch's reads. STRATA_FILL_SERIAL=1 keeps the serial fill.

Tested on the B70: cold page cache, Coder fill 23.4 GiB 76.2 s -> 18.8 s;
launch to first token Coder 82 s -> 26 s, IQ2_XS 120 s -> 41 s; outputs
identical to the serial fill.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… read (Niko1221#471)

The engine's `strata serve: prompt ...` line now says "= 0 reused + 12288 of 98179 read in ... (cancelled)" with the
rate from the tokens read, and DONE gets a 15th field: the prompt tokens actually read. The server's /metrics history
and totals count reused + read for a cancelled read (prompt_total keeps the whole prompt); an older engine's DONE line
(no such field) is read as before.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Engine: when the MTP draft head's subset does not fit, it prints what the head needs, the VRAM free, and the smaller
shipped subsets (--draft-vocab cyrillic / en, their MiB from the head's real row size) with how to pick one
(START-HERE.bat / ./setup.sh --draft-vocab en, kept in the model's config). Text only, on the failure path.
Server: a start that ended that way carries the engine's hint (or the same advice for an older engine) in its
"the engine exited before it was ready" error. Setup: a card under 14 GB with the default subset gets a note
recommending --draft-vocab en; nothing is changed, an explicit or saved choice gets no note, and the golden
configs are unchanged.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…o1221#475)

In a git clone they run `git pull --ff-only` (a zip copy is told to download the new zip), then `setup.py --update`:
what a plain start does to an install before it starts the model - the Python packages, the ready-made engine when
MIN_ENGINE is newer (a compiled one when its source changed), each model's config upgrades and its draft subset, the
Windows AMD runtime - and nothing else: no question, no model files, no start. Both scripts run from one parsed
block, since the pull can rewrite the running script. Documented in docs/INSTALL.md (Updating) and the README.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…profile-save, Niko1221#477)

`--expert-profile-save P` (--serve) writes the adaptive tier's learning as a profile in tools/make_profile.py's
format, byte for byte (read_expert_profile reads it back): every pair ranked - the experts resident in VRAM now
first, then the routing counted since the start (the decayed usage summed at each adapt step, before its decay),
then the profile the engine started from. On QUIT / stdin EOF and every --expert-profile-save-every minutes
(default 10, 0 = exit only) between requests, through <P>.tmp and a rename. Without the flag nothing is counted or
written: `heat` stays empty and every added loop is over it.

Server: "expert_profile_save" (and "expert_profile_save_every") in the config pass the flags, and the next start
begins from the saved file instead of the config's --expert-profile when its header matches that profile's
layers x experts. tests/core/expert_profile_save_test.cpp (CPU only) checks the order and the bytes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…, the draft-head hint (Niko1221#474), UPDATE.bat (Niko1221#475), --expert-profile-save (Niko1221#477); setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1)

Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32
values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync),
then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in
registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue)
and one down launch into the per-slot rows the existing combine reads.  Experts are read where they are (cache
slot or streamed ring slot, a pointer per expert): no gather.

Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with
STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU).  Everything else - the default, native/IQ packs, small
chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch).

tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS
1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5.  One layer at 2048 tokens x top 10
over 512 experts: MMQ path 11.12 ms, fused 3.13 ms.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…e layer's experts) by default

The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch
waited for slots its own first batch frees.  With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM
ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides).

RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 ->
2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%.  1024 slots: 4K +24-27%, but the 32K chunk falls
from 8192 to 6144 (+6%).  The default path's output is identical to 0.1.35's on both prompts.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
… hold only what the fused kernels need

In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the
grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H
and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the
GDN set does).  A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for
stream_all_min() - 1 tokens.  The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default
stays 512).  Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events.

The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert
of the layer, routed or not - at a streamed chunk's size nearly all are.

RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots):
  4K:   MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046)
  32K:  MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153)
  128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346
        (8192), + 1024-slot ring 2,392 (6144)
A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per
8K tokens.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…ault (two layers' experts)

With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots,
under the 90% cap).  RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating
pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to
6144 and is still faster.  Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656
(+21.5%), 128K 2,128 -> 2,454 (+15.3%).  The ring does not change the numbers (the same output as at 512 slots).
Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…y argmax, the same ids and tokens

- qsa_block_topk: calls of up to 16 queries (decode windows) run on a cluster of 8 CTAs per query
  (block_topk_cluster_kernel): keys in shared memory, digit histograms pushed through distributed shared memory,
  one cluster barrier per radix pass, a parallel digit search, cluster-wide offsets for the ordered emit. The
  threshold and the cut at it are the one-CTA kernels' (same tie rule, NaN keys 0). STRATA_QSA_CLUSTER=0: off.
  RTX 5070 per call: 21.9 -> 15.6 us at 32K, 58 -> 18 at 128K, 200 -> 22 at 262K.
- sample_tokens greedy without penalties: a cluster of 8 CTAs per row (sampler_greedy_cluster_kernel), four loads
  in flight per thread, results merged in CTA 0 (lowest index on ties, NaN never taken, 0 when none).
  STRATA_ARGMAX_MULTI=0: off. 39.6 -> 5.9 us per call at 248,320 logits.
- Both launch through cudaLaunchKernelEx with the cluster attribute (captured in the decode graphs) only where
  they run: CUDA, compute capability 9+ (STRATA_EMULATE_CC respected), the kernel's own PTX/SASS at sm_90+ (an
  older build's JIT copy is a trap), room for one cluster - checked once per device. HIP and sm_75/86/89 keep the
  one-CTA kernels.
- decode_cluster_parity (ctest; skips below sm_90): bitwise vs the one-CTA dispatch and the reference kernel over
  434 top-k cases (contexts 7..262,144, every tail weight, 1-16 queries, ties / NaN / +-inf / -0 / denormal / all
  NaN) and 216 argmax cases, plus a CUDA graph capturing both (first calls inside the capture); --bench timings.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1)

src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert -
cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of
the native packs.  The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign
tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit
loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which
multiply at m16n8k16).  The activations are int8 per 32 values in natural order (quantize_act_native); fragments by
ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add.  A work item is 256
weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128).

Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the
FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers.  sm_75, HIP and the
default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy
tokens, 4K and 32K).

tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision
reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 /
2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5),
IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 /
2.72e-2 (4.4e-5); both tile shapes.  One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192
tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6,
12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5.  MMQ's products alone (no gathers) at
3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them.
(The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be
queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0.
prefill_fused_moe_test's timing part has the same pattern.)

End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt
tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K
1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%).  The IQ3
packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s
of 3.4 s), so faster kernels change little.  Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at
162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89,
IQ3_S 31 / 256.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…(STRATA_PF_FUSED=0 off), the native IQ packs' fused kernels opt-in (=1), cluster decode kernels; the fused test waits for its uploads

The ring rule: 1024 slots for the Q2_0 pack's fused experts (P3), 512 for native packs with fused layers.
Quality (teacher-forced vs the FP16 path, Q2_0): fused as close as MMQ at 8K, closer at 32K (phaseA-tests s18-tf).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
device, expert_source, native_head, prefill gemm, prefill.cpp, generate.cpp
merged from dpct of aeb35be and d9ab843 with sycl/tools/merge_upstream.py (the
3-way merge restricted to files upstream changed, dpct kernel-name hashes
canonicalized to the port's). Version 0.1.35.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…ls) and their switches

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
The PCIe probe sets the share of the missed experts the GPU reads over
the link (--pcie-frac): the native default 0.55 from 20 GB/s up, less
on a slower link (Niko1221#44). On Windows HIP it timed its copies with events
that do not bracket them there: an RX 6800 read between 3,348 and
25,999 GB/s in 24 engine starts, so every link kept the x16 share. A
standalone copy of the probe on that card: its events measure 0.13 ms
for the 1 GiB, the host clock 43 ms (25.1 GB/s).

On Windows HIP the probe now times the same copies on the host, between
two synchronizes (1 GiB, tens of milliseconds): the same RX 6800 on
PCIe 4.0 x16 reads 27.2 GB/s in every run, which keeps 0.55 there. On
a PCIe 4.0 x8 or 3.0 x16 link the share now shrinks as Niko1221#44 intended.
CUDA and Linux keep the events.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…DGET)

On Windows, hipMemGetInfo (ROCclr's PAL backend) reports the card's size
minus this process's own allocations; what the desktop and other
programs hold on the card is not subtracted. On a card that also drives
the desktop the engine counts the ~930 MiB those hold as free, and
--expert-cache auto fills the card past what Windows lets the process
keep in VRAM: WDDM moves part of it out to system RAM. On an RX 6800
(HIP SDK 7.2) with setup's configuration: a 9.83 GiB cache, the card at
16132 of its 16368 MiB, 929 MiB more of the engine in system RAM than
with this change, 30.1-31.0 tok/s decode. On a 0.1.30 build,
--vram-reserve-mib 1500 left the card as full (16037 MiB), since the
reserve is taken from the same figure.

Windows gives every process a video memory budget that does account
for the others (DXGI QueryVideoMemoryInfo). On that card it is 799 MiB
below the card's size while the process is small and about 1.8 GiB
below once it holds 11 GiB; allocations past it still succeed.
cudaMemGetInfo, which the HIP shim maps to hipMemGetInfo, now lowers
hipMemGetInfo's free figure by what the budget withholds (the card's
size minus the budget), on Windows only. What is sized from it then
stays within the budget: --expert-cache auto sizes the cache from it,
and the existing check after the slots are written shrinks it ("only
8 MiB free once the slots are written ... shrinking"): an 8.05 GiB
cache, the card at 15230 MiB, 40.6-41.5 tok/s.
Lowering HIP's own figure, instead of taking the budget minus DXGI's
usage, keeps memory the HIP runtime has freed but still caches counted
as free, as HIP counts it.

dxgi.dll is loaded at run time, so nothing new is linked; the adapter
is found by the device's LUID. Without dxgi.dll, the adapter or a
budget below the card's size, hipMemGetInfo's figure stands. The first
query logs the budget once. STRATA_WDDM_BUDGET=0 turns it off. CUDA
and Linux builds are unchanged.

Claude-Session: https://claude.ai/code/session_015Ld5YDzuKW1hgMVwwgRZDP
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…g forever (Niko1221#481)

A request whose engine prints nothing for engine_silence_s (default 300 s; while a prompt is read, a chunk may take
3x the previous one's time, the first up to its tokens at 50 tok/s more) ends with EngineSilent (an EngineDied): the
engine is killed and the next request starts it again. The STOP drain, under the request FIFO, is no longer an
untimed wait. "engine_silence_s": 0 keeps the old behaviour. Tests: unit (silence, PP chunks, first chunk, STOP
never acknowledged, 0, config) and over HTTP with a fake engine process (non-stream, stream, stop token).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…0 log lines (Niko1221#496)

Every start failure, not only the draft-head case (Niko1221#474), now carries the engine's own reason: the last 20 lines
this start wrote to the engine log (never an earlier start's). The failed start also closes its pipes and its
handle on the log. Tests: the tail itself, and a fake engine that exits with a reason in its log.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
imanu86 added a commit to imanu86/Strata-2080Ti that referenced this pull request Oct 5, 2026
…n batch

Porta Niko1221#904 di Hardin22 (commit 8e1324a e 1d4f833): rami
paralleli nel grafo della finestra di verifica (STRATA_DF_BRANCH, attivo su
CUDA), accodamento KV e post-operazioni PLE in batch (STRATA_DF_KVAPP,
STRATA_DF_PLE, STRATA_DF_QB_Q), PDL solo da sm_90 (STRATA_DF_PDL, non scatta
sul 2080 Ti). Applicata a tre vie senza conflitti sul codice del fork.
Verifica sul 2080 Ti: teacher forcing a posizionamento fisso identico bit per
bit al Daily 0.1.39 su codice e spiegazione; a 131k L01 47,0 / L02 43,3 t/s
contro 45,9 / 41,9.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@blange48

blange48 commented Oct 5, 2026

Copy link
Copy Markdown
Contributor

Feedback on a 4-way split (rig: 4 x RTX 5080 16 GB (PCIe Gen3 x16, sm_120), i9-10920X, Linux, IQ3_S, 262K context, int8 K/V, layer split 12,24,36 with --trim-stage-weights, --batch 8 --batch-groups 4, MTP spec 4): applied on top of 0.1.39 together with other patches, in production since today.

  • Exact: tools/batch_test.py gives slots identical to solo with 8 and 3 slots (greedy, --pcie-frac 0 --adapt-every 1000000, STRATA_IQ_MT_MIN=1), and parking_test.py is identical.

  • Speed through the HTTP server, all switches on (default) vs STRATA_DF_PDL=0 STRATA_DF_BRANCH=0 STRATA_DF_QB_Q4=0 STRATA_DF_KVAPP=0 STRATA_DF_PLE=0, total tok/s:

    Concurrent off on
    1 115.2 122.1
    2 151.6 150.6
    4 252.2 251.3
    8 367.1 372.1

So about +5 % for one request (the solo path with MTP drafts) and neutral to +1 % in batch windows, consistent with your +3 %. These are single runs; our run-to-run noise is about ±3 %.

  • No NVRM/Xid messages in dmesg across the engine starts and stops.

…Linux)

With the mixer's side branches in the verify graph, a reviewer on Linux (open
kernel modules 610.43.02 with GSP, 2x RTX 3090 and also 2x RTX 2080 Ti on
NVLink) saw a stage's window stop inside its graph, at the GDN layer's fork,
when an NVML query (nvidia-smi, any field) ran between requests: 80/80 clean
requests with STRATA_DF_BRANCH=0, 30/30 clean without the query, stalls with
both. No PCIe replays, Xid or AER. On Windows (4060 Ti + 5080) 15 long
requests with a query after each did not stall. Until the cause is known the
branches are opt-in (STRATA_DF_BRANCH=1); they were worth ~2% of decode there.
The other switches of this PR keep their defaults.
@Hardin22

Hardin22 commented Oct 5, 2026

Copy link
Copy Markdown
Contributor Author

Thank you for running it on 4x 5080 in production. It's useful to see slots identical to solo with batch 8, and +5% on single requests.

One change since then: STRATA_DF_BRANCH is now opt-in (f89d31f). On #905, adambenhassen bisected a stall on Linux (open modules 610.43.02, GSP) to the branches in the verify graph combined with an NVML query (nvidia-smi, any field) between requests:

  • the window stops inside its graph at the GDN fork;
  • 80/80 clean with STRATA_DF_BRANCH=0, 30/30 clean without the query;
  • on 2x 3090 and on 2x 2080 Ti over NVLink;
  • not seen on Windows.

PDL, the Q4_0 KV batching and the PLE post-ops keep their defaults. If your monitoring polls NVML while the engine serves, I'd turn the branches off until we know more.

imanu86 added a commit to imanu86/Strata-2080Ti that referenced this pull request Oct 5, 2026
Sorgente congelata 14433b0 (r2: Niko1221#904, Encapsulate, Niko1221#910, Q2_0 esatto, Niko1221#930, Niko1221#949, Niko1221#965), eseguibile shipping b0edf8455031; manifest 0.1.39 conservato come -0139, manifest del 5/10 storici, check_source legge prima source-manifest.json.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
@Niko1221 Niko1221 closed this Oct 6, 2026
@Niko1221

Niko1221 commented Oct 6, 2026

Copy link
Copy Markdown
Owner

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.

architectds pushed a commit to architectds/Strata that referenced this pull request Oct 6, 2026
--pipeline-windows and --adapt-async 1 excluded each other: the decode loop of
--pipeline-windows 2 adapts with windows in flight, and the asynchronous tier
assumed nothing in flight at a tick. Now --pipeline-windows 2 --adapt-async 1
run together. --pipeline-windows 1 still turns the asynchronous tier off (said
once), and so do exchange buffers that are not page-locked (that check is
kept: beside the windows every copy is queued by the loop's thread, and a
pageable source would make it a staged, synchronous copy there).

In the pipelined loop the tier's four steps advance at every verdict (after
the drafter's chain is launched, off its critical path) and, while a round is
in flight, every 0.5 ms between verdicts without counting a window. Windows
are in flight at every tick (window K+1 on the first card while the last
verifies or commits K), each one waiting on the loop's thread at every layer
(its doorbell) and planning every layer from the residency table as it is
when that layer is served. Two rules:

  - no CUDA call of the tier reaches a card while a window there waits on
    the loop's thread. A copy queued to a card whose window spins on a host
    flag can block inside the driver until that window ends (issue Niko1221#31), and
    a thread blocked there can hold up the loop the window waits for. So
    beside the windows the job thread does host work alone (the choice, the
    bounce copies, the RAM moves), and the loop's thread queues each card's
    copies itself at that card's next gap - a moment its stage has no window
    in flight: before a window is launched there, and at every tick;
  - a gap of a stage after a step is also the step's fence for that card:
    every window planned there before the step has completed. Step 2's
    copies into the evicted slots go at a gap after step 2 (a window planned
    before may still read the old expert from its slot), and step 3's RAM
    moves (a new RamFence state) wait for a gap after step 3 on every card
    whose layers exchanged (a window planned before may still pull an
    incoming expert over PCIe from the RAM place its evicted one takes; with
    STRATA_EXCHANGE_ROTATE the flip that hands that place over comes after
    them too). No events, no host waits. The CPU's reads need no fence: the
    pool runs inside the loop's service calls.

The device copies of the residency table are not uploaded meanwhile (a
synchronous copy; the pipelined verifiers publish every row whatever they
say): once, when the loop has drained. At the loop's end (after pl_drain)
what still waits for a gap is queued and the round in flight goes on unfenced
at the next tick, the next request's drain. An error ends the engine as
before; die() and the loop's end wait only for the job's host work. The
serial loop's asynchronous tier is unchanged.

STRATA_PIPELINE_ADAPT_ASYNC=0, read with STRATA_PIPELINE_DEBUG=1 like the
other pipeline test variables, keeps the blocking tier beside the pipelined
windows (--adapt-async 1 is then off for the run, said once at startup).

Also, as run() does: the pipelined Verifier::service's timeout prints the
verifier's diag before it releases the GPU waits and the trace after it, with
STRATA_VERIFY_TRACE=1.

Measured on 0.1.40, 4060 Ti (layers 0-19) + 5080 (20-47), i9-14900KF, 32 GB
of RAM, the resident RAM mode with the whole copy in RAM, Swift 1.5 IQ3_XXS at
160K (q4_0 KV), the stock draft layer, --trim-stage-weights, greedy, five
prompts x two rounds, decode tok/s:
- --pipeline-windows 2, blocking tier (0.1.40 as released): 93.5, 94.1
- --adapt-async 1, serial loop (0.1.40 as released): 97.4, 102.1
- both (this): 110.2, 111.3, i.e. +18% over the pipeline alone and +11% over
  the async tier alone.
Stress, both on, STRATA_PIPELINE_FORCE_MISS=2 STRATA_PIPELINE_THETA=0: 12
sequential 1500-token coding requests at temperature 0.6 with an nvidia-smi
query after each (as in the stall report on Niko1221#905, since traced to Niko1221#904's
graph branches): no timeout, no degenerate answer; the tier ran 1552 rounds
(105,435 experts swapped in) at ~40 ms per round. The asynchronous tier is
not bit-exact by design (Niko1221#876), so there is no exact comparison; on 0.1.40
the serial loop itself also gives a different greedy text for a different
window size (--spec 4 vs 2: 2/4 prompts identical).

Checked: builds (MSVC + CUDA 13.3, sm_89/sm_120).
(cherry picked from commit ad4eaaa)
zsariboga pushed a commit to zsariboga/Strata that referenced this pull request Oct 7, 2026
A decode window is a long chain of small kernels, and on a fast card many of
them are over before the next one has fetched its weights. With programmatic
dependent launch the window's quantizations and dense projections (multi-column
MMVQ, multi-row BF16 GEMV) start while the kernel before them is still running:
they load their first weights, then wait (griddepcontrol.wait) for its output.
Nothing is written before the wait, so every value is the one an ordinary
launch computes.

- pdl.hpp: pdl_wait / pdl_trigger and launch_pdl. The attribute is used only
  inside the verify window's capture (pdl_scope), with a CUDA 12.3+ runtime, on
  sm_90+ (STRATA_EMULATE_CC counts), for a kernel whose loaded code is sm_90+
  (ptxVersion and binaryVersion >= 90: older PTX, JIT-compiled, has no wait and
  would race), and only when every node the capturing stream's next node
  depends on is a kernel. So there is never a programmatic edge after a memcpy,
  memset, host or event node. After a join of branches each joined kernel gets
  a programmatic edge (the wait covers them all); STRATA_DF_PDL=2 allows a
  single predecessor only.
- On CUDA the PDL kernels' activations and outputs are no longer __restrict__
  (the compiler could otherwise load them before the wait); the weights keep
  it. For every architecture: the host pass must see the device pass's
  signature, or MSVC rejects the template stubs.
- HIP: ordinary launches, no-op device functions, the kernels' plain loops.
- copy_rows_strided replaces the q/gate split's cudaMemcpy2DAsync on the
  batched QSA path: a memcpy node would end the PDL chain there. A copy: the
  same bytes, PDL on or off.
- The profiler's stamps pass the early launch on inside the window.
- STRATA_VERIFY_NODES also counts the programmatic edges and those that are not
  kernel to kernel (0 expected).
- STRATA_DF_PDL=0: ordinary launches.

Byte-identical: PDL changes when a kernel starts, not what it computes. Off on
pre-sm_90 cards (the 4060 Ti) and on HIP.

Checked: pdl_parity (new test) captures a verify-shaped chain (quantize,
small-K and large-K MMVQ, BF16 GEMV, the shared Q8_1 buffer rewritten, a
memcpy node, a fork and a join) with and without PDL: bitwise equal outputs
over 200 replays with every buffer NaN-poisoned before each; programmatic
edges only kernel to kernel, and never into a node that also follows a
non-kernel node; present on sm_90+, absent elsewhere. In our fork the same
kernels gave identical windows over full requests (per-layer hashes of the MoE
inputs and the final residual).

Measured together with the next commit (PDL is on the 5080 only: the 4060 Ti, sm_89,
has none); the numbers are there.

PR Niko1221#904's PDL commit ported onto 0.1.40 (8 conflicts resolved against 0.1.40's EXACT_T bf16 kernels, Niko1221#783's rows rule, TSUM and the w2/y2 pair outputs; launch_pdl needs the defaulted arguments spelled out). pdl_parity passes on sm_120; no window-time change here.
Niko1221 added a commit that referenced this pull request Oct 7, 2026
…y path, PLE post-ops in one pass (all opt-in)

Ported from #904 onto the 0.1.40 verify window. Every switch is OFF by default, so
default output and the default launch sequence are unchanged:

- STRATA_DF_PDL=1: programmatic dependent launch (pdl.hpp) for the window's
  quantizations and dense projections (multi-column MMVQ, multi-row BF16 GEMV);
  sm_90+ only, CUDA 12.3+; HIP and older cards keep ordinary launches.
- STRATA_DF_PLE=1: the window's hyper-connection write and PLE post-ops in one
  pass, which also writes the commit's history snapshots (native_gr_post_multi
  with strides, native_ple_postops_batch_snap).
- STRATA_DF_QB_Q4=1: a Q4_0 KV cache takes the batched QSA query path too.
- STRATA_DF_BRANCH=1: the mixer's side branches inside the window graph.
  Left opt-in because the author saw a stall on Linux with NVML queries.

The q/gate split now uses copy_rows_strided (a kernel) instead of a
cudaMemcpy2DAsync so a PDL chain stays kernel to kernel; the bytes are the same.

Left out: STRATA_DF_KVAPP (kv_append_q4_step_multi). 0.1.40 already appends a
window's K/V cells in one launch (kv_append_q4_steps, #783 PR-d).

Merged by hand with the 0.1.40 forms of the same code: the exact-T BF16
multi-token kernel, the PAIR/TS MMVQ variants, QFUSE (q8_attn), the fused
norm+rope and the batched residual write.

Co-authored-by: Francesco Albano <119073350+Hardin22@users.noreply.github.com>

(PR #904)
@Niko1221

Niko1221 commented Oct 7, 2026

Copy link
Copy Markdown
Owner

Thanks, Hardin22. The PDL, PLE post-op, Q4_0 query-path and graph-branch parts are in 0.1.40.2, credited as co-author, all opt-in (STRATA_DF_PDL, STRATA_DF_PLE, STRATA_DF_QB_Q4, STRATA_DF_BRANCH). The Q4_0 KV append part was already in 0.1.40 (kv_append_q4_steps), so it was dropped. Defaults stay off until a 10-pair A/B on sm_90+ hardware says otherwise.

0xrjman pushed a commit to 0xrjman/Strata that referenced this pull request Oct 7, 2026
…y path, PLE post-ops in one pass (all opt-in)

Ported from Niko1221#904 onto the 0.1.40 verify window. Every switch is OFF by default, so
default output and the default launch sequence are unchanged:

- STRATA_DF_PDL=1: programmatic dependent launch (pdl.hpp) for the window's
  quantizations and dense projections (multi-column MMVQ, multi-row BF16 GEMV);
  sm_90+ only, CUDA 12.3+; HIP and older cards keep ordinary launches.
- STRATA_DF_PLE=1: the window's hyper-connection write and PLE post-ops in one
  pass, which also writes the commit's history snapshots (native_gr_post_multi
  with strides, native_ple_postops_batch_snap).
- STRATA_DF_QB_Q4=1: a Q4_0 KV cache takes the batched QSA query path too.
- STRATA_DF_BRANCH=1: the mixer's side branches inside the window graph.
  Left opt-in because the author saw a stall on Linux with NVML queries.

The q/gate split now uses copy_rows_strided (a kernel) instead of a
cudaMemcpy2DAsync so a PDL chain stays kernel to kernel; the bytes are the same.

Left out: STRATA_DF_KVAPP (kv_append_q4_step_multi). 0.1.40 already appends a
window's K/V cells in one launch (kv_append_q4_steps, Niko1221#783 PR-d).

Merged by hand with the 0.1.40 forms of the same code: the exact-T BF16
multi-token kernel, the PAIR/TS MMVQ variants, QFUSE (q8_attn), the fused
norm+rope and the batched residual write.

Co-authored-by: Francesco Albano <119073350+Hardin22@users.noreply.github.com>

(PR Niko1221#904)
Niko1221 added a commit that referenced this pull request Oct 8, 2026
#1469, #1470)

2e4ddf6 (the #904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
W1nge pushed a commit to W1nge/Strata that referenced this pull request Oct 8, 2026
Niko1221#1469, Niko1221#1470)

2e4ddf6 (the Niko1221#904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
Niko1221 added a commit that referenced this pull request Oct 8, 2026
#1469, #1470)

2e4ddf6 (the #904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
zsariboga pushed a commit to zsariboga/Strata that referenced this pull request Oct 8, 2026
Niko1221#1469, Niko1221#1470)

2e4ddf6 (the Niko1221#904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
imanu86 added a commit to imanu86/Strata-2080Ti that referenced this pull request Oct 11, 2026
…su CUDA come nel Daily

Il Daily (be4cf1c, port del PR Niko1221#904) li accendeva di default su CUDA: rami laterali
del mixer nel grafo della finestra (+2,4% a 131k sulle due schede) e passata unica
dei post-proiezione PLE, bit per bit il ciclo per token (+3,3%). In 0.1.42 sono
opt-in: DF_BRANCH per lo stallo su Linux (Niko1221#905, mai visto su Windows).

DF_BRANCH: di default solo con CUDA su Windows, Linux e HIP restano opt-in.
DF_PLE: di default con CUDA, HIP opt-in. =0 li spegne entrambi.
Compilato si, provato no.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
giridhart pushed a commit to giridhart/Strata that referenced this pull request Oct 15, 2026
Niko1221#1469, Niko1221#1470)

9ba9d57 (the Niko1221#904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
giridhart pushed a commit to giridhart/Strata that referenced this pull request Oct 15, 2026
Niko1221#1469, Niko1221#1470)

9ba9d57 (the Niko1221#904 port) removed __restrict__ from the activation pointers of
native_quantize_q8_1_kernel, bf16_f32_mmvf_multi_kernel and
native_mmvq_multi_kernel on every CUDA architecture, and turned on the weight
prefetch before pdl_wait(). Both only matter for PDL (sm_90+). On sm_61 the
missing restrict loses the read-only data path (LDG.CI) and decode halves
(36.4 -> 17.7 tok/s on a Tesla P40, 32-36 -> 17.7-18.6 on a GTX 1080 + 1070).

STRATA_PDL_RESTRICT is __restrict__ and kPdlPrefetch false again for
__CUDA_ARCH__ < 700 in the device pass. sm_70 and newer compile to the same
SASS as before (cuobjdump -sass of native_mmvq.cu and native_bf16.cu at sm_70,
86, 90 and 120 is identical); the host pass is unchanged.

Reported and bisected by paulhothersall and lineape.
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.