Repository navigation
Conversation
demetree
pushed a commit
to demetree/Strata
that referenced
this pull request
Oct 10, 2026
…, crashes in record_round) Ported the eager-replay refactor the earlier branch carried for mtp.cpp: record_prefill / record_round / record_step extracted from capture_prefill / capture_prefill_dev / capture_round / capture_step, the capture_* guarded with eager_replay(), and the launch sites in prefill() and draft() replaying the body when there is no graph backend. Upstream's cs_ = nullptr (load() creates its own queue) and the dhead_type_ line are both kept. State on the B70: the graph-capture failure is gone - 'window and draft graphs captured in 0 ms' - and record_prefill works (the prompt path completes). record_round then crashes with 0xC0000005 in sycl::platform's copy constructor on a worker thread, reproduced with --spec 2 (so capture_step is not involved) and with --spec 4. Needs a symbolised stack before it goes anywhere near a pull request; PR Niko1221#1890 still describes the --mtp gap as open.
demetree
force-pushed
the
intel-vram-headroom-dxgi
branch
2 times, most recently
from
October 11, 2026 01:06
b8280a9 to
5ab7177
Compare
added 7 commits
October 11, 2026 15:15
… 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.
…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.
…M figure (Niko1221#1549) upstream main now carries the SYCL port (a1c47bc), built and measured on Linux/Level Zero, where oneMKL and command graphs both work. This is the Windows/OpenCL half: the pieces that backend needs and that a Linux build never exercises, plus the expert-cache sizing fix that came out of it. Cherry-picked from the earlier branch (f3a882c, ddb610d): * STRATA_EXPERT_CACHE_HEADROOM=<fraction> - the opt-in A/B knob, default 0 everywhere * the live free-VRAM figure: src/core/expert_cache.cpp's own_gpu_local() asks DXGI for this process's dedicated VRAM and the OS's budget for it, and device_free_bytes() prefers it over STRATA_DEVICE_FREE_MIB Then the Windows portability layer, one file at a time, because the upstream migration was done on Linux and several of its copies only build there: * pinned.dp.cpp: restore the STRATA_USE_CUDA guard around the cudaDeviceProp block * verify_kernels.dp.cpp: rename the template parameter OUT -> WRITE_OUT (a Windows SDK macro) * gguf_expert_source.cpp: Windows file I/O - the #ifdef _WIN32 guard and the portable_open / portable_pread / portable_close helpers, on top of upstream's new chunked mirror * generate.cpp: rename the `small` identifier (also an SDK macro) and replace memmem with std::search. The identifier only - the earlier fixups.py did a global replace that also rewrote comments and strings ("A small header binds" -> "A small_chunk header binds"), which is a bug in that tool and is not repeated here * CMakeLists.txt: compiler-ID gating (/fp:precise and /W3 for MSVC, -fp-model=precise elsewhere), STRATA_SYCL_JIT_STATELESS (the IGC >4 GiB surface-offset bug, measured on an A770), the oneMKL stub for an installer-free conda install, and the DLLs that must sit beside strata.exe * gemm.dp.cpp: keep upstream's staged_gemm (the oneMKL >4 GiB fix) AND our plain-SYCL fallback, gating the oneMKL attempt on onemkl_usable() so an OpenCL run goes straight to the fallback. Upstream has no fallback at all: on Linux/Level Zero oneMKL works, so there is nothing to fall back to. On OpenCL oneMKL is not usable, so without this the prompt path has no expert GEMM and the engine dies on the first request with "oneMKL ... does not support this API" (status 999) * verify.cpp: restore the STRATA_VERIFY_EAGER guard on capture_commit - without it an OpenCL run reaches begin_recording and dies on "opencl backend is not supported by SYCL Graph extension" Verified on the B70: the engine sizes its cache from the live figure (23.84 GiB read where the constant gave 31.00), picks 11,348 slots / 21.63 GiB - exactly the pre-rebase measurement - fills the profile mirror, and runs the prompt path through the plain-SYCL fallback with oneMKL bypassed. Not done: the MTP drafter still captures command graphs unconditionally (src/core/mtp.cpp's capture_prefill / capture_prefill_dev / capture_round / capture_step). The earlier branch refactored those into record_* bodies with replaying capture_* wrappers under STRATA_VERIFY_EAGER; that refactor is not in this commit, so a run with --mtp still fails at "mtp: begin capture" on OpenCL. A run without --mtp reaches the same class of refusal from the mirror sizing and needs its own look.
generate.cpp's post-load check read dpct's get_memory_info() directly, which on this backend reports a figure that moves with nothing. So it could print "31744 MiB of VRAM free with everything loaded" and pass while the card was in fact full - which is how a --max-context 262144 config filled the card to zero and then died in queue creation (0xC0000005 in sycl::platform's copy constructor on a worker thread, i.e. an allocation that cannot be satisfied). own_gpu_local() moves out of the function's static scope so the header can declare it. This removes the disagreement between the two figures but does NOT by itself stop the crash: the expert cache is sized before the session's KV is allocated, so the live figure read at sizing time cannot see the KV a long-context config goes on to reserve. That ordering is the remaining problem and is not addressed here. Measured on the B70, Coder IQ1_M, 262,144-context config, --spec 4 --mtp: builds clean, and the serve-side line now reports a figure that tracks the card instead of the constant. The crash still occurs; with --vram-reserve-mib 4096 the server does come up, but that is the workaround the engine already suggests, not a fix.
generate.cpp:4315 read dpct's get_memory_info() directly while the auto branch three screens up calls strata::core::device_free_bytes(), so on this backend one path sized from the live figure and the other from a figure that moves with nothing. Third of the three call sites fixed alongside the serve-side check (5ab7177). The crash this was chasing is NOT fixed: with --max-context 131072 --kv-resident 131072 --spec 4 --mtp the engine still dies with 0xC0000005 in sycl::platform's copy constructor on a worker thread after 'window and draft graphs captured in 0 ms', on the same command that completes on a build with fprintf markers inside record_round. The figures now agree, which was the point of this change, but the underlying failure remains undiagnosed and is the next thing to run down.
…ng a null graph (Niko1221#1549) On the OpenCL backend there is no SYCL Graph extension, so STRATA_VERIFY_EAGER runs the verify window and the commit's kernel sequence directly on the queue. run() already did that for the window (record_window), but capture_commit() answered "captured, nothing to do" for the commit while commit() went on to launch the graph it had been told would be there: verify.hpp:405 command_graph_exec_ptr commit_exec_ = nullptr; never assigned verify.cpp:1842 if (STRATA_VERIFY_EAGER) return true; "captured" verify.cpp commit(): (cs_)->ext_oneapi_graph(*commit_exec_) a null graph commit() died with 0xC0000005 inside itself on a worker thread - the innermost frame is Verifier::commit in verify.cpp.obj at +1268 bytes, resolved from a /MAP of the same binary. The comment above the launch already said the replay's outputs are never read; nothing replayed them. capture_commit's body is now Verifier::record_commit, and the two callers are: capture_commit() begin_recording -> record_commit -> end_recording -> finalize commit() STRATA_VERIFY_EAGER ? record_commit() : ext_oneapi_graph(*commit_exec_) The body refreshes everything it reads - the mapped commit_ upload first, then the layer snapshots - so a replay sees the same inputs the graph would have, and it is safe to run once per window. h_commit_ is written by commit() before the launch, as before. What this leaves alone: the default path records the same kernels in the same order into the same graph and finalizes it the same way - only the body moved into a function the eager path can call. The window's replay (record_window), the MTP drafter's replays (record_round/record_step/record_prefill) and the one-token self-commit are unchanged. Measured on the Arc Pro B70, 32 GB, Windows 11 / OpenCL, engine 0.1.43 + the Windows port, coder-IQ1_M, 2047-token prompt, --spec 4 --spec-min-p 0.5, STRATA_VERIFY_DEVICE_PLAN=1 STRATA_VERIFY_NO_HOST=1 STRATA_VERIFY_EAGER=1 STRATA_MIRROR_MIB=4096, greedy: before exit -1073741819 (0xC0000005) after "expert cache 12288 slots, 23.42 GiB" after exit 0; prefill 2047 tokens in 25252 ms -> 81.1 tok/s decode 128 tokens in 1921 ms -> 66.6 tok/s, speculation 3.91 tokens per round (33 rounds of 6, 96 of 96 drafts accepted) Built for SYCL on Windows. I have no CUDA or HIP toolchain on this machine, so those two backends are not built or run by me - the file touches both but the change is inside a constant that is false unless STRATA_VERIFY_EAGER is set, which only the SYCL port's launcher does.
…1549) The comment on own_gpu_local() still described the port's older two-source behaviour: get_memory_info() answering from STRATA_DEVICE_FREE_MIB, a setup-time constant. That override is gone (git grep finds the variable nowhere in the engine now), and what the backend actually answers without the DXGI figure is 0 - which sizes the cache to nothing rather than to a constant. Comment only: no code changed, so the binary and every measurement are untouched.
demetree
force-pushed
the
intel-vram-headroom-dxgi
branch
from
October 12, 2026 11:45
cb003e0 to
d2a7dda
Compare
This was referenced Oct 12, 2026
added 2 commits
October 12, 2026 19:32
Two things that came out of measuring the port on an Arc Pro B70 on Windows. **1. `strata-sycl.bat` gains `STRATA_SYCL_LEVEL_ZERO=1`, off by default.** The port assumed the OpenCL backend because `sycl-ls` never listed a Level Zero device on this card. That turned out to be a graphics-driver version, not the card: the Intel Arc **Pro** line's newest (32.0.101.8976) does not expose Level Zero through any SYCL adapter at all, while the **generic** Arc line's 32.0.101.9034 does. With it, `ext_intel_free_memory` and `ext_oneapi_graph` both appear, which is exactly what the OpenCL workarounds existed for. The settings it turns on, all measured on that card, 2047-token prompt, coder-IQ1_M: | | OpenCL (eager) | LZ v2 | LZ v1 | |---|---|---|---| | decode, 128 tokens | 66.6 tok/s | 87.9 tok/s | **108.8 tok/s** | | prefill, 2047 tokens | 81.1 tok/s | 64.5 tok/s | 81.1 tok/s | | verify graph capture | skipped | 9,495 ms | **276 ms** | `SYCL_UR_USE_LEVEL_ZERO_V2=0` is the v1 adapter: +24% decode over v2 and 34x faster graph capture. Niko1221#1940 pins v2 on for graph replay (intel/llvm#23073); on this card that trade is the wrong one. It stays opt-in and off by default because it is not reliable yet: through the serve request path the device-plan verify window occasionally never gets its doorbell and the engine stalls ("waiting for the GPU to finish the window; the GPU rang 0"). The engine itself is stable in the CLI and through the launcher (512 decode tokens at 74.5 tok/s, prefill 121.4 tok/s against OpenCL's 81.1), so the trigger is narrower than "Level Zero is broken", but it is not fixed, and a default that stalls after a few requests is worse than a default that is merely slower. **2. `LEVEL_ZERO_FINDINGS.md` records how that was established**, including the parts that turned out to be wrong. The two most useful entries for anyone reading this PR: - The Level Zero **driver** on this card was never broken. Walking the Level Zero API directly, with no SYCL or UR in the way, gives a driver handle, a device handle and full device properties (`TYPE_GPU`, `0x8086`/`0xe223`, 31.2 GiB). The failure was entirely in the loader and adapter above it. - `ur_adapter_level_zero.dll` imports `UMF.dll` and the `VCRUNTIME140*` set, and `sycl/CMakeLists.txt` stages neither beside `strata.exe`. The adapter therefore fails to load, the UR loader reports only OpenCL, and `ONEAPI_DEVICE_SELECTOR=level_zero:0` kills the process with `0xC0000409` before `main` prints anything. That is a real gap in the copy list, and the silence is what makes it expensive: every Level Zero adapter failure looks like "there is no Level Zero here". **This is the piece of the PR I would most want reviewed.** ## What this leaves alone The default path is untouched — `STRATA_SYCL_LEVEL_ZERO` unset is exactly the previous behaviour, including `ONEAPI_DEVICE_SELECTOR`, `SYCL_UR_USE_LEVEL_ZERO_V2` and `ZES_ENABLE_SYSMAN` staying unset. Every variable is only written when it is not already defined, so an explicit `ONEAPI_DEVICE_SELECTOR` for a two-card split or an OpenCL run still wins. No engine source changed. Built and run on SYCL/Windows only; I have no CUDA, HIP or Linux toolchain here.
…rify window
With Level Zero selected, a verify window run as a replayed SYCL command graph never raises the
doorbell flag the host polls for: the engine stalls reporting that the GPU "rang 0" with all three
flags still zero, at KV position ~84 - about 20 windows into a request - and it reproduces every
time through the serve request path.
The same workload with the eager replay (no graph, the window's body re-run on the queue) is stable:
35 requests and three 900-token generations, 995 verify windows, no stall. So the doorbell write
happens inside the replayed graph and is either not executed or not visible to the host's polling
read on Level Zero.
Ruled out by direct measurement before that: the device plan being off (which makes decode 0.95
tok/s), the launcher's environment, the driver-search variable, the v1/v2 adapter choice, prompt
length, and --mtp.
strata-sycl.bat therefore sets STRATA_VERIFY_EAGER=1 whenever it selects Level Zero. An explicit
STRATA_VERIFY_EAGER=0 in the environment is the only way back to the graph path, and that is now a
debug setting.
Measured, Arc Pro B70, driver 32.0.101.9034, coder-IQ1_M, 2047-token prompt, greedy, MTP on:
LZ + graphs decode 74.5 tok/s (CLI) stalls after ~20 windows
LZ + eager decode 63.3 tok/s (CLI) 995 windows, no stall
OpenCL decode 66.6 tok/s (CLI)
LZ prefill 120.6 tok/s against OpenCL's 81.1 (+49%)
That last row is with oneMKL refusing in all three cases, so the prefill win is the plain-SYCL
fallback GEMM itself running faster on Level Zero. The Level Zero win is prefill, not decode.
LEVEL_ZERO_FINDINGS.md has every number, including the ones that were wrong, and the stall-dump
analysis that turned out to be a dead end: the engine's stall dump carries no stack memory, so its
31 threads cannot be resolved beyond "every one of them is in ntdll.dll, blocked in a wait". A stall
dump with MiniDumpWithDataSegs would have answered this in one step, which is worth fixing in
whatever writes these.
demetree
force-pushed
the
intel-vram-headroom-dxgi
branch
from
October 13, 2026 00:05
60fb14e to
7e3ab52
Compare
…c Arc line
Two things in LEVEL_ZERO_FINDINGS.md.
**1. A dedicated section for the driver decision, near the top rather than buried in the
chronology.** The generic Arc line is a prerequisite, not a detail, so it now says what it is, what
the evidence is, what it costs and how to go back:
- the Arc Pro line's newest build, 32.0.101.8976, does not expose Level Zero through any SYCL
adapter, and it is the newest build that line has - so there was nothing to update to
- what differs between the two builds is the Level Zero ICD (ze_intel_gpu64.dll) the driver
ships: on 8976 the Level Zero driver answers every direct API call correctly (1 driver,
TYPE_GPU, 0x8086/0xe223, 31.2 GiB) and no SYCL adapter can drive it; on 9034 they can
- what it costs: ISV certification and the Pro control panel, both irrelevant for inference
- what it buys: 120.6 tok/s prefill against OpenCL's 81.1
- how to check which line is installed, and that reverting is a normal Arc Pro install whose
only consequence is that Level Zero stops enumerating and the launcher falls back to OpenCL
**2. The TL;DR and the surrounding sections were stale after the stall was root-caused**, and said
the two things that had since changed:
- "Level Zero is not yet safe as a default" and "until that is fixed" - it is fixed, and the cause
is command graphs, so they are now off
- "STRATA_VERIFY_EAGER must be UNSET" on Level Zero - the opposite is true, and it is the fix
- "what this retires, once the stall is fixed" retired STRATA_VERIFY_EAGER along with the other
two - it does not retire, because the graphs it avoids are what stalls the window
Also added the exact two commands that answer "is this machine in the working state", since the
answer was previously spread across four sections.
Documentation only: no code changed.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What this is
mainnow carries the SYCL port, built and measured on Linux with Level Zero, where oneMKL and command graphs both work. This is the Windows half — the pieces Windows needs that a Linux build never exercises, plus the expert-cache sizing fix that came out of building it.Rebased onto
7695859a(the 0.1.43 release); the first version of this PR was againsta1c47bc9. All measurements below are from the rebased tree.Single card, all measurements: Arc Pro B70, 32 GB, Windows 11 24H2, driver 32.0.101.8976 (Arc Pro line) / 32.0.101.9034 (generic), OpenCL + eager replay, engine 0.1.43.
The eight commits
223ee707fe7aaa449f1b6a579188d55199664a1d6cdf81f5d2a7ddac1baf4c95What it changes
1. The Windows build layer —
sycl/CMakeLists.txt,sycl/tools/build.bat,sycl/tools/onemath_config.hpp,sycl/serve/strata-sycl.bat. Compiler-ID gating (/fp:preciseand/W3for MSVC,-fp-model=preciseelsewhere — the precise model is load-bearing:quantize_act_paritygoes from byte-exact to 303k mismatches without it),STRATA_SYCL_JIT_STATELESS(the IGC >4 GiB surface-offset bug), the oneMKL stub for an installer-free conda install, and the DLLs besidestrata.exe.2. Portability fixes in the migrated copies — the
STRATA_USE_CUDAguard inpinned.dp.cpp,OUT→WRITE_OUTinverify_kernels.dp.cpp, Windows file I/O ingguf_expert_source.cpp, thesmall/memmemfixes ingenerate.cpp, and the plain-SYCL fallback GEMM ingemm.dp.cppgated ononemkl_usable().3. The live free-VRAM figure — on OpenCL there is no free-VRAM query at all (
ext_intel_free_memoryneeds Level Zero sysman), soget_memory_infoanswers 0, which sizes the cache to nothing.own_gpu_local()asks DXGI. Same binary, free card, 2047-token prompt: 0 slots / refused → 12,288 slots / 23.42 GiB, prompt 81.1 tok/s, decode 66.6 tok/s.4.
6cdf81f5, submitted separately as #2106 — completesSTRATA_VERIFY_EAGERfor the commit path.main's eager support covers the window (record_window,run()'s guard) but not the commit;mtp.cpp'sprepare_chaindocuments the gap out loud.5.
1baf4c95, the Level Zero diagnosis — opt-inSTRATA_SYCL_LEVEL_ZERO=1in the launcher, off by default, plusLEVEL_ZERO_FINDINGS.mdrecording how it was established. Off by default because the serve request path stalls intermittently; when it runs it is 108.8 tok/s decode and 81.1 prefill against OpenCL's 66.6 / 81.1, with the v1 Level Zero adapter beating v2 by +24% decode and 34x faster graph capture.The one thing I would most want reviewed
ur_adapter_level_zero.dllimportsUMF.dlland theVCRUNTIME140*set, andsycl/CMakeLists.txtstages neither besidestrata.exe. The adapter fails to load, the UR loader reports only OpenCL, andONEAPI_DEVICE_SELECTOR=level_zero:0kills the process with0xC0000409beforemainprints anything. The silence is the expensive part: every Level Zero adapter failure looks like "there is no Level Zero here", which is what sent me down the OpenCL path in the first place and cost days. The workaround in this PR isLibrary\binfirst onPATH; the proper fix is to extend the copy list.Second, and related: the v1 Level Zero adapter beats v2 decisively on this card (+24% decode, graph capture 276 ms vs 9,495 ms). #1940 pins v2 on for graph replay (intel/llvm#23073); on a B70 with this workload that trade is the wrong one.
What this leaves alone
Every new feature is opt-in and the default path is byte-identical unless a variable is set:
STRATA_EXPERT_CACHE_HEADROOMdefaults to 0 everywhere,STRATA_SYCL_LEVEL_ZEROis unset by default, andSTRATA_VERIFY_EAGERis only read where it already was. Every launcher variable is only written when not already defined, so an explicitONEAPI_DEVICE_SELECTORfor a two-card split still wins. No engine source changed in1baf4c95.What I could not verify — please read before reviewing
9f1b6a57(pinned.dp.cpp's unguarded CUDA calls, theOUTmacro inverify_kernels.dp.cpp). So it cannot be built standalone here.gemm.dp.cpp:1031), the likely unsupported type; that is a hypothesis, not a measurement. So the prompt path is on the plain-SYCL fallback and the prefill figures are not oneMKL-backed.LEVEL_ZERO_FINDINGS.md.Suggested review order
gemm.dp.cpp— the plain-SYCL fallback is what makes prompts run at all here, and it is the biggest piece.verify.cpp— thecapture_commiteager guard (sycl verify: complete STRATA_VERIFY_EAGER for the commit path #2106 is the same change standalone).expert_cache.cpp— the DXGI figure, since it changes sizing behaviour on every SYCL backend.CMakeLists.txt— the copy list above, and the compiler-ID gating.I would rather split this:
223ee707+d2a7ddacare one coherent knob,6cdf81f5is #2106, and1baf4c95is really a bug report about the copy list with an opt-in switch attached.