Skip to content

Intel arc 0.1.39 + perf fixes - #809

Closed
maxfridbe wants to merge 863 commits into
Niko1221:mainfrom
maxfridbe:intel-arc-0.1.39
Closed

maxfridbe wants to merge 863 commits into
Niko1221:mainfrom
maxfridbe:intel-arc-0.1.39

Conversation

@maxfridbe

Copy link
Copy Markdown

Follow-up to #423 on top of the 0.1.39 release. Every change is in sycl/ or the Intel docs; no shared upstream file is touched.

1. The 0.1.39 SYCL port did not build (fixes)

Two upstream changes merged for the release did not reach the port's copies of the files they touch. On the B70 (oneAPI 2026.1, AOT bmg-g31) sycl/ failed with 3 errors:

With those two 3-way merged, 0.1.39 builds with 0 errors. On the card, 21 of 25 kernel tests pass, and benchy matches the 0.1.38 port's speeds. Of the four failures:

Not ported yet: #559's batch slots (--batch / "parallel"). The port's verify.cpp / verify.hpp do not have them. It is opt-in, so default use is not affected.

2. Prompt-path speedups

Measured on a B70 with each model's served config. Decode speed is unchanged.

  • Lend mirror. The slots a prompt borrows have their experts re-read from the GGUF for every chunk and again for the refill. Those experts are now kept in pinned RAM and DMA'd. Same bytes, identical outputs. About 2 GB more RAM and about 2 s more at start. STRATA_LEND_MIRROR=0 turns it off.
  • QSA block scores as oneMKL GEMM tiles. The old kernel used one sub-group per (query, block) pair, so its cost grew with the square of the context. Now the pooled keys are multiplied against the batch's indexer queries in fp32 tiles, followed by a small relu-sum kernel. 7-8x faster at 40K-256K, relative error 2-4e-7 vs FP64. STRATA_SELECT_GEMM=0 restores the old kernel.
  • Prompt attention scores one work-item per cell (128-cell chunks), instead of a 5-step shuffle tree per head. 1.5x faster. Decode keeps the old kernel. STRATA_ATTN_PERCELL=0 restores the old kernel.

The last two are FP32 in a different summation order, so they are not bitwise, like the CUDA build's tensor-core paths. Greedy outputs follow the same text and part ways at a near-tie after tens of tokens.

Time to first token, Coder IQ1_M on its served 256K config (measured on the 0.1.38 base):

Prompt Before After
2,185 tokens 4.9 s 4.6 s
40K 40.1 s 33.6 s
128K 140.9 s 111.2 s
256K 328.8 s 235.1 s

IQ2_XS and Swift 1.5 at 128K: 191 -> 156 s.

