sycl prefill: DMA an expert straight from the arena only when it is USM host memory (A750 prompt deadlock)

SYCL has no cudaHostRegister, so the expert arena is pageable memory even where the engine calls it pinned. A copy->memcpy of a 2.3 MB blob
from it goes through the Level Zero driver's staging path on the calling thread, which never returned on some prompts (A750, default
flags, a 2,744-token docs excerpt: the thread spun in sched_yield inside libze_intel_gpu with every engine idle, 6 of 6 runs; also seen
in a 5.3K-token needle prompt of the task suite). Only USM host allocations (the --stream-experts pinned mirror) are DMA'd directly now
(answer cached per 2 MiB bucket); every other expert goes through the stager's pinned buffers like an unpinned one.
STRATA_PF_DMA_PAGEABLE=1 restores the old direct copy.
Checks: the repro completes 3 of 3 (0 of 6 before); identical first-token logits and tokens, old vs new: A750 (frozen share, 1.4K, 2.5K and
4K prompts), B70 stock config and B70 without --stream-experts (there the new path is also faster: prefill 680-740 vs 77-456 tok/s); A750 4K
prefill 79.4 vs 79.3 tok/s, B70 stock 1246 vs 1240.
Adds sycl/tools/test_prompt_deadlock.sh with the 2,744-token prompt (passes on the fix, times out on the old build).

sycl/src/prefill/prefill.cpp

@@ -51,6 +51,7 @@
 #include <memory>
 #include <mutex>
 #include <thread>
+#include <unordered_map>
 #include <vector>
 
 #ifndef STRATA_PREFILL_MMQ
@@ -97,6 +98,32 @@ namespace strata::prefill {
 namespace {
 
 using Clock = std::chrono::steady_clock;
+
+// 0.1.48, A750 deadlock: SYCL has no cudaHostRegister, so the expert arena is ordinary pageable memory even where the engine
+// calls it "pinned" (pinned.dp.cpp registers nothing).  A DMA straight from it, `copy->memcpy(stage, arena_ptr, 2.3 MB)`, makes the
+// Level Zero driver copy through its own staging buffers on the calling thread, and on some prompts that wait never ends (the
+// thread spins in sched_yield inside libze_intel_gpu with the GPU idle; tools: /data/148-intel notes).  Only a USM host
+// allocation (the --stream-experts pinned mirror, the B70's) is DMA'd directly; every other expert is copied into one of the
+// stager's pinned buffers by its threads first, the path unpinned experts already take.  STRATA_PF_DMA_PAGEABLE=1 restores the
+// direct copy from pageable memory.
+inline bool dma_direct(const core::ExpertSource* src, int64_t layer, int64_t expert, const void* blob) {
+    if (!src->pinned(layer, expert)) return false;
+    static const bool pageable_ok = [] { const char* v = std::getenv("STRATA_PF_DMA_PAGEABLE"); return v && v[0] == '1'; }();
+    if (pageable_ok) return true;
+    if (blob == nullptr) return false;
+    // the answer per 2 MiB bucket (a query costs ~2 us, 10k+ experts a chunk): "host USM" only when both ends of the bucket are, so a bucket
+    // that straddles two allocations is never taken for USM
+    static thread_local std::unordered_map<uintptr_t, bool> usm_bucket;
+    const uintptr_t key = (uintptr_t) blob >> 21;
+    const auto it = usm_bucket.find(key);
+    if (it != usm_bucket.end()) return it->second;
+    const auto ctx = dpct::get_in_order_queue().get_context();
+    const bool usm = sycl::get_pointer_type((const void*) (key << 21), ctx) == sycl::usm::alloc::host &&
+                     sycl::get_pointer_type((const void*) (((key + 1) << 21) - 1), ctx) == sycl::usm::alloc::host &&
+                     sycl::get_pointer_type(blob, ctx) == sycl::usm::alloc::host;
+    usm_bucket.emplace(key, usm);
+    return usm;
+}
 constexpr float EPS = 1e-6f;
 constexpr int64_t N = 2560, HC = 4, D = N * HC, LR = 320, K = 10, NE = 512;
 constexpr int64_t C = 10240, ZV = 6144, HV = 48;
@@ -2550,7 +2577,7 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                                 continue;
                             }
                         }
