Repository navigation
Conversation
The migration turned `q = cudaStreamQuery(cs)` into `DPCT_CHECK_ERROR(cs->ext_oneapi_empty())`. That macro yields 0 whenever
the call does not throw, whatever the queue holds. The guard `q != 1 && *seq < want` (verify.cpp, session.cpp) therefore
fired on any ring that took more than 2 ms and reported "never rang (graph finished)" while the graph was still running.
With every expert resident the host does not wait per layer (STRATA_VERIFY_NO_HOST), so the path measured on the B70 never
reached it. With host-mirrored experts the first layer whose GPU time passes 2 ms does: layer 15 of the Coder on one Arc Pro
B60, every time.
q is now 0 when the queue is empty and 1 while it is running, as the guard and the trace comment ("1 = still running") expect.
fixups.py gets the same rewrite (entry 8), so a fresh dpct run of session.cpp produces it too.
Measured on one Arc Pro B60 (4,042 of 12,288 Coder IQ1_M experts mirrored in RAM, 128 greedy tokens, AOT bmg-g21): before,
exit 1 at layer 15 in 5 of 5 tries; after, exit 0 in 7 of 7 runs, at 11.2-11.9 tok/s with coherent text.
|
(The testing and write-up below were done by Claude, on my machine. Follow-up: a fix for the flag-visibility cost is in #889.) Thanks for this fix, and for the B60 notes in #868. I can confirm both on a second B60, plus one result on the Setup: one Arc Pro B60 24 GB ( Before this PR: without Same three requests in every mode (short answer, 300-token answer, 21.5K-token prompt), thinking off:
So on this card the 200 bound looks fast only because the device stops waiting before the host's data is visible: whenever the host path is actually used, the output collapses into a single repeated token (this is with a larger, more host-dependent mirror than the Coder run in your notes, which may be why you did not see it). It seems consistent with your reading that the remaining cost is the GPU not seeing the host's flag store in time, rather than the spin bound itself. Happy to rerun anything on this card if it helps chase that. Two smaller observations, in case they are useful:
Output quality was very good in my own tool-calling task set, so thank you to everyone working on the Arc port. |
The 0.1.39 paragraph said nobody had run the 0.1.39 port on an Arc. Two Arc Pro B60 reports (#867, #870) did, with the two sycl/ fixes now in 0.1.40 (#784 the compile fix, #866 the ring wait). Say that, and that an unpatched 0.1.39 port did not compile against its own engine sources. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
|
Fixed in 0.1.40: merged, so a ring wait over 2 ms is no longer reported as finished (the fixup in sycl/tools/fixups.py carries it for a re-migration). Compile-checked with oneAPI 2026.1; the two-B60 reports are the runtime evidence. |
and Niko1221#1054 fixed Upstream main's own SYCL build had stopped compiling (fused_gr, expert_source, prefill/gemm against 0.1.40's headers): sycl/ still held 0.1.39 copies. Re-migrated with dpct; 3-way merged with base 5047172, the commit the port was migrated from (the branch point dropped 0.1.39b's 88 commits silently); histogram diffs; 175 conflicts and three hand-ported files resolved; dpct's mistranslations in the new code fixed (docs/INTEL.md, 0.1.40). - Niko1221#866: the per-layer ring waits treated every layer slower than 2 ms as 'never rang' (DPCT_CHECK_ERROR of ext_oneapi_empty() is always 0) whenever experts were mirrored and STRATA_VERIFY_NO_HOST was off. - Niko1221#1054: with a layer split the host mirror holds only the first GPU's layers. Its pinned memory belongs to that GPU's context; the later GPU's cache fill copied from it (two B70s hung at startup), and every later-stage expert was mirrored for nothing (13-18 GiB of RAM). - mmvq_multi_parity: the negative control has no power on the port's wide kernels (one layout at every column count), reported like gfx906's wave64 case. - vmm.cpp: upstream's no-VMM branch; gemm_bf16_parity (cuBLAS) and native_expert_bench (raw CUDA) not built. - Engine version 0.1.40-sycl (setup.py's MIN_ENGINE). ctest 34/34 on the B70, no engine resets. IQ2_XS greedy output identical to the 0.1.39 port; a K=29 layer split over B70 + B65 matches one card token for token. Decode 63.5 -> 68.0 tok/s (one card), 57.9 -> 62.7 (split). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
and Niko1221#1054 fixed main's SYCL build had stopped compiling: sycl/ still held its 0.1.39 copies. Re-migrated with dpct and 3-way merged against 5047172, the commit the port was migrated from. 175 conflicts and three hand-ported files are resolved, and dpct's mistranslations in the new code are fixed (docs/INTEL.md, 0.1.40). - Niko1221#866: the per-layer ring waits counted any layer slower than 2 ms as "never rang". - Niko1221#1054: with a layer split, the host mirror holds only the first GPU's layers. The later GPU no longer reads another context's pinned memory, and 13-18 GiB of RAM is no longer spent. - Engine version 0.1.40-sycl. ctest 34/34 on a B70. IQ2_XS greedy output is identical to the 0.1.39 port, and a K=29 split matches one card. Decode 63.5 -> 68.0 tok/s on one card, 57.9 -> 62.7 on the split. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Rb4SRxFuLpffQAAXBf4HxA
The 0.1.39 paragraph said nobody had run the 0.1.39 port on an Arc. Two Arc Pro B60 reports (Niko1221#867, Niko1221#870) did, with the two sycl/ fixes now in 0.1.40 (Niko1221#784 the compile fix, Niko1221#866 the ring wait). Say that, and that an unpatched 0.1.39 port did not compile against its own engine sources.
The 0.1.39 paragraph said nobody had run the 0.1.39 port on an Arc. Two Arc Pro B60 reports (Niko1221#867, Niko1221#870) did, with the two sycl/ fixes now in 0.1.40 (Niko1221#784 the compile fix, Niko1221#866 the ring wait). Say that, and that an unpatched 0.1.39 port did not compile against its own engine sources.
3 participants
The SYCL port's per-layer ring wait gives up too early. With experts mirrored in host RAM, a layer whose GPU time passes 2 ms is reported as
layer N never rang (graph finished)while the graph is still running. On one Arc Pro B60 the Coder IQ1_M run stopped at layer 15 every time.Where it broke
The CUDA code asks the stream whether it is still running, and only gives up when it is not (
src/core/verify.cpp:1458):The migration turned the query into
const dpct::err0 q = DPCT_CHECK_ERROR(cs_->ext_oneapi_empty()).DPCT_CHECK_ERRORevaluates to 0 whenever the expression does not throw, soqis always 0,q != 1is always true, and any ring that takes longer than 2 ms is a failure. The same pattern is insession_run_token(sycl/src/core/session.cpp) andVerifier::run(sycl/src/core/verify.cpp).With every expert in VRAM the host does not wait layer by layer (
STRATA_VERIFY_NO_HOSTwaits for the whole window), so the paths measured on the B70 never reached it. With part of the experts mirrored in RAM the host does wait, and the first layer slow enough to trip it was layer 15 of the Coder.The change
qis now 0 when the queue is empty and 1 while it is running, which is what the guard and the trace comment ("1 = still running") already expect:Two sites, 8 lines added and 3 removed in the engine files, plus a
fixups.pyentry (#8) that produces the same code from a fresh dpct run, so a re-migration keeps it.Verified
One Arc Pro B60 24 GB, Coder IQ1_M, 4,042 of 12,288 experts mirrored in pinned RAM,
--max-context 8192 --kv int8 --spec 4, 128 greedy tokens, AOTbmg-g21, Level Zero V2, NEO 26.09.37435.12, oneAPI 2026.1.1, Ryzen 5 5600 with 64 GB RAM:verify: layer 15 never rang (graph finished), 5 of 5 tries;The all-resident path is unchanged (two B60s,
--layer-split, every expert in VRAM: Coder 56.6 tok/s, IQ2_XS 58.6 to 61.3, same as before the change). Setup and numbers are in the results report, #867.Notes
verify.cppby maxious's fork (maxious/Strata_SYCL, branchb70, "DPCT_CHECK_ERRORdiscards the value of its expression", fixed there as an explicitidle = cs_->ext_oneapi_empty()). Thesession.cppcopy (session_run_token) is still unfixed in that fork and inmaxfridbe/Strata_B70at the time of writing. This PR covers both sites on upstream's tree, with a reproduction on a card.sycl::atomic_fence(seq_cst, system)instrata_spin_pause(). Before this fix it fails at layer 15 the same way, so it is not the cause of this failure. After this fix the run with the fence did 11.4 tok/s and without it 11.2 to 11.9, so it made no difference here. It is not part of this PR.wait for rings, and it falls with the device-side spin boundkSpinMax(25.1 tok/s at 2,000, 28.4 at 200). That looks like a separate problem, the GPU not seeing the host's flag store, and I have not changed it.src/core/verify.cpphas two moreq != cudaErrorNotReadywaits (:1983,:2098) with no SYCL copy yet. They need the same treatment when ported.6f32ec0sycl/does not build without it.