Repository navigation
Conversation
ba2e1af to
3e85d6e
Compare
|
Tested #1390 (73acfb5, on main 0.1.40.3) on 2x Arc Pro B70, Windows 11, driver 32.0.101.8864, official oneAPI 2026.1.1 offline installer, AOT bmg-g31. On this machine the driver does have a Level Zero adapter: sycl-ls lists both cards as level_zero:gpu (UR over Level-Zero V2) next to opencl:gpu. So the Windows path runs with command graphs and oneMKL here, with no eager replay and no GEMM fallback. The missing adapter in the PR description may be the conda toolchain rather than the driver. |
…x two migrated-pointer bugs found by TeppeiLan1104 (Niko1221#1390) TeppeiLan1104 tested this branch on 2x Arc Pro B70 (Windows 11, driver 32.0.101.8864, official oneAPI 2026.1.1, AOT bmg-g31) and found five problems on top of it. Three of them are in the engine and two are in the port's tooling, so they are taken here; the other two - the mirror allocated in 8 GiB pieces past max_mem_alloc_size, and the setup_intel/CMake changes for the official installer's clang-cl mode - need a machine this port cannot be tested on (a 32.9 GiB mirror, a second card), so they are left to their branch, which is the better place for them anyway. Their comment on Niko1221#1390 has the measurements. **The PLE staging line.** Splitting capture_commit() into record_commit() dropped `if (ok && ss.ple.ready() && ple_stage()) copy_indexed(ss.ple.hist, hist_snap_, HS, commit_ + 1, HS, cs_);` at the end of the body. hist_snap_ is what stage_inputs() reads the history back out of, so without it a window restores a stale history. Their Coder IQ1_M greedy output changes without the line (hash 8769f8bd4768) and is byte-identical to the 0.1.39 build's with it (376092cc1dd5), at the same speed. Restored verbatim from origin/main. On this machine (eager replay, one window per request) the line is not reached: hist_snap_ is written in stage_inputs() and only read back by the *batched* commit (verify.cpp:2720), which needs captured graphs. Checked with --batch-mtp 2 and STRATA_DEC_BATCH=0 - both run the unbatched path here, so the fix is verified by their hashes on a graph backend and by code reading here, not by a local A/B. The commit body otherwise matched upstream's. **ext_stream_ was not "no stream".** Upstream's `cudaStream_t ext_stream_ = nullptr` migrated to `dpct::queue_ptr ext_stream_ = &dpct::get_in_order_queue();`, which is *evaluated* when the Verifier is constructed: "none" became "the constructing device's in-order queue". A later stage of a layer split is constructed on CUDA0 and init'ed on its own card, and init()'s `if (ext_stream_ != &dpct::get_in_order_queue()) cs_ = ext_stream_;` then took CUDA0's queue: stage [17, 48) had cs_ on 0000:0d:00.0 with sh_cs_/copy_ on 0000:b0:00.0, which fails the first window ("cannot submit to a queue with a dependency from a graph that is associated with a different context", or UR_RESULT_ERROR_DEVICE_LOST with STRATA_SH_STREAM=0). nullptr = none again, and both tests against the in-order queue become tests for nullptr. Reported upstream separately - this one is not Windows-specific. **The host-mirror slice was unbounded.** resident_plan() found its slice of the mirror table from the pointer alone (`res_layer >= g_mirror_res`), which compares addresses, not layers: a layer-split stage whose own residency table happens to sit above this one read a slice far outside the table, which shows up as a device lost with experts short of the second card. The difference in int32 entries is the layer index, and it is now checked against the table's own layer count, which the setter takes. One card and eager replay cannot reach it either (the mirror table is null when nothing misses). sycl/tools/fixups.py carries both migrated-pointer corrections so a re-migration keeps them. quantize_act_parity and sampler_parity: 0 failures. tools/test_setup_intel_windows, test_setup_sycl, test_setup_amd and test_setup_choices pass.
|
Thank you - this was a genuinely useful review, and two of the three engine fixes you found are now on Taken, and why:
Left for your branch, deliberately:
So please do open the PR against
One measurement of yours corrects something I reported: I filed #1549 for the prompt path collapsing to 1.03 tok/s Also worth saying plainly: 82 tok/s on one B70 with graphs against 76.4 here for the same model and prompt length is |
|
Correction to my review, and one correction to yours - both about context size on a 32 GB card. I measured the resident-KV sweep you did not ask about, on my side, to find where the card stops taking a bigger
Medians of three interleaved runs, 128 greedy tokens after a 16,386-token prompt. Two things I did not expect: The resident KV does not really compete with the expert cache. The expert cache is 12,288 slots / 23.42 GiB in all 131,072 is the ceiling, and past it the card is not merely slower - the prompt never finishes. At If you have a moment on your setup, Not asking you to run anything before the PR - the PR is yours to open whenever, and the mirror piece and the clang-cl |
|
Level Zero on Windows: root cause found, one artifact missing @TeppeiLan1104 — I chased down why 1. UR never even loads its Level Zero adapters. The dependency chain is 2. UR then crashes with 0xC0000005 (Event 1000: faulting module 3. The driver never installs its own Level Zero files. 4. Registering the paths manually gets the loader to find and load the ICD — and then: Calling the ICD's exported table getters directly ( 5. The ICD itself is fine. Driving it with the old v1.2.3 loader works: But v1.2.3 exports only So the missing piece is a loader in the window: >= 1.8.0 (exports Could you tell me:
My driver is 32.0.101.8976 (2026-08-27) against your 8864, so I suspect 8976 swapped in a V2 loader without a matching V2 ICD. If that holds it is worth reporting to Intel, since it would break Level Zero for anyone on 8976 on Windows. Meanwhile this branch stays on OpenCL/eager on my box: |
Re-applied on top of upstream 0.1.40.2 (the release merged most of the earlier forward-port, and has its own STRATA_VERIFY_EAGER window replay, so this branch carries the Windows deltas and nothing else). Build and setup - setup.py names the Intel Arc on Windows (display adapters + the display-class registry's real VRAM, which WMI caps at 4 GB) and --backend sycl hands over to sycl/setup_intel.py there too instead of stopping at "no Windows setup yet". - sycl/setup_intel.py: the native exe, sycl/serve/strata-sycl.bat, host paths, the AMD-prebuilt/HIP-device bypass, and STRATA_DEVICE_FREE_MIB/TOTAL_MIB (the free query needs Level Zero sysman; kept across a rewrite when detection flakes). - sycl/CMakeLists.txt: IntelLLVM/MSVC flag branches (/fp:precise is load-bearing), installer-free oneMKL (headers + mkl_sycl_blas.lib, default beside the profile), and the SYCL/UR/OpenCL/oneMKL DLLs staged beside strata.exe. - sycl/tools/build.bat imports the VS environment itself and finds libircmt beside the compiler through the build dir's CMakeCache when the shell has no icx on PATH (otherwise the link fails LNK1104). Port fixes the new release needs on Windows - verify_kernels.dp.cpp: template <bool OUT> is a Windows SDK macro (minwindef.h `#define OUT`, reached through <sycl/sycl.hpp>), so the parameter vanished and every use was a parse error. Renamed WRITE_OUT. - generate.cpp: memmem is POSIX (std::search over the same bytes); `small` is rpcndr.h's `#define small char`. sampler_parity.cpp: `near` is windef.h's. - pinned.dp.cpp: the Niko1221#243 shared-memory budget query is behind CUDA headers now. - gguf_expert_source.cpp: positional reads through OVERLAPPED (no pread here). - sampler_parity: fixture 18's graph half is skipped without the Graph extension instead of killing the process; sampler_parity reports 0 failures on OpenCL. Without command graphs (the OpenCL adapter has none) - verify.cpp: record_commit() is the commit body, replayed by commit(); eager replay refuses a host-served expert plan loudly (its mid-record waits expire into stale plans - silently wrong, not slow). - mtp.cpp: capture_round/capture_step/capture_prefill bodies are record_*() functions the drafter replays; prepare_chain refuses (its events come from a recorded graph). - gemm.dp.cpp: oneMKL SYCL BLAS has no OpenCL Xe2 backend and its failed attempt poisons the queue, so the prompt GEMM is a plain SYCL kernel there (beta = 0 no longer reads the caller's uninitialized Y). Measured (Arc Pro B70 32 GB, Win10, driver 32.0.101.8976, oneAPI 2026.1.1, OpenCL, STRATA_VERIFY_EAGER=1, Coder IQ1_M all resident): quantize_act_parity and sampler_parity byte-exact; decode 47.7-50.5 tok/s with the MTP draft layer (Linux B70 with graphs: 78.2), 24.4 tok/s without, prompt 65-68 tok/s warm. The profile puts 24 ms of a 65 ms window in the GDN hyper-connection read, nearly flat from T=2 to T=6: dispatch without graphs, not arithmetic.
… it costs 47.7-50.5 tok/s with the draft layer, 24.4 without, 66-68 prompt, parity exact - against 74.7-77.4 for the previous 0.1.39-based Windows build of this port, so the 0.1.40 changes do not help on a backend without command graphs. The stage profile puts 24 of 65 ms in the GDN hyper-connection read, flat from T=2 to T=6: dispatch, not arithmetic. Also the three Windows-SDK macro names that eat identifiers.
strata.exe statically imports libmmd.dll (Intel's own runtime, beside icx), and the UR adapters load ur_loader.dll / ur_win_proxy_loader.dll from beside the exe. A build made from a clean tree had none of the three, so the engine ended with 0xC0000135 (STATUS_DLL_NOT_FOUND) before printing anything - it only ran here because an earlier build directory had them. Found by building the previous 0.1.39-based branch from scratch for an A/B and starting it. Also updates the Windows section with the A/B: the same test on that previous build runs 63.9-67.1 tok/s back to back against 47.8-51.8 here, with 80% of its drafts accepted against 72%.
… and a correction The earlier numbers came from 155-token answers through the API, which are not comparable on this backend: the same build and mode gave 48.9 and 67.5 tok/s on two consecutive 128-token runs, and a 155-token run contradicted a 256-token one. Interleaved medians over 256-token greedy runs instead: | | 0.1.39-sycl | 0.1.40-sycl | |---|---|---| | decode, 256 tokens, --spec 4 --mtp | 69.7 | 77.3 | | decode, 128 tokens after a 2,048-token prompt | 37.6 | 39.6 | | decode, 256 tokens, --spec 2 (no draft layer) | 53.6 | 22.3 | | prompt reading, 2,048 tokens | 14.5 tok/s | 14.5 tok/s | So with the draft layer this release is about 11% faster here, and the one reproducible regression is --spec 2 without it. The "not a speed-up" claim is withdrawn.
oneMKL SYCL BLAS has no OpenCL Xe2 backend, so on Windows every prompt GEMM runs through the port's plain SYCL kernel. That kernel gives each work-item its own four columns of W, so it reads the whole weight matrix once PER ROW: a 2,048-row chunk of one 2,560 x 4,096 projection moved 86 GB of weights (T x N x K x 2 bytes). Measured on the B70: a 2,048-token prompt read at 8.0-14.5 tok/s, where docs/INTEL.md's Linux rows read at 790-1,002. fallback_gemm_tiled gives a work-group a 64 x 64 corner of Y and walks K in 16-wide tiles, staging the activations and the weights in local memory, with a 4 x 4 register corner per thread. Each weight element is read once per 64 rows and one load feeds 16 multiply-adds instead of one. The dispatcher keeps the old kernel for T < 64, which is what a decode window asks for (its per-row weight reads are cache hits there, and a 64-row tile would leave most of the block idle) - decode is unchanged, measured. Both kernels accumulate over k in ascending order and spell the fused multiply-add out (sycl::fma) so the compiler cannot round them differently from each other. STRATA_FALLBACK_SELFTEST=1 runs both over one problem whose rows, columns and k count all cross the tile edges and reports the bitwise difference: 0 of 4,189 values, largest 0. Measured (Arc Pro B70 32 GB, Windows, OpenCL, Coder IQ1_M, greedy, interleaved medians over 2 runs of a 2,048-token prompt): prefill 8.0-14.5 -> 71-73 tok/s, decode 42.2 -> 43.8 tok/s (unchanged). quantize_act_parity and sampler_parity: 0 failures. Note on method: the model's own greedy output is not a check for this. With a 2,048-token prompt the engine is not run-to-run deterministic on this backend at all - the same binary diverges at token 60 - so the self-test, not the answer, is what shows the two kernels agree.
…here 71-73 tok/s on a 2,048-token prompt (was 8.0-14.5), proven bit-identical to the kernel it replaces by STRATA_FALLBACK_SELFTEST (0 of 4,189 values). Also: with a long prompt the engine is not run-to-run deterministic on this backend, and the web app answers temperature 0 with its own sampling defaults - so neither is a check for a change here.
The A/B for the tiled fallback (and the way out if a card turns up that the
tiled kernel behaves badly on). Measured in one binary on the 0.1.40.3 engine,
medians of two, Coder IQ1_M, greedy, 2,048-token prompt:
prefill 73.5 tok/s tiled -> 14.5 naive (5.1x)
decode 25.4 tok/s tiled -> 25.4 naive (unchanged, as designed: a decode
window's rows are below the tile, so it takes the naive kernel)
Also re-applied on top of 0.1.40.3, which merged upstream's Intel setup text
and added the Arc A750 (i915) path; its STRATA_SYCL_BIN now stays posix, and
docs/INTEL.md keeps both the A750 and the Windows sections.
…67% at 2,048 context tokens) Upstream 0.1.40.3 made the device's spin bound on host flags a build option and defaulted it to 20,000 only for a *bmg AOT* build, 2,000,000 otherwise - the A-series JIT build, where the GPU genuinely waits on slow per-layer CPU expert work. A B70 built without AOT (JIT) therefore inherited the A-series value, and on this branch's OpenCL backend the waits are the other kind: a kernel's writes to host-mapped memory are not reliably visible while a window runs, so they are expected to run out and give up - every extra read is time spent on nothing. Same binary, one -DSTRATA_SYCL_SPIN_MAX apart, 128 greedy tokens, medians of two interleaved runs on an Arc Pro B70: | context | 2,000,000 | 20,000 | | |---|---|---|---| | 3-token prompt | 40.6 tok/s | 59.0 | +46% | | 2,048-token prompt | 25.6 tok/s | 42.9 | +67% | More of it the longer the context, which is what more expired waits looks like. With 20,000 this release decodes *faster* than 0.1.40.2 did (42.9 against 40.8 at the 2,048-token context, 59.0 against 51.0 short), so the regression was entirely this default and not the rest of 0.1.40.3. The carve-out is on Windows rather than on the card, because that is what the property is: the OpenCL backend's handshake, not which Arc is in the machine. A user who wants the long bound still passes -DSTRATA_SYCL_SPIN_MAX=... . quantize_act_parity and sampler_parity: 0 failures. Through the API, 2+2 answers 4 at 36 tok/s.
…drafts nothing) The bound added in 129c499 is right for the port's own configuration, and the shape of the thing turned out to be sharper than that commit says: what the bound decides is whether speculation works at all, not how long a window costs. A window that spins on a host flag it cannot see runs out and goes on without what it was waiting for. That is harmless when what it wanted was a plan it already has, and ruinous when it was a draft. Same build, one -DSTRATA_SYCL_SPIN_MAX apart, Coder IQ1_M with --ple-gguf, 256 greedy tokens after a 2,048-token prompt, medians of two interleaved runs: --spec 4 --mtp 2,000,000: 26.2 tok/s, 2.28 tokens/round, 73 of 126 accepted 20,000: 45.4 tok/s, 4.19 tokens/round, 99 of 99 accepted --spec 2 2,000,000: 68.8 tok/s, 2.17 tokens/round (2 suffix windows) 20,000: 21.8 tok/s, 1.03 tokens/round (0 drafted) --spec 4, no MTP 2,000,000: 69.7 tok/s 20,000: 18.9 tok/s With the MTP drafter on the short bound is both faster and better, because the MTP record path has already published what the window needs. Without --mtp the suffix drafter's own waits run out inside 20,000, it produces no draft at all, every round verifies a single token, and decode is 3.7x slower - silently, since the answer is still right. So the default stays 20,000 (it is what setup's --spec 4 --mtp wants) and the wrapper says so on stderr when the engine is run without --mtp, pointing at -DSTRATA_SYCL_SPIN_MAX=2000000 for that case. The upstream report is revised to match: one constant cannot serve both kinds of wait, because the drafting waits want the long bound and the window's want the short one.
The mirror that holds the experts which did not fit in VRAM took its default size from /proc/meminfo, which does not exist on Windows, so `cap` was 0 and nothing was ever mirrored here. A model that does not fit entirely in VRAM then either refused to run - "6889 experts are neither in VRAM nor mirrored; with STRATA_VERIFY_NO_HOST the device plan cannot run them" - or, with a host plan, read every miss from the SSD. GlobalMemoryStatusEx is the same reading the process memory stats above already take on this platform. With it, IQ3_XXS (39.97 GiB of experts against the B70's 32) runs: 17,687 of 24,576 experts in VRAM and the other 6,889 (11.15 GiB) in pinned host memory, which the GPU reads over PCIe. Measured on an Arc Pro B70, Coder IQ1_M vs IQ3_XXS, greedy, medians where noted: | | decode after a 2,048-token prompt | decode, 3-token prompt | prompt reading, 2,048 tokens | |---|---|---|---| | Coder IQ1_M (23.42 GiB experts, all resident) | 43.6 tok/s | 76.4 tok/s | 74.7 tok/s | | IQ3_XXS (39.97 GiB experts, 72% resident) | 30.0 tok/s | 19.2 tok/s | 1.03 tok/s | The prompt path is the interesting half: 1.03 tok/s. Decode copes with the PCIe misses (the device-built verify plan points the expert kernels at the mirror), but the prompt's GEMMs do not, and a 2,048-token prompt takes 33 minutes. So on a 32 GB card a model that cannot fit entirely in VRAM is decode-usable and prompt-useless, which is worth knowing before choosing one - the port's own configuration (all experts resident) is the one that reads a prompt quickly. quantize_act_parity and sampler_parity: 0 failures. tools/test_setup_intel_windows, test_setup_sycl, test_setup_amd and test_setup_choices pass.
The Monitor tab showed a name and a VRAM size for an Intel Arc on Windows and nothing else: GPU load, VRAM, temperature, power and PCIe were all empty tiles, because the port's reader (_XeGpu) has sysfs to read on Linux and nothing here - no sysfs, no user-mode Level Zero adapter, so no sysman either. Windows does keep counters for every WDDM adapter, and two of them answer what was missing: pdh.dll's "GPU Engine" (per process, per engine) and "GPU Adapter Memory" (dedicated bytes in use). The card is the adapter holding the most dedicated memory - the discrete one a model runs on - and every instance is filtered to its LUID, so an iGPU beside it does not add to the numbers. Measured on an Arc Pro B70 through the API: 2.9% of the card idle, 100% while decoding, and 30.01 of 32.00 GiB of dedicated memory in use - the 23.4 GiB of experts, the KV and the driver's own buffers. The engine's own VRAM figures agree (12,288 of 12,288 experts resident). Two traps in this path, both found by measuring rather than reading: - PDH's wildcard API answers PDH_INVALID_ARGUMENT on this machine for every path, core counters included: PdhExpandWildCardPathW and PdhGetFormattedCounterArrayW. The instance list therefore comes from `Get-Counter -ListSet` and each instance gets its own counter. Sampling is then pure PDH - 3.8 ms to collect the 392 instances of the two sets and 0.4 ms to read them - against a PowerShell call per sample. - A byte counter must be asked for as PDH_FMT_LARGE. With PDH_FMT_DOUBLE, PDH fills the union with a double, whose bits read as an integer are nonsense: 4.4 GB came out as 4,437,051,947 GB. An instance belongs to the process that made it, so the list is rebuilt when the card reads idle (a process that started since the last look has none yet) and at most every 30 s, the cost of one `Get-Counter -ListSet` (~100 ms warm). The reader is built once too: Telemetry.sample() asks ok() every second, and a fresh reader threw the bound instances away between samples. Temperature, power and the PCIe link have no Windows counter at all, so those tiles stay empty - now with the dashboard saying why rather than showing a bare dash (the same gpu_note an AMD card on Windows already used). serve/test_server.py: 262 tests pass. No GPU and no network are needed for them.
…x two migrated-pointer bugs found by TeppeiLan1104 (Niko1221#1390) TeppeiLan1104 tested this branch on 2x Arc Pro B70 (Windows 11, driver 32.0.101.8864, official oneAPI 2026.1.1, AOT bmg-g31) and found five problems on top of it. Three of them are in the engine and two are in the port's tooling, so they are taken here; the other two - the mirror allocated in 8 GiB pieces past max_mem_alloc_size, and the setup_intel/CMake changes for the official installer's clang-cl mode - need a machine this port cannot be tested on (a 32.9 GiB mirror, a second card), so they are left to their branch, which is the better place for them anyway. Their comment on Niko1221#1390 has the measurements. **The PLE staging line.** Splitting capture_commit() into record_commit() dropped `if (ok && ss.ple.ready() && ple_stage()) copy_indexed(ss.ple.hist, hist_snap_, HS, commit_ + 1, HS, cs_);` at the end of the body. hist_snap_ is what stage_inputs() reads the history back out of, so without it a window restores a stale history. Their Coder IQ1_M greedy output changes without the line (hash 8769f8bd4768) and is byte-identical to the 0.1.39 build's with it (376092cc1dd5), at the same speed. Restored verbatim from origin/main. On this machine (eager replay, one window per request) the line is not reached: hist_snap_ is written in stage_inputs() and only read back by the *batched* commit (verify.cpp:2720), which needs captured graphs. Checked with --batch-mtp 2 and STRATA_DEC_BATCH=0 - both run the unbatched path here, so the fix is verified by their hashes on a graph backend and by code reading here, not by a local A/B. The commit body otherwise matched upstream's. **ext_stream_ was not "no stream".** Upstream's `cudaStream_t ext_stream_ = nullptr` migrated to `dpct::queue_ptr ext_stream_ = &dpct::get_in_order_queue();`, which is *evaluated* when the Verifier is constructed: "none" became "the constructing device's in-order queue". A later stage of a layer split is constructed on CUDA0 and init'ed on its own card, and init()'s `if (ext_stream_ != &dpct::get_in_order_queue()) cs_ = ext_stream_;` then took CUDA0's queue: stage [17, 48) had cs_ on 0000:0d:00.0 with sh_cs_/copy_ on 0000:b0:00.0, which fails the first window ("cannot submit to a queue with a dependency from a graph that is associated with a different context", or UR_RESULT_ERROR_DEVICE_LOST with STRATA_SH_STREAM=0). nullptr = none again, and both tests against the in-order queue become tests for nullptr. Reported upstream separately - this one is not Windows-specific. **The host-mirror slice was unbounded.** resident_plan() found its slice of the mirror table from the pointer alone (`res_layer >= g_mirror_res`), which compares addresses, not layers: a layer-split stage whose own residency table happens to sit above this one read a slice far outside the table, which shows up as a device lost with experts short of the second card. The difference in int32 entries is the layer index, and it is now checked against the table's own layer count, which the setter takes. One card and eager replay cannot reach it either (the mirror table is null when nothing misses). sycl/tools/fixups.py carries both migrated-pointer corrections so a re-migration keeps them. quantize_act_parity and sampler_parity: 0 failures. tools/test_setup_intel_windows, test_setup_sycl, test_setup_amd and test_setup_choices pass.
…arly free; 196,608 stops) Measured with all experts resident so the KV cache has what the card has left, --kv-resident the only variable, 16,386-token prompt, 128 greedy tokens after it, medians of three interleaved runs. | --kv-resident | KV streams? | prompt reading | decode after it | |---|---|---|---| | 32,768 | yes | 178.4 tok/s | 40.7 tok/s | | 65,536 | yes | 178.6 | 42.1 | | 131,072 | no - all fits | 181.6 | 41.2 | | 196,608 | - | never finished (>25 min) | - | Within the range that works the resident KV is nearly free: the expert cache is 12,288 slots / 23.42 GiB in every run and the prompt borrows about the same slots throughout (1,036 / 1,036 / 969), and decode does not fall off (40.7 / 42.1 / 41.2 is inside the run-to-run noise, runs 39.2-45.9). The port's existing "about 3.6 GiB of VRAM to context costs ~6% of decode" note does not apply to a KV reservation that fits - the expert pool is sized from the free VRAM, so the KV is what gets what is left, not the reverse. Past the ceiling the prompt stops rather than slowing: at 196,608 the prefill ran past 25 minutes with the card at 100% and 30.83 of 32.00 GiB dedicated memory in use (the Windows GPU counters, added in 6300b1c), the log stopped at the prompt GEMM's banner and the expert cache still fully resident. Same shape as the prompt collapse in Niko1221#1549, reached by over-reserving the KV instead of by not fitting the experts. Also recorded, both against the figures elsewhere in this file: prompt reading is 178-186 tok/s on a 16,386-token prompt against 74.7 at 2,048 (a short prompt is the worst case), and decode after a 16,386-token prompt is 40.7-42.1 tok/s against 43.6 at 2,048.
…the build Upstream now picks the bound per device at run time (strata::spin_max(queue), sycl/include/strata/sycl_doorbell.hpp): 20,000 reads on a card under the xe driver, 2,000,000 everywhere else, Windows included. intel_gpu_driver() reads /sys/class/drm, so on Windows it returns nothing and every card there gets the long bound whatever the drafting configuration - which leaves the configuration setup writes, --spec 4 --mtp, at 1.18 tokens per round on this card. The launcher is the place that knows the configuration, so it chooses: STRATA_SPIN_MAX=20000 when --mtp is in the arguments and the working adapter's name is a B-series Arc; the default is left alone otherwise. An A-series card keeps the long bound even with --mtp, because there the GPU genuinely waits on the host's per-layer CPU expert work and a short bound made it go on with the experts' outputs missing. A STRATA_SPIN_MAX already in the environment wins, and -DSTRATA_SYCL_SPIN_MAX still fixes one bound for a whole build. This replaces the previous commit's stderr warning, whose premise - a build pinned to 20,000 - upstream's per-device change removed. Measured after the rebase onto that change, so against the run-time choice and not the build option: the same binary, one STRATA_SPIN_MAX apart, Coder IQ1_M, --spec 4 --mtp, 256 greedy tokens after a 2,049-token prompt, medians of three interleaved runs - 44.1 against 24.9 tok/s (+77%), 2.15 against 1.18 tokens per round, and 204 of 217 rounds accepting no draft at all with the long bound against 48 of 120 with the short one. Prefill is unchanged by the bound (72.9 tok/s both ways), so this is the decode path alone. quantize_act_parity and sampler_parity are byte-exact on the rebased tree (0 failures each).
3981f19 to
a99943f
Compare
Niko1221#1549) `auto` filled all but the --vram-reserve-mib reserve, which on a 32 GB card leaves the prompt path with nowhere to go: under WDDM an over-subscribed allocation does not fail, it pages to system memory and crawls. With the prompt path borrowing cache slots (the default with a profile) P5 reserves nothing for it, so the reserve was all there was. Measured on an Arc Pro B70 (32 GB, IQ3_XXS pack, OpenCL backend, eager, driver 32.0.101.8976), 48-token prompt, the same 3,668 gate/up and 3,668 down naive GEMM calls at every cache size: auto -> 17,687 slots (28.82 GiB): 54.05 ms per expert GEMM call, the prompt reads 1,171 s, ~15 GiB of GPU memory held in shared memory 14,688 slots (23.83 GiB, 7.17 GiB kept free): 2.48 ms per call, 12.8 s the cliff is between 14,704 and 15,382 slots; above it the times are noisy but every size up to 14,704 is stable at ~10 s Both sizes give identical greedy output tokens on a scrambled 48-token prompt, so this is a performance cut, not a behaviour change. The Coder config still sizes its whole 12,288 pairs resident, so a model that fits is unaffected. `auto` now keeps about a sixth of what is free for the session's other buffers, and falls back to the old sizing when that would leave no room at all. An explicit --expert-cache is the user's own budget and is not cut.
… HIP (Niko1221#1549) The B70 finding in the previous commit needs an A/B on hardware this has not been measured on, and the repo's rule is that a default change is made where it was measured. So the share of free device memory that `--expert-cache auto` keeps for the session's other buffers is now one helper in include/strata/core/expert_cache.hpp, read by both the CUDA/HIP and the SYCL sizing: STRATA_EXPERT_CACHE_HEADROOM=<fraction> "1/6", "0.17", or "0" for the old sizing The default is a sixth where it has been measured (the SYCL port, Arc Pro B70 32 GB, IQ3_XXS, OpenCL backend, eager, driver 32.0.101.8976) and zero on CUDA and HIP, so every other build's default path is byte-identical to the last release. The measurements and the A/B are in the helper's comment. Verified on the B70: default 14,688 slots (23.83 GiB), 0 -> 17,687 slots (28.82 GiB) exactly as before, 1/4 -> 13,088 slots (21.25 GiB). The same 14,688-slot cut gives identical greedy output tokens to the uncut one on a scrambled prompt, so the headroom costs no accuracy. The CUDA and HIP sources have the change but this machine has no nvcc or hipcc, so those two backends have not been compiled here - the SYCL port builds clean, and the CUDA/HIP lines are the same statement against the same helper.
c0a749f to
f3a882c
Compare
…e card is full (Niko1221#1549) The prompt path's per-expert GEMM cost went from 2.3 ms at a 23.83 GiB expert cache to 54-122 ms at 28.82 GiB on an Arc Pro B70 - same calls, same shapes - and the phase timing charged both expert GEMM phases with it. This times the engine's own Gemm::f16 at the prompt path's shapes and routed row counts, with the card empty and again after allocating and touching 28 GiB of it: gate/up 1280x2560 empty 1467 us/call 28 GiB 1467 us/call (1.00x) down 2560x640 empty 380 us/call 28 GiB 382 us/call (1.00x) So a full card does not slow this kernel, and the 111 ms per expert pair the slow arm charged to those phases was not GEMM work. That is what narrowed Niko1221#1549: the slowdown is spread across most kernels (gdn 55x, the QSA projections 42x, the GDN output projection 53x), not the expert path. Y is sized for the wider of the two outputs - a run sized for gate/up alone overruns it on the down product. vram_pressure [calls=800] [fill_gib=28]; on Windows it needs ZEL_LIBRARY_PATH pointed at a Level Zero loader to start, as the engine does.
…penCL) (Niko1221#1549) The IQ3_S rows above are the Linux VM (xe, Level Zero), where a 4K prompt reads 980-1,002 tok/s with the same mirror design. On Windows with the OpenCL backend the same model falls off a cliff, and the reason is not the mirror: 35% of the pack mirrored over PCIe costs nothing (2.05 ms per expert against 2.01 fully resident), and the slow arm streams fewer experts than the fast one. Same calls, same shapes, identical output tokens - 3,668 gate/up and 3,668 down GEMM calls at every cache size, and sycl/probe/vram_pressure.cpp says a full card does not even slow the GEMM. Measured, Arc Pro B70, Windows, OpenCL, eager, driver 32.0.101.8976, Flash-Next IQ3_XXS, --prefill 4096, greedy, 64-token prompt: 17,687 slots (28.82 GiB, auto) 426 s at 111 ms per expert GEMM pair; 15,382 slots 472 s at 122 ms; 14,704 slots (23.87 GiB) 9.7 s at 2.32 ms; 10,698 slots 10.3 s at 2.39 ms. Every size up to 14,704 is stable at ~10 s and every size from 15,382 up takes minutes. Nothing stops it, because the engine's post-write free-VRAM check cannot fire here: STRATA_DEVICE_FREE_MIB makes dpct's get_memory_info return a setup-time constant - (32 - 1.0) * 1024 = 31,744 MiB - before the allocation and after it alike, and the Linux DRM correction does not run on Windows. Printed on the B70: 'post-write free 31.00 GiB, want 2048 MiB, cache 17687 slots (28.82 GiB)'. On CUDA cudaMemGetInfo is live and the check works, which is why the RTX 5070 Ti A/B found no difference at 1/6. Under STRATA_PREFILL_SYNC=1 the whole prefill is 36x slower with the same work and the slowdown is spread across most kernels rather than the expert path - gdn 55x, the QSA projections 42x, the GDN output projection 53x, the two expert GEMMs 39x. What the GPU is doing with the last 2 GiB is unexplained and the section says so.
…1549) On the OpenCL backend there is no free-VRAM query at all: `ext_intel_free_memory` needs Level Zero sysman, which this driver does not expose, so dpct's `get_memory_info` answers from STRATA_DEVICE_FREE_MIB - a setup-time constant, `(vram_gb - 1.0) * 1024` = 31,744 MiB here, written by setup_intel.py from the registry. It is the same number before and after the weights, the native head and the MTP draft layer are resident, because those are allocated before the expert cache is sized and the constant cannot see them. Two consequences, both measured on an Arc Pro B70: - The engine sized `auto` against the whole card and then allocated on top of it. 28.82 GiB of experts plus the engine's own 6.8 GiB of weights, KV, draft head and buffers is 35.6 GiB against a 31.16 GiB card, so the driver evicted and every kernel that touched an evicted page crawled - `gdn` 55x, the QSA projections 42x, the GDN output projection 53x, the expert GEMMs 39x, the whole prompt 36x. - The engine's post-write free check could never fire, since it compares that constant to a threshold. DXGI has the query this backend lacks. `IDXGIAdapter3::QueryVideoMemoryInfo` gives this process's dedicated VRAM and the OS's budget for it, which already excludes whatever else holds the card (the display compositor was 0.88 GiB of the 7.5 GiB gap; the engine's own footprint was the rest). Both are live. Measured in `sycl/probe/vram_pressure.cpp`, which now prints them: DXGI (before anything): this process uses 0.07 GiB of a 31.04 GiB budget (30.97 GiB left) DXGI (after 28 GiB fill): this process uses 28.07 GiB of a 30.97 GiB budget ( 2.89 GiB left) where dpct's figure gave 0.00 and an error at the same moment. `src/core/expert_cache.cpp`'s `own_gpu_local()` reads them and `device_free_bytes()` prefers that figure over the constant, keeping a 256 MiB margin for what the cache sizing does not price. STRATA_DXGI_FREE=0 takes the old path, which is the A/B. The helper lives on the port's side and is Windows-only; every other platform keeps the sizing it had. With it, `auto` is told 23.84 GiB rather than 31.00 and sizes the Coder 11,348 slots (21.63 GiB) instead of 12,288 (23.42 GiB). That smaller cache is faster on both paths, three interleaved reps each with a 30-second settle, same binary, 2,048-token prompt and 128 greedy tokens: | | slots | prompt | decode | |---|---|---|---| | constant + the 1/6 headroom | 12,288 | 73.8 tok/s | 13.0 tok/s | | the DXGI figure | 11,348 | 113.1 tok/s | 56.3 tok/s | So the headroom's default is now 0 everywhere: it was standing in for the missing query, not for a real requirement. Kept as a pure A/B knob, since a card whose free figure is *not* live (or a user who sets STRATA_DXGI_FREE=0) still wants a way to take a share back. The B70 measurement where 1/6 helped, the cliff table and the 5070 Ti null result are unchanged and still in docs/EXPERT_CACHE_HEADROOM.md.
|
Superseded by #1890, which rebases this work onto the 0.1.43 release's SYCL port and adds the OpenCL expert-GEMM fallback, the eager commit path and the live free-VRAM figure. |
What
Native Windows support for the SYCL port (
setup --backend sycl): no Docker, the engine is built from source withIntel oneAPI (
sycl\tools\build.bat) and run directly (sycl\serve\strata-sycl.bat). Plus the fixes the port needsto compile on Windows at all, and the no-command-graph paths ΓÇö this Windows driver exposes no user-mode Level Zero
adapter, so the OpenCL adapter has no Graph extension.
Replaces #1295: that branch predates 0.1.40.2, and the release merged most of its forward-port (and has its own
STRATA_VERIFY_EAGERwindow replay), so this one carries the Windows deltas only ΓÇö 4 commits.Measured
Arc Pro B70 32 GB, Windows 10, driver 32.0.101.8976, i9-9900 / 64 GB, conda-forge
dpcpp_win-642026.1.1, OpenCLbackend,
STRATA_VERIFY_EAGER=1. Coder IQ1_M, 12,288/12,288 experts resident,--stream-experts, INT8 KV, greedy.Each row is this branch's engine (0.1.40-sycl) against the previous 0.1.39-based build of this port, run
interleaved from the command line (v1, v2, v1, v2, …), which is how the numbers in docs/INTEL.md were taken:
| | 0.1.39-sycl | 0.1.40-sycl (this PR) | Linux / Level Zero / graphs |
|---|---|---|---|
| decode, 256 tokens,
--spec 4 --mtp, 3 pairs | 69.6-69.9 (median 69.7) | 77.2-77.3 (median 77.3) | 78.2 || decode, 128 tokens after a 2,048-token prompt, 2 pairs | 37.6 | 38.2-41.1 (median 39.6) | ΓÇö |
| decode, 256 tokens,
--spec 2(no draft layer), 3 pairs | 53.4-53.9 (median 53.6) | 22.0-22.9 (median 22.3) | 33.6 || prompt reading, 2,048 tokens | 8.0-14.5 tok/s | 71-73 tok/s | 790-1,002 tok/s |
|
quantize_act_parity --selftest,sampler_parity| byte-exact, 0 failures | byte-exact, 0 failures | ΓÇö |Two traps in measuring this backend, both worth passing on. With a 2,048-token prompt the engine is not
run-to-run deterministic: the same binary diverges at token 60 and its state hash differs between two runs. And the
web app answers a
temperature: 0.0request with its own Chat defaults (temperature 1.0, top_p 0.95), so a "greedy"run through it is sampled. The kernel-level self-tests and the CLI's
--greedyare the checks that hold.And short runs are not comparable on this backend. The same build and mode gave 48.9 and
67.5 tok/s on two consecutive 128-token runs, and a 155-token answer through the API contradicted a 256-token run by
25%. My first pass at this table used those short runs and wrongly read as a ~25% regression; interleaved medians over
256 tokens are the numbers above.
The one reproducible regression is
--spec 2ΓÇö 22.3 against 53.6 tok/s, interleaved, same card and flags. It isthe no-draft-layer path, so it is either the drafter-free window or the suffix drafter's cost; not attributed.
STRATA_VERIFY_PROFILE=1puts 24 of a 65 ms 4-token window in the GDN hyper-connection read, nearly flat from T=2(23.6 ms) to T=6 (25.6 ms): a per-layer fixed cost of ~0.5 ms over 48 layers ΓÇö kernel dispatch, not arithmetic.
sycl-lslists onlyopencl:gpuhere, so there is neither the Graph extension nor sysman's free-VRAM query.Two things a fresh Windows build needs, both found by building it clean
The DLLs to start at all.
strata.exestatically importslibmmd.dll(Intel's runtime, besideicx), and theUR adapters load
ur_loader.dll/ur_win_proxy_loader.dllfrom beside the exe. Without those three staged theengine exits
0xC0000135(STATUS_DLL_NOT_FOUND) before printing anything ΓÇö it only ran in my earlier builddirectory by accident.
sycl/CMakeLists.txtnow stages all three with the rest.sycl/tools/build.batfindslibircmt.libthrough the build dir's CMakeCache when the shell has noicxonPATH, so the script works from a plain prompt; otherwise the link fails
LNK1104.Three compile bugs in the port's own code
<sycl/sycl.hpp>reaches the Windows SDK, whose headers defineOUT(minwindef.h),small(rpcndr.h) andnear(windef.h) as macros. An identifier with one of those names is macro-expanded away:
sycl/src/kernels/cuda/verify_kernels.dp.cpp:template <bool OUT>loses its parameter, so every use becomeserror: expected expression. RenamedWRITE_OUT.sycl/src/program/generate.cpp:const int64_t smallbecomesconst int64_t char;memmemis POSIX (nowstd::searchover the same bytes).sycl/src/kernels/sampler_parity.cpp:const bool near.Without command graphs
STRATA_VERIFY_EAGER=1(set bysycl/serve\strata-sycl.bat) replays the window, the commit and the drafter'schains on the queue instead of launching captured graphs:
record_commit()is the commit graph's body;capture_*()record it where there is a graph backend.The
capture_round/capture_step/capture_prefillbodies becamerecord_*()functions the drafterreplays.
prepare_chain()refuses: its per-step events only exist inside a recorded graph.Eager replay refuses a host-served expert plan instead of running it ΓÇö the mid-record waits expire into stale
plans, which is silently wrong rather than slow.
oneMKL SYCL BLAS has no OpenCL Xe2 backend and the failed attempt poisons the queue, so the prompt GEMM is a
plain SYCL kernel there - and the first version of that kernel read the whole weight matrix once per row (a
2,048-row chunk of one 2,560 x 4,096 projection moved 86 GB of weights).
fallback_gemm_tiledblocks a 64 x 64corner of Y per work-group and walks K in 16-wide tiles with the operands in local memory: **8.0-14.5 -> 71-73
tok/s** on a 2,048-token prompt, decode unchanged. Same arithmetic - both kernels accumulate over k in ascending
order with the fma spelled out - and
STRATA_FALLBACK_SELFTEST=1runs both over a problem that crosses every tileedge and prints the bitwise difference: 0 of 4,189 values, largest 0.
sampler_parity's fixture 18 skips its graph half without the extension instead of ending the process.Also
Setup names the Intel Arc on Windows (display adapters + the display-class registry's 64-bit VRAM, which is the
true VRAM ΓÇö WMI's
AdapterRAMstops at 4 GB), and--backend syclhands over tosycl/setup_intel.pythere too.sycl/CMakeLists.txt: IntelLLVM / MSVC flag branches (/fp:preciseis load-bearing ΓÇöquantize_act_paritygoesfrom byte-exact to 303k mismatches without it), installer-free oneMKL, and the runtime DLLs staged beside
strata.exe.GgufExpertSourcereads at an offset throughOVERLAPPED(there is nopreadon Windows).The free-VRAM query needs Level Zero sysman, so setup writes
STRATA_DEVICE_FREE_MIB/_TOTAL_MIBfrom theregistry, and keeps them across a rewrite when detection flakes.
The expert mirror on Windows (a model that does not fit in VRAM)
The mirror's default size came from
/proc/meminfo, which does not exist here, so nothing was ever mirrored onWindows: a model that does not fit entirely in VRAM refused to run ("6889 experts are neither in VRAM nor mirrored").
GlobalMemoryStatusExis the same reading the process stats already take there, so with it IQ3_XXS runs on the B70(17,687 of 24,576 experts in VRAM, the other 6,889 - 11.15 GiB - in pinned host memory the GPU reads over PCIe).
Decode absorbs the misses; the prompt does not, and not at the bandwidth bound (535 GiB over PCIe in 1,988 s is
~2% of the link). Reported with the profile that would settle it as
#1549.
The Monitor tab on this card
GPU load and VRAM come from the counters Windows keeps for every WDDM adapter -
GPU Engine(per process, per engine,summed per card as Task Manager does) and
GPU Adapter Memory(dedicated bytes) - filtered to the LUID of the adapterholding the most dedicated memory, so an iGPU beside the card does not add to the numbers. Measured through the API on
the B70: 2.9% idle, 100% while decoding, 30.01 of 32.00 GiB in use (the 23.4 GiB of experts, the KV and the driver's
own buffers), which agrees with the engine's own "12,288 of 12,288 experts resident".
Temperature, power and the PCIe link have no Windows counter at all, so those three tiles stay empty - now with the
dashboard saying why instead of showing a bare dash.
Two traps in that path, both found by measuring rather than reading, both in the commit message: PDH's wildcard API
answers
PDH_INVALID_ARGUMENTon this machine for every path (core counters included), so instances come fromGet-Counter -ListSetand each gets its own counter; and a byte counter must be asked for asPDH_FMT_LARGE, becausewith
PDH_FMT_DOUBLEPDH fills the union with a double and 4.4 GB comes out as 4,437,051,947 GB.Limits
Needs an admin account for VS Build Tools and the Intel driver; the conda + VS + PyPI/NuGet oneMKL pieces install
per-user.
A model whose experts do not all fit needs a RAM mirror (
STRATA_MIRROR_MIB) or all-resident sizing; IQ3_XXS(17,687/24,576 experts resident here) stalls without one.
Windows-SDK macro names have to be avoided as identifiers in the port.
Tests:
tools/test_setup_intel_windows.py(new),tools/test_setup_sycl.py,tools/test_setup_amd.py,tools/test_setup_choices.pyΓÇö 87 pass, no GPU and no network.