-                        if (!m.src->pinned(l, e)) {
+                        if (!dma_direct(m.src, l, e, b)) {
                             job = (int) js.size();
                             js.push_back({b, (size_t) lay0.blob_bytes(l)});
                         }
@@ -3809,7 +3836,7 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                             for (size_t j = 0; j < order.size(); ++j) {
                                 const int32_t e = order[j];
                                 if (m.host_res && m.cache && m.host_res[(size_t) l * m.g->n_expert + e] >= 0) continue;
-                                if (m.src->pinned(l, e)) continue;
+                                { const uint8_t* pb = m.src->pinned(l, e) ? m.src->blob(l, e) : nullptr; if (dma_direct(m.src, l, e, pb)) continue; }
                                 job_of[j] = (int) js.size();
                                 if (gsrc != nullptr) {   // SYCL port: --stream-experts - the stager's thread reads it from the GGUF
                                 js.push_back({nullptr, (size_t) lay.blob_bytes(l), nullptr, (int32_t) l, e, gsrc});
@@ -3834,7 +3861,8 @@ bool Prefill::run_impl(const int64_t *tokens, int64_t n, int64_t pos0,
                             const int sl = stage_next;
                             stage_next = (stage_next + 1) % stage_slots();
                             const auto th = Clock::now();
-                            const bool pinned = m.src->pinned(l, e);   // pinned: never transient
+                            const uint8_t* pb0 = m.src->pinned(l, e) ? m.src->blob(l, e) : nullptr;
+                            const bool pinned = dma_direct(m.src, l, e, pb0);   // pinned: never transient
                             const uint8_t* b = pinned ? m.src->blob(l, e) : nullptr;
                             if (pinned && !b) { err = "prefill: expert source has no blob"; return false; }
                             if (pinned) {

sycl/tools/data/prompt_deadlock_2744.tok

@@ -0,0 +1 @@
+248045,846,198,8917,5437,537,279,2614,9417,47801,303,4097,16916,3387,13,271,16,21,42,11,220,19,17,13,16,19,1411,220,17,21,13,18,15,9537,506,220,18,17,42,553,357,4220,9640,25,279,1788,7620,440,1132,411,3289,5324,11,264,7478,4560,791,198,2835,11,220,16,16,13139,303,69905,1906,1238,264,41784,799,11,279,4560,579,1510,39242,20759,7561,1396,9640,1301,303,2400,318,2083,66420,198,29567,553,8618,4560,15339,264,41784,220,19,42,9640,1118,11,1521,279,1118,9640,1238,264,1151,33094,777,318,263,1782,1118,198,23693,12389,279,1788,13139,6348,220,24,20,4,26126,314,498,16,13,15,15,19,11,220,16,13,15,18,17,60,321,498,15,13,24,23,16,11,220,16,13,15,16,20,14545,279,49219,1732,1754,975,2957,198,63,14,15318,5285,76,9586,63,6,82,2081,1395,4871,321,2050,5331,7611,30494,599,14217,1990,364,220,18,274,11,321,264,6505,303,864,2892,9640,1990,198,3923,2081,11567,466,2050,55557,557,11878,321,16384,1495,318,18,314,279,220,21,24,13139,11,660,220,18,17,42,440,279,3289,383,26,8341,303,11,220,18,17,42,198,70,1821,220,16,13,15,17,15,321,220,16,13,15,16,20,553,328,9333,220,20,1,369,279,11292,7101,314,1510,759,27270,42247,46,20668,63,3644,220,20,318,4058,1510,6457,4337,76465,49461,11481,1102,28,17,15,19,23,63,1174,198,1,25902,1,369,279,4560,1973,1070,318,2720,220,19,579,38171,16,16,20,16,16184,4565,553,8261,513,1711,5212,26,279,11013,369,279,50717,2977,314,198,1719,220,16,16,974,593,383,39875,440,264,220,24,20,4,26062,9581,26,328,69,2225,1,14159,279,13139,303,864,279,3289,557,10281,25,271,91,10885,735,6560,735,974,1411,383,735,974,593,383,735,10281,735,198,91,4277,25,91,4277,91,4277,25,91,4277,25,91,4277,25,91,198,91,220,19,11,15,24,21,735,3644,220,20,735,220,18,13,22,15,18,1411,220,18,13,21,19,16,274,735,220,16,13,15,16,23,498,16,13,15,16,21,11,220,16,13,15,17,15,60,735,220,16,16,314,220,16,16,735,198,91,220,19,11,15,24,21,735,16184,735,220,20,13,24,16,19,1411,220,20,13,23,17,24,274,735,220,16,13,15,16,19,498,16,13,15,16,15,11,220,16,13,15,16,22,60,735,220,16,15,314,220,16,16,735,198,91,220,18,17,11,22,21,23,735,3644,220,20,735,220,17,21,13,16,20,22,1411,220,17,20,13,20,23,24,274,735,220,16,13,15,17,16,498,16,13,15,17,15,11,220,16,13,15,17,18,60,735,220,16,16,314,220,16,16,735,198,91,220,18,17,11,22,21,23,735,16184,735,220,19,19,13,15,23,17,1411,220,19,18,13,19,22,24,274,735,220,16,13,15,16,20,498,16,13,15,16,17,11,220,16,13,15,16,22,60,735,220,16,16,314,220,16,16,735,198,91,220,16,17,23,11,15,15,15,735,3644,220,20,735,220,16,15,22,13,16,19,24,1411,220,16,15,19,13,23,21,21,274,735,220,16,13,15,17,16,498,16,13,15,17,15,11,220,16,13,15,17,18,60,735,220,16,16,314,220,16,16,735,198,91,220,16,17,23,11,15,15,15,735,16184,735,220,16,24,18,13,19,18,24,1411,220,16,24,16,13,15,20,20,274,735,220,16,13,15,16,19,498,16,13,15,16,17,11,220,16,13,15,16,21,60,735,220,16,16,314,220,16,16,735,271,4272,883,220,17,4,383,1844,314,3644,220,20,321,883,220,16,13,20,4,383,279,16184,13,2332,220,18,21,81787,9141,5777,6738,13224,279,6000,12041,6805,883,198,15,13,20,21,274,364,264,220,18,17,42,9640,26,220,15,13,20,22,274,318,2720,220,20,8,321,220,15,13,21,15,274,318,25902,8,998,16384,13,271,5170,33686,65427,2272,279,1788,1542,1973,279,3289,11,220,16,17,49219,314,220,16,13,20,42,12,17,19,42,10885,25,75862,20819,48,19,6645,50,220,15,13,15,15,21,22,2977,318,15,13,15,19,19,1866,8,383,279,198,25902,321,220,15,13,15,15,18,24,318,15,13,15,17,22,8,383,1844,314,3644,220,20,26,75862,27325,19,9802,6645,43,220,15,13,15,15,21,18,318,15,13,15,19,18,8,321,220,15,13,15,15,19,15,318,15,13,15,16,19,1174,279,1788,1844,3817,303,660,314,198,47092,318,565,2060,279,220,16,17,49219,279,1118,3817,4829,1132,1330,12103,2663,553,10790,220,20,579,30147,7249,513,220,15,13,15,15,21,23,318,59865,19,6645,50,8,321,198,15,13,15,15,18,16,318,48,19,9802,6645,43,8,2272,279,16184,11,264,2086,2222,421,369,9711,364,5281,321,524,430,264,3798,364,411,799,13,8020,279,198,3765,7658,3817,557,7463,25,264,1248,9191,321,70665,477,998,524,16384,13,271,10743,321,1542,1439,11988,303,279,77528,1870,6025,25,271,71093,898,198,6031,706,1137,5653,1870,12,2082,1137,5480,17603,1824,16884,29542,290,17603,1824,16884,6860,12756,198,65590,1137,1877,43774,1870,12,2082,471,49,359,2082,1824,16884,8102,25148,290,91,2670,12756,20973,1137,2911,10020,2149,9094,198,1891,585,303,220,16,220,17,220,18,26,635,12619,5653,12,2082,14,2082,1824,16884,29542,290,1137,26157,13279,26,12619,5653,12,2082,14,2082,1824,16884,29542,290,1137,26157,26,2725,198,71093,271,4699,27502,1542,6473,279,21966,38372,21853,314,7840,8213,1238,1330,7948,8339,11,279,1788,69777,2763,1528,1396,1542,25,198,63,312,26157,63,411,1752,11,1510,312,26157,13279,63,38171,16,16,20,16,579,1581,61369,27718,6457,4337,2559,30237,41832,28,17,7561,1510,6457,4337,2559,30237,8827,56,28,16,57208,561,3289,369,1301,198,12802,791,1817,11,748,279,1330,513,8387,8213,13,1049,2942,799,6000,579,72632,321,7230,11,524,264,4220,9640,13,271,550,4744,321,13092,6343,271,48887,264,7021,1558,3641,321,1141,380,4100,15135,1608,279,6149,7141,13,198,2381,2441,895,9242,35494,18,62,5955,50,11,1732,498,1719,11135,23967,Line truncated

sycl/tools/test_prompt_deadlock.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+# Regression test for the Arc A750 prompt deadlock (0.1.48): a 2,744-token prompt (a docs excerpt in chat framing, token ids in
+# data/prompt_deadlock_2744.tok) made the prompt path hang in the Level Zero driver (the thread spun in sched_yield inside
+# libze_intel_gpu while the GPU was idle) when the expert arena was DMA'd straight from pageable memory.  The engine now copies
+# every expert that is not USM host memory through the stager's pinned buffers.
+#
+#   STRATA_BIN=build-jit/strata MODEL_ARGS="--pack ... --native ... --ple-gguf ... --mtp ... --kv int8 --max-context 8192 ..." \
+#       sycl/tools/test_prompt_deadlock.sh
+#
+# MODEL_ARGS are the engine arguments of the box (everything but --tokens-file/--max-new); the A750 needs its usual environment
+# (render group, ONEAPI_DEVICE_SELECTOR, IGC_EnableDPEmulation=1, ...).  Passes when the prefill line is printed within
+# DEADLOCK_TIMEOUT seconds (default 900; the prompt takes about 60 s on an A750 after the start-up).  Exit 0 pass, 1 fail, 2 usage.
+set -u
+HERE=$(cd "$(dirname "$0")" && pwd)
+BIN=${STRATA_BIN:-}
+[ -n "$BIN" ] && [ -x "$BIN" ] && [ -n "${MODEL_ARGS:-}" ] || { echo "usage: STRATA_BIN=<strata> MODEL_ARGS=\"...\" $0"; exit 2; }
+OUT=$(mktemp)
+timeout "${DEADLOCK_TIMEOUT:-900}" "$BIN" $MODEL_ARGS --tokens-file "$HERE/data/prompt_deadlock_2744.tok" --max-new 4 --greedy > "$OUT" 2>&1
+rc=$?
+if grep -q "^prefill  *2743 tokens" "$OUT"; then
+    echo "PASS: the 2,743-token prompt was read ($(grep -E '^prefill' "$OUT" | head -1 | tr -s ' '))"
+    rm -f "$OUT"; exit 0
+fi
+echo "FAIL: no prefill line (engine exit code $rc); output kept in $OUT"
+tail -5 "$OUT"
+exit 1