Pull requests / #2433
#2433 SYCL: a pipelined window still on the GPU is polled again, not failed
open · @deriknel · 0 Kommentare · Auf GitHub
BenchmarksServer & APIMulti-GPUNVIDIA / CUDAModels & quantsWindows
Beschreibung
## Title Issue: none filed. Found running `--pipeline-windows 2` on two Arc Pro B70s. ## Summary On the SYCL build, `--pipeline-windows 2` fails the first request with `verify: error`, and the engine then aborts: ``` strata serve: verify: error terminate called after throwing an instance of 'sycl::_V1::exception' what(): level_zero backend failed with error: 37 (UR_RESULT_ERROR_UNINITIALIZED) ``` `Verifier::done()` still has dpct's placeholder for CUDA's error check. It treats every query result except "complete" as an error, and `dpct::sycl_event_query` returns 1 for a command that is still running. So the first window that has not finished when the loop polls it fails the request, and tearing the window down while the GPU still runs it aborts the engine. This change ports the CUDA branch: a window still on the GPU is polled again. ## What changed `sycl/src/core/verify.cpp`, `Verifier::done()` only: - A query of 1 (still running) returns false without setting `err`, as the `cudaErrorNotReady` branch does in `src/core/verify.cpp:3553-3571`. - While it waits, the queue is flushed every 2 ms (`ext_oneapi_empty()` in place of `cudaStreamQuery`), and the 20 s timeout from #267 stays. - The DPCT1000/1001/1009 placeholder comments go. A SYCL error arrives as an exception, which the function's existing `catch` handles. Every caller of `done()` (`sycl/src/program/generate.cpp:7470, 10132, 10153, 10779, 11003`) polls a window started with `pl_launch`, so the serial loop does not change. ## Extra Notes - Reproduced on 0.1.43.1: two starts with `--pipeline-windows 2`, and both failed on the first request as above. `done()` is the same in 0.1.44.1 (296b59d2). - With this change alone the pipelined loop runs, but its greedy text forks from the serial loop's within a few dozen tokens. The fix for that is a SYCL port of c81f9984, which I will send as a separate PR. - Tested with both changes on 0.1.44.1 (296b59d2): 2x Arc Pro B70, Swift 1.5 Flash-Next IQ3_S, layer split 26, every expert in VRAM, `--spec 4 --spec-min-p 0.6`, int8 KV. - Greedy text identical to the serial loop on 5 prompts (prose, code, maths, an edit, a story), in each of 3 comparisons. - Decode 79.6 / 113.8 / 89.0 tok/s (prose / code / a long answer), against 75.0 / 99.3 / 79.1 with the pipeline off. - A ~148K-token prompt (answer found, 78 tok/s decode) followed by 20 prompts of ~13K tokens, 4 runs on a fresh engine, no errors; a 109-request soak with no failures. - Built and tested the SYCL target only; the change is in `sycl/`. We investigated (myself and Claude Opus 5.5 in Claude Code) and I reviewed the findings and ran every test on my two Arc Pro B70s on my CachyOS machine.
Mehr auf der Site
Links zu Install, Modellen, Releases.