Things tried that did not help are recorded in sycl/TODO.md: 15 attention variants, including two XMX versions, and expert grouping on the GPU (its "6.5%" was mostly the profiler's own cost).

3. benchy v1: a standard bench for Intel results

sycl/benchy.sh runs every set-up model on fixed prompts (sycl/bench/v1: 20 / 2,185 / 8K / 40K / 128K / 256K tokens) with its own serve config, from a cold page cache. It reports:

  • speed: PP, TTFT, TG, drafts accepted;
  • resources: experts in VRAM and in the RAM mirror, VRAM, RAM, SSD reads, load time, power;
  • the system and each config's args.

docs/INTEL_PERFORMANCE.md holds every measured number: current results, dated history, the B50 / B580 / 2x B70 results from #423, and "Submitting numbers", which asks testers to post benchy's report plus their config.

4. Smaller

  • setup_intel.py --check sized each model by the first card only. With two cards it now reports what fits across both.
  • strata-sycl.sh: with no ONEAPI_DEVICE_SELECTOR set, a --layer-split gets level_zero:gpu (the image pins level_zero:0). A set selector is still passed through.
  • Duplicate stage_room fix removed: where INTEL B70 32gb support #423's fix and yours overlapped, this keeps yours.
  • Docs: an accuracy pass on docs/INTEL.md (build section, switches, a corrected "gather-bound" claim), plus the probe programs and three kernel benches (attn_bench, sel_scores_bench, xmx_int8_bench).

Testing

  • The 0.1.39 build fix: built AOT for the B70, kernel tests run, benchy v1 on all three models.
  • This exact branch: a full B70 run (build, ctest, benchy v1) is in progress. I will post its report as a comment.

Niko1221 and others added 30 commits October 2, 2026 12:43
…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
…p-experts (Niko1221#505)

The server window said "loading the experts into RAM (about N GB)" whatever the flags. Now: --mmap-experts maps them
from the model files (the OS file cache reads them, nothing is loaded up front; its "experts via mmap" line triggers
the message when the model has no PLE line), --resident-experts copies the ones the GPU does not hold, and
--resident-budget-gib N copies up to N GiB. setup's start banner says "maps" for the plain mapped mode too. Tests:
the words per flag set, and the narrator over a mapped and an arena log.

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

hf() takes the host from HF_ENDPOINT (default huggingface.co, a trailing slash dropped); the pinned revisions, the
fallback to main and the SHA-256 checks are unchanged. tools/mtp_fetch.py reads the MTP tensors from the same host.
Step 5 prints the folder setup expects the model files in and how to use files downloaded by hand (that folder with
their original names, or --gguf-dir), and the mirror when one is set. INSTALL.md says the same. Tests: the endpoint
in hf()/hf_unpinned()/mtp_fetch, and step 5's words over a mocked install. 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
…ds it; no budget on a --gpus start (Niko1221#498)

The engine refuses --resident-budget-gib with a layer split, so setup always dropped UD-Q4_K_XL to one GPU. With RAM
>= the GGUF size + UNSLOTH_RAM_LEFT_GB (~135 GB) and 2+ usable cards, setup now asks (one GPU stays the default and
--yes keeps it; an explicit --gpus is honoured; an explicit --resident-budget-gib keeps one GPU) and the split writes
"gpu": [..], "layer_split": "auto" without the budget: all the experts loaded into RAM at start, the files through the
OS file cache (the config Niko1221#498 measured: 2x RTX 3090, 31 -> 64-78 tok/s). Latent bug: START-HERE --gpus 0,1 (or the
start-time offer, or a hand-edited config) on a UD-Q4 config kept the budget and the engine exited with code 2; the
split now drops it, or setup stops before saving when the RAM is too small. docs/UNSLOTH_Q4.md says so. Tests in
tools/test_setup_unsloth.py; 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
The engine already leaves --vram-reserve-mib MiB of VRAM free (default 700). INSTALL.md now says what it does, its
default, and how to set it (the flag, or "--vram-reserve-mib", "2048" in the config's args). Setup's new flag writes
it into the config only when given (one value: it replaces the images' own 700, with a note when it is less than the
GPU image encoder needs); on a start of an installed model it is saved into that model's args, and a new copy set up
from an earlier install keeps it. The default config is unchanged (golden test). Tests in test_setup_choices.py.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…are, the samples logged

The probe timed one 1 GiB burst; an x16 PCIe 4 laptop (RTX A3000) read 18.5, 6.9 and 5.8 GB/s
in three starts, and the low reading set pcie_frac (0.12-0.15 where calibrate.py found 0.35).
The same 1 GiB is now four bursts of 256 MiB, each timed on its own (events between them on
the stream; on Windows HIP the host clock, Niko1221#377), and the fastest is the link's figure: a
disturbance only ever slows a copy, while the median would still follow one lasting through
half the bursts. The probe copies no more than before.

The share was base above 20 GB/s and base * bw / 26 below (19.9 GB/s -> 0.42, 20.0 -> 0.55;
4.0 -> 0.08, 3.9 -> 0). It is now base * min(1, bw / 20): continuous, unchanged from 20 GB/s
up (an x16 PCIe 4/5 link reads ~26-28 GB/s, so the default stays 0.55 there), in proportion to
the link below, so the link's time on its share stays about what the x16 share costs. The rule
is one function for CUDA0 and the layer split's stages.

The log names every burst: "PCIe probe: 28.9 GB/s host->device (best of 28.9 28.9 28.9 28.9)"
(the RTX 5070 here; pcie_frac 0.55 as before).
Noted, not changed: a given --pcie-frac is every stage's share and skips the stages' probes.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…#380) and the host-timed PCIe probe (Niko1221#377)

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…MD counts the desktop's VRAM (Niko1221#380 Niko1221#377 Niko1221#497), a steadier PCIe probe (Niko1221#485), fixes Niko1221#496 Niko1221#495 Niko1221#498 Niko1221#505 Niko1221#493; setup requires it

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
…iko1221#374 Niko1221#363 Niko1221#407 Niko1221#413 onto 0.1.35

One dpct run of upstream main + the seven PRs, 3-way merged into the files they
touch (sycl/tools/merge_upstream.py); their header changes are forked into
sycl/include until upstream merges them.

- Niko1221#385 stager waits for a buffer's previous DMA on every job (port: polled seq)
- Niko1221#463 decode waits for an adaptive swap before reading residency
- Niko1221#453 drafter's batched K/V on the KV-streaming ring
- Niko1221#374 first chunk's PLE rows read beside layer 0; port fixes: the next read
  starts at c0 + T (the port's first chunk is shorter), and the upload is
  marked by a polled sequence number instead of a host event wait, which
  deadlocked prompts past ~4K tokens under the Level Zero v2 adapter
- Niko1221#363 group stride + fused SwiGLU/q8_1 for the PCIe expert call, re-applied
  to the port's lane kernels (fp contract off, as upstream's HIP form)
- Niko1221#407 --adapt-tuned (opt-in), Niko1221#413 key-head DeltaNet recurrence (opt-in via
  STRATA_GDN_KEYHEAD=1: 8% slower on the B70; its NVIDIA gate and SM-holding
  parity bench not ported)

Tested on the B70: native_grouped_parity, iq_multi, gr, qsa, kv_hybrid,
quantize_act, qsa_prompt_attn parity pass; Coder 19/2,184-token, IQ2_XS and
40K outputs identical to 0.1.33; Coder 2,184-token prompt 784 -> 825 tok/s;
128K int8 KV streaming 895 / 67.0 tok/s; IQ2_XS decode +1.7% (Niko1221#363).

Co-authored-by: sergqwer, constantindjonkam, architectds, BlueKingMuch (upstream PR authors)
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…er_intel)

- engine (serve mode): GEN key logprobs=K (0..20) writes "LP logprob id:lp ..."
  after each "T id", from the verify window's head logits (Verifier::copy_logits;
  before temperature, penalties and sampling); absent key, no LP lines
- server_intel.py: pairs LP lines with their tokens (StrataEngine._pump), passes
  logprobs/top_logprobs to the engine, and fills OpenAI's choices[].logprobs
  (content / reasoning_content) in streamed chunks and the collected reply; the
  stop token is not listed. serve/ itself is unchanged.

Tested: a fake engine on hardac (streamed, not streamed, without logprobs,
stop token excluded); on the B70 with the IQ2_XS: same answer with and without,
"Paris" -0.0019 with top-1 equal to the chosen token, no measurable time cost.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… wait, released by one event

In the streamed walk every expert put a cross-stream wait (its copy) and an event record (its ring slot back) on the
compute stream around a ~3 us gather. Under WDDM each pair left ~10 us of GPU idle even when the copy had long
landed: an nsys trace of a 32K prompt showed 226 ms of idle between back-to-back gathers while the host was 3-11 ms
ahead of the GPU. Now a group of 16 is gathered by one launch (copy16_group_kernel) after one wait on its last
streamed copy - the copy stream is in order - and one event releases all of its ring slots: the issuer waits on
used[used_of[slot]]. A ring entry the routing skipped inside an open group first gathers what the group holds so
far, so no more than a group's entries are ever held back from the issuer. The same bytes in the same group slots.
STRATA_PREFILL_GROUP_GATHER=0: one at a time, as before.

IQ2_XS, RTX 5090, 262K context: 32K prompt -8.6%, 8K -10%, 2K -12.6%; first-token logits bit-identical at 2K, 8K,
16K and 32K and with a 96-slot ring.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
… 256 at a time

Layer 1 adds the per-layer embedding to every token of a chunk, so its rows must be in. A later chunk's rows were
already read on a thread during the chunk before; the first chunk's were read synchronously before layer 0, so the
GPU idled while the SSD read them (an nsys trace: ~245 ms of a single-chunk 32K prompt). Now the first chunk's
gather starts on a thread too and lands just before layer 1 (ple_land), which also starts the next chunk's.

The reads themselves went 64 at a time: a 32K prompt is 512,720 rows in 213,358 4 KiB reads, ~620K IOPS at queue
depth 64. --ple-inflight now defaults to 256, where the drive saturates (1024 is no faster). Decode reads 16 rows a
token and is unchanged. The same rows: the same results.

IQ2_XS, RTX 5090: a 32K prompt 5986 / 6013 -> 5843 / 5808 ms, 8K 1714 -> 1624 ms; first-token logits bit-identical.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
…e checked

From an audit: the upload of a chunk's rows, the event recorded after it and the wait on the other buffer's event
before the SSD thread refills it were unchecked, so a failed record or wait let the thread write a buffer an
upload might still read. A failure now ends the chunk with the CUDA error.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
…e previous DMA from it

Stager::work waited for a buffer's DMA only from the generation's kRing-th job on (j >= kRing); a generation's first
kRing jobs wrote their buffers straight away. The previous generation's last DMAs from those buffers are only
known to be done where something else waits for them: the routing sync before a staged layer's start(), or the
sync at the end of a chunk with an MTP drafter or stages. A stream_all chunk that ends without one (no --mtp), or a
ring entry the routing skipped (its copy is never waited for), leaves a window where the next chunk's first jobs
overwrite a pinned buffer a DMA may still read. The event is now waited for on every job; an event never recorded,
or already complete, returns at once.

IQ2_XS with --mmap-experts --resident-budget-gib 16 (12,144 of 48,668 streamed experts through the stager): the
same tokens on a 32K prompt, prefill 6,491 vs 6,546 ms.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
… and STRATA_DRAFT_TIMING

With KV streaming (--kv-resident) the MTP drafter's K/V is a ring of its window (kv_mode 2: page p in slot p % n_slots
over a host copy), and draft_kv refused it, so every prompt chunk fell back to the drafter's own pass - one graph per
few rows. The batched path takes the ring with its own page table and host copy (as a streamed main layer appends);
the cells it writes are those the window can still reach, which the ring holds, and one batch never spans more than
the ring's slots, so no two cells of a batch share a slot. STRATA_MTP_BATCH_RING=0: the drafter's own pass (A/B).
STRATA_DRAFT_TIMING=1 (debug): the pass's time, synced.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…g on pinned-host <-> pageable copies

A queue memcpy/memset with neither end in device memory (USM host pinned <-> pageable host) runs on the B70's
copy engine, which hung on it: dmesg "Engine reset: engine_class=bcs", and after repeated runs the device was
declared wedged (needs a reboot). conversation_snapshot_test's fill() into its pinned KV pools hit it every run
(the urEventWait "hang" was the driver resetting the engine under the copy), and so would the engine's
conversation_kv_save / restore of a --kv-resident session, whose pools are pinned host memory (the opt-in
conversation cache, Niko1221#596). transfer() and the test's fill / fill_after now memcpy / memset on the CPU when
sycl::get_pointer_type says both ends are host or pageable; copies with a device end stay on the queue.
On the B70: conversation_snapshot_test 2169 checks passed, no engine resets before or after.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
@maxfridbe

Copy link
Copy Markdown
Author

Update: conversation_snapshot_test is fixed in 5780652 and passes on the B70 (2,169 checks, no copy-engine resets). A queue memcpy with no device end (pinned host ↔ pageable host) hangs the Arc's copy engine (Engine reset: engine_class=bcs; repeated runs wedged the card). Host-only copies now go through the CPU. The engine's conversation_kv_save/restore had the same copy for --kv-resident sessions, used by the opt-in conversation cache (#596). Now 22 of 25 kernel tests pass.

… as upstream does; s2 grouped kernels bitwise

- iq_parity ran as `iq_parity --selftest`, which read "--selftest" as its fixture directory and failed. Registered as
  the top-level CMakeLists does: an iq_parity_fixtures setup test (tools/iq_fixture.py, when a python3 with numpy
  exists) and iq_parity on its output, both a skip on exit 3. On the B70 with generated fixtures: all 10 types ok
  (dequant rel 0, mmvq rel 4.6-5.7e-3).
- ple_parity needs ggml's captured block (bench/micro/ple_in.bin, ple_out.bin), which the repository does not
  carry: registered only where the captures exist (STRATA_PLE_FIXTURE_DIR), instead of failing everywhere.
- s2_expert_grouped_parity: the new transposed grouped kernels write `acc += dw * dx * v` in one expression, which
  -fp-model=precise contracts to an FMA; the previous kernels compute the product in chunk_dot and add it in the
  caller, which it does not (nvcc contracts both, so they are bitwise equal on CUDA). The last bit differed (both
  within 1e-7 of the host reference). No contraction in gu/down_grouped_t_kernel; fixups.py keeps it. Also an
  S2_DEBUG=1 print of the first differences in the test.

ctest on the B70: 23 of 23 pass (iq_parity not registered in the image: no numpy), no engine resets.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
@maxfridbe

Copy link
Copy Markdown
Author

b56c45e: iq_parity and ple_parity are registered as the top-level CMakeLists does. The port ran iq_parity --selftest, which read the flag as its fixture dir. With generated fixtures, all 10 types pass on the B70. ple_parity needs the bench/micro captures, which aren't in the repo.
s2_expert_grouped_parity: under -fp-model=precise the new grouped kernels' acc += dw * dx * v became an FMA, while the old kernels' chunk_dot product plus a separate add did not. nvcc contracts both, which is why it was bitwise on CUDA. The new kernels no longer contract, so the test is bitwise here too.
5780652 (earlier): conversation_snapshot_test. Host-only queue copies hung the Arc's copy engine; they now use the CPU.

maxfridbe and others added 3 commits October 4, 2026 21:59
… host, B70 and B65

- build.sh sources setvars.sh with -u off: under -u it ended the script outside the image.
- STRATA_SYCL_RUNNER: native (inside the toolbox) or distrobox:<name> (from the host), for
  sycl/serve/strata-sycl.sh and benchy; docker stays the default. The image's single-card
  selector (level_zero:0) is kept unless ONEAPI_DEVICE_SELECTOR is set.
- benchy records the card(s) each run held its VRAM on ("Ran on"), and says that the power
  column sums every card's energy counter when there are several.
- gpustat: card names from lspci -vmm (-mm's fields did not split), the B65's id.
- INTEL.md: the toolbox path, including the Level Zero loader it needs (>= 1.21; Debian 13's
  1.20.6 fails the first KV allocation with UR_RESULT_ERROR_UNSUPPORTED_FEATURE).
- INTEL_PERFORMANCE.md: benchy v1 on a Ryzen 9 9950X host, B70 and B65 at PCIe Gen5 x8. The
  B70 against the earlier Gen3 x8 host: IQ2_XS/Swift decode +15-30% and long prompts ~25%
  faster, Coder prompt +15-40%. The B65 runs at 70-77% of the B70 at ~70% of its power.
  The earlier host's other tables are kept under their own heading until re-run here.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
- strata-sycl.sh and benchy run the strata-sycl-dev image with podman, SELinux labels off for
  the data root's mount (an unlabeled data disk is unreadable to a confined container).
- benchy runs as root; podman's commands (image check, run, kill) run as the user who ran
  sudo, whose store holds the rootless image.
- INTEL.md: building and running with Podman. Its benchy rows match the toolbox run to within
  noise (9 rows: decode within 0.1 tok/s, prompt within 0.3%).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…benchy passes a config's env

- benchy hands the engine the config's "env" map as serve/server.py does (it dropped it), which
  is also how its A/B variants set STRATA_PREFILL_STREAM_ALL / STRATA_PREFILL_TIMING.
- Re-run on the Ryzen 9 9950X host, B70 at Gen5 x8, the dev image under Podman (2026-10-04):
  the 32K test config (Coder 2,185 tokens at 980-984 tok/s), long contexts by KV format,
  stream-all against routed-only, borrowing on/off, the API, start time, VRAM/RAM, and the
  prompt-path time shares.
- Findings: k8v4 is now the fastest Coder KV format at 128K-256K on both counts (1,506 / 65.1
  tok/s at 128K against int8's 1,327 / 63.9); stream-all and routed-only read prompts at the same
  rate since the lend mirror; borrowing keeps decode 2-4% ahead; logprobs cost ~6% (the old
  "no measurable time" was within the earlier host's noise).
- The earlier host's tables move to History.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
@joeyjoe02

Copy link
Copy Markdown
Contributor

The testing and write-up below were done by Claude on my machine, at my request. I'm not a developer; I'm posting it because it may help with the conversation cache on Arc.

Tested this branch (baa84d9) on one Arc Pro B60 24 GB (8086:e211, OCuLink PCIe 4.0 x4, NEO 26.31, oneAPI 2026.1.1, AOT bmg-g21), with #866 and #889 applied on top. Flash-Next GSQ-RCO IQ2_XS, --max-context 65536 --kv int8 --spec 4 with MTP, --stream-experts --vram-reserve-mib 1024, 11,935 experts resident and 12,643 in the RAM mirror, per-layer waits (no STRATA_VERIFY_NO_HOST).

Good news first:

  • Builds with 0 errors; conversation_snapshot_test passes on the B60 (2,169 checks).
  • Prompt reading is faster than the 7ba023e port: a 21.5K-token cold prompt 29.9 s vs 35.2 s. Decode unchanged (~22 tok/s).

The opt-in conversation cache still produces token-0 output on this card (--conversation-cache-mib 8192 --conversation-cache-slots 4):

Engine log lines around the first bad reply, if useful:

conversation cache: parked 23036 tokens in 150.9 ms; parked=4 bytes=3798174316 evictions=5 snapshot_bytes=1695236392 reused_kv_bytes=343928448
prompt 18889 tokens = 18870 reused + 19 read in 688 ms (27.6 tok/s), 256 generated in 5848 ms (43.8 tok/s), drafts accepted 126 of 126, 6 checkpoints

(that pair is from the 7ba023e run, where I first saw it; this branch shows the same pattern, the bad replies all 256 generated ... drafts accepted 126 of 126.)

Happy to rerun anything on the B60 (with STRATA_VERIFY_TRACE or similar) if it helps narrow it down. For now I'm running with the conversation cache off.

🤖 Generated with Claude Code

maxfridbe and others added 5 commits October 5, 2026 08:24
…EMMs, int8 linears, fused decode kernels)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
… fp16 wins on large expert groups, int8 1.6-2.4x

On the B70 at the prompt path's shapes: dense projections (4,096 rows) oneMKL fp16 128-147 TFLOP/s, oneDNN fp16
1.00-1.14x, oneDNN int8 1.6-2.4x; expert gate/up at 256-512 rows oneMKL 28-32 TFLOP/s against oneDNN fp16's 54-56.
TODO: oneDNN fp16 for expert groups >= 128 rows (same precision) before any int8 work.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
…1.39 (B60 runs, int8 DPAS benches, spin pricing)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
….2-1.9x in isolation, not a decode win on 32 GB

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

From maxious/Strata_SYCL (branch b70-intel-arc-0139, their experiments 26-30 and 42; the code is theirs):
- STRATA_MMVQ_PREUNPACK=1: each Q6_K / Q5_K dense matrix unpacked once to signed bytes, so decode skips the
  per-token 6-bit unpack (native_mmvq, native_dense); STRATA_MMVQ_LOOP=1: the wide Q6_K kernel's 5+ columns as a
  loop. Both off by default: the default path is unchanged (B70, Coder k8v4: 916 / 75.2 tok/s at 2,185 tokens
  against 918 / 75.2, the same drafts accepted; 40K 1,571 / 69.1 against 1,571 / 69.2).
- On a 32 GB card the pre-unpacked copies cost ~2.5 GiB of VRAM: 1,313 fewer cache slots, decode 72.0 (-4%), as
  their own engine A/B found. For cards with VRAM to spare.
- q6k_preunpack_parity / q5k_preunpack_parity (pass on the B65), q6k/q5k_preunpack_bench, xmx_mmvq_bench (their
  int8 DPAS matvec: 1.2-1.9x the shipped kernel at decode widths on a B65; bench only).
- dpct::error_string / last_error (DPCT_CHECK_ERROR keeps the exception's message) and the standard sub-group
  helpers in sycl_math.hpp, which their files use; added beside, nothing replaced.
- TODO: their branch's other switches A/B'd on the B70 (none beats this port; their default build drafts once and
  then decodes without speculation at 40K, and STRATA_PREFILL_INT8 does not fit the 256K config's buffers).

Co-Authored-By: maxious <maxious@lambdacomplex.org>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
@joeyjoe02

Copy link
Copy Markdown
Contributor

This testing and write-up were done by Claude on my machine, at my request. I'm not a developer; I'm posting it because the benchy v1 tables here have no B60 rows yet.

benchy v1 on one Arc Pro B60 24 GB (8086:e211, OCuLink PCIe 4.0 x4, Ryzen AI 9 HX 470, NEO 26.31, oneAPI 2026.1.1, AOT bmg-g21). Original Flash-Next IQ2_XS, --max-context 65536 --kv int8 --spec 4 with MTP, --stream-experts --vram-reserve-mib 1024: 11,935 experts in VRAM, the rest in the RAM mirror, per-layer waits (no STRATA_VERIFY_NO_HOST). The v1 prompts and run shape from perf_matrix.py (fresh engine, --max-new 256 --greedy --stats), two runs per size, the mean shown (the two runs agree within ~1% on PP and within ~10% on TG, which moves with draft acceptance).

Input tokens This branch (baa84d9) PP / TTFT 7ba023e port PP / TTFT TG this branch TG 7ba023e
20 28.4 / 0.9 22.9 / 1.0 31.7 31.6
2,185 359.3 / 6.3 322.1 / 7.0 38.5 35.0
8,000 684.6 / 11.9 575.3 / 14.1 36.2 38.0
40,000 809.8 / 49.7 666.6 / 60.2 32.8 33.6

So on the B60 this branch reads prompts ~19-22% faster at 8K-40K, and TG is the same within noise.

For context, #868 has 2x B60 (--layer-split) results and a single-B60 run on 6f32ec0 (Coder IQ1_M, 11.9 tok/s with experts mirrored). These are single-B60 numbers on this branch with benchy v1, so they line up with the B70 / B65 tables.

Differences from a standard benchy run, so the numbers are read correctly:

Correction to my comment above: the "~22 tok/s" decode figure there came from a short synthetic prompt where drafts are rarely accepted. On these prompts it's 32-38 tok/s, and real agent conversations on this card average ~41 tok/s.

🤖 Generated with Claude Code

…3% of one)

Every process's fdinfo lists each card it has open; the VRAM and busy sums now keep the clients whose drm-pdev is the card being reported.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
@TheEmeryverse

Copy link
Copy Markdown

Tested the PR #809 branch on an Arc B580 (G21) with Qwen3.8 IQ2_XS.

Startup: The checked-in wrapper forces STRATA_VERIFY_NO_HOST=1 and doesn’t set STRATA_MIRROR_MIB. Startup reported 21,216 experts neither resident nor mirrored. With our B580 wrapper (STRATA_MIRROR_MIB=8192, leaving STRATA_VERIFY_NO_HOST unset), 5,989 experts were mirrored in 8 GiB and the model served a request. Since both settings differed, this doesn’t isolate the effect of the verification flag.

Benchmarks: In one synthetic QSA run at 32,768 cells / 256 queries, the current kernel took 5.895 ms; TF32 GEMM took 0.336 ms (17.52×), with relative score error of 2.25e-7 and identical selected IDs. At 128K context, across five timing repetitions, per-cell attention took 0.981 ms versus 0.990 ms for the engine kernel, with exact output.

Across two valid runs per configuration on a 2,185-token prompt with 256 greedy output tokens, the candidate stack (PR #809 plus #866/#889 and local B580 safety changes) averaged 314.01 tok/s prefill and 31.18 tok/s decode, versus 311.96 and 31.18 for the existing B580 build; TTFT was 7.138 s versus 7.184 s. These small differences are within noise and don’t establish an end-to-end gain.

The up-projection XMX GEMM completed but was 0.25–0.33× as fast as dequantize plus oneMKL. The int8 DPAS prototype was 0.56–0.73× as fast as the reference, including input quantization. Q6_K pre-unpack parity, MMVQ, and down-projection XMX runs encountered UR_RESULT_ERROR_DEVICE_LOST; Xe engine resets were logged during these tests. B580 testing of the experimental paths is still in progress.

@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.

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.