[Engine] Intel (SYCL) 0.1.56: DFlash2 accepts 0 of every draft with a Q6_K output head (IQ3_S): the draft head is gathered as rows from the STRATA_Q6_PLANAR planes

sycl/src/core/dflash2.cpp の Dflash2Drafter::bind から始め、MTP パスが sycl/src/core/mtp.cpp で平面 Q6_K 行をどのように収集するかと、sycl/src/core/native_head.cpp でのレイアウト処理を比較します。SYCL ビルドと Q6_K 出力ヘッドを使って再現し、ドラフトヘッドが平面レイアウトから収集されることと、ドラフト受け入れ数がゼロのままにならないことを確認します。

Area

Engine

Desktop app, Terminal or source?

Source (./setup.sh --backend sycl, engine built with sycl/tools/build.sh in the strata-sycl-dev image, AOT bmg-g31)

What happened

On an Arc Pro B70 with the Qwen3.8-Flash-Next IQ3_S native pack (ISTA-DASLab GSQ-RCO GGUF, output.weight is Q6_K), --draft dflash2 starts, prints DFlash2 drafter on, and then every request logs drafts accepted 0 of N. Decode drops from about 52 tok/s with the MTP layer to 11-12 tok/s: each round drafts a window, the verifier rejects all of it, one token comes out. The committed text is the greedy text (identical to the MTP arm in 20 of 20 A/B pairs), so only the drafter's logits are wrong.

Cause. sycl/src/core/dflash2.cpp builds its draft head with kernels::gather_rows(head->weights(), row_bytes, ...). On the SYCL engine a Q6_K output head is converted at load to the planar layout (sycl/src/core/native_head.cpp -> native_q6_k_planar_adopt, STRATA_Q6_PLANAR default on; the log says the output head (248320 x 2560) is stored as 16-byte-aligned planes). Gathering GGUF rows out of a planar buffer gives a garbage head, and native_mmvq(14, ...) then serves it as planar as well. The MTP draft head on the same path (sycl/src/core/mtp.cpp: "STRATA_Q6_PLANAR: the LM head is stored as planes; the subset is gathered from them, planar") already handles this with native_q6_k_planar_gather; the DFlash2 port does not.

Only Q6_K heads are made planar, so a Q8_0-head model is fine: the Unsloth UD-Q4_K_XL in #5138 gets 83-85 % acceptance on the same card. The draft-head size in the log tells the two apart: 212.9 MiB here (Q6_K, 2100 B/row) vs 275.7 MiB in #5138 (Q8_0, 2720 B/row).

Ruled out first (one fresh engine each, same prompts, all 0 of N): --kv fp16 instead of int8; no control vector; STRATA_PROMPT_ATTN_ESIMD=off; an engine built with STRATA_SYCL_IQ_DECODE_FAST=OFF.

Proof. STRATA_Q6_PLANAR=0 alone: the same requests go to 6 of 6, 97 of 117, 65 of 98, 68 of 99, 66 of 90 accepted, 50-73 tok/s.

Fix. Gather the subset from the planes when the head is planar, as mtp.cpp does. With it the engine logs STRATA_Q6_PLANAR: the DFlash2 draft head (106299 x 2560) is gathered from the planar head and the numbers below follow.

diff --git a/sycl/src/core/dflash2.cpp b/sycl/src/core/dflash2.cpp
index 0ab23a15..79b51891 100644
--- a/sycl/src/core/dflash2.cpp
+++ b/sycl/src/core/dflash2.cpp
@@ -300,7 +300,19 @@ bool Dflash2Drafter::bind(const WeightTable& wt, const NativeHead* head, const s
             return false;
         }
         cs_->memcpy(dv, vocab.data(), vocab.size() * 4).wait();
-        kernels::gather_rows((const uint8_t*) hw, row_bytes, dv, (int64_t) vocab.size(), dhead_, cs_);
+        bool planar_gathered = false;
+        if (head->type() == 14 && kernels::native_q6_k_is_planar(hw)) {
+            // STRATA_Q6_PLANAR (local patch, see the lab notes): the output head is stored as 16-byte-aligned planes (native_head.cpp), so its
+            // rows cannot be gathered as GGUF rows; the subset is gathered from the planes, planar, as the MTP draft head is (mtp.cpp).
+            // Gathered as rows, the draft head is garbage and the engine accepts 0 of every draft (Arc Pro B70, IQ3_S pack, Q6_K head).
+            void* pd = kernels::native_q6_k_planar_gather(hw, (int) (row_bytes / 210) * 256, dv, (int64_t) vocab.size(), "DFlash2 draft head");
+            if (pd != nullptr) {
+                sycl::free(dhead_, *cs_);
+                dhead_ = (uint8_t*) pd;
+                planar_gathered = true;
+            }
+        }
+        if (!planar_gathered) kernels::gather_rows((const uint8_t*) hw, row_bytes, dv, (int64_t) vocab.size(), dhead_, cs_);
         cs_->wait();
         sycl::free(dv, *cs_);
         vram_ += (uint64_t) ((int64_t) vocab.size() * row_bytes);

(The registry entry of the gathered head is not removed in the destructor, same as the MTP head today.)

Result with the fix, MTP vs DFlash2 on the same config, greedy, 256 new tokens, alternating fresh engines, medians of 3 rounds (chat and the 6,345-token prompt) and 2 rounds (the 21,102-token prompt). MTP is --mtp rt --spec 4 --spec-min-p 0.5; DFlash2 is what strata_control/drafting.py picks for SYCL, --spec 6 --draft-conf 0.6, no MTP layer loaded:

Request	MTP decode	DFlash2 decode	accepted MTP -> DFlash2
chat, code answer (34-token prompt)	59.0 tok/s	68.5 tok/s	81 % -> 87 %
chat, prose (27)	52.8	54.8	56 % -> 64 %
chat, prose (29)	51.5	52.7	58 % -> 64 %
after a 6,345-token prompt	48.6	49.2	64 % -> 70 %
after a 21,102-token prompt	47.5	48.0	65 % -> 70 %

Prompt reading unchanged (1608 vs 1607 tok/s at 6.3k, 1757 vs 1761 at 21k). The drafter takes 807 MiB and leaves 37 more expert slots (12141 vs 12104) than the MTP layer. Every run was deterministic across rounds.

Two side notes:

1. STRATA_Q6_PLANAR=0 cannot be passed through sycl/serve/strata-sycl.sh: the launcher drops every STRATA_* value of exactly 0 (#4546). For the test I let that one variable through locally.
2. A stock Intel install never reaches this engine because the serve layer falls back to MTP first (the two bugs in #5138: the drafter check reads the launcher script, and /work paths are checked on the host). Once those are fixed, Auto starts DFlash2 and every IQ3_S user gets the 0-acceptance engine.

Smallest steps or command that reproduce it

# Flash-Next IQ3_S native pack (Q6_K output head), engine 0.1.56 SYCL, one card, greedy
strata drafting fetch --data-dir <data> --config <config.json> --mode dflash2
build-sycl-aot/strata --serve --pack <data>/packs/iq3_s --native <IQ3_S shard 1> --ple-gguf <IQ3_S shard 2> \
  --expert-cache auto --max-context 262144 --kv int8 --stream-experts --prefill 4096 --kv-resident 32768 --vram-reserve-mib 2048 \
  --draft dflash2 --draft-model <data>/dflash2/qwen3.8-flash-next-dflash2.q4_0.pack --draft-fallback none --mtp <data>/mtp/rt --spec 6 --draft-conf 0.6
# any chat request at temperature 0 -> "drafts accepted 0 of N"; with STRATA_Q6_PLANAR=0 in the engine's environment -> 60-85 % accepted

Reproduction status

Reproduced repeatedly on one machine (every request of 11 engine starts: the default, four variants, and the planar-off and patched runs)

Strata version

v0.2.10, engine 0.1.56 (strata 0.1.56, built from the tag with sycl/tools/build.sh, AOT bmg-g31)

Graphics card, memory and system

Intel Arc Pro B70 32 GB, xe driver, Debian 13 (kernel 7.2), oneAPI 2026.1 in the strata-sycl-dev image, 62 GB RAM

Model, size and settings

Qwen3.8-Flash-Next-GSQ-RCO IQ3_S (ISTA-DASLab, 2 shards, native pack from tools/iq_pack.py), 262144 context, --kv int8 --kv-resident 32768 --prefill 4096 --stream-experts --vram-reserve-mib 2048, control vector experimental-speed-projection layers 4-44, STRATA_PREFILL_MMQ=off, one request at a time

Logs

Stock engine:

strata generate: native pack: /work/<data>/packs/iq3_s experts (largest blob 2.66 MB), token embedding IQ4_XS in mapped host memory (322 MiB, 0.4 s)
strata: STRATA_Q6_PLANAR: the output head (248320 x 2560) is stored as 16-byte-aligned planes (same bytes)
strata dflash2: 106299-token draft head (212.9 MiB), drafter 807 MiB of VRAM (Q4_0 weights)
strata serve: DFlash2 drafter on (/work/<data>/dflash2/qwen3.8-flash-next-dflash2.q4_0.pack, graphs captured in 15 ms); windows up to 8 tokens; no MTP layer is loaded
strata serve: prompt 34 tokens = 0 reused + 34 read in 406 ms (83.8 tok/s), 256 generated in 23543 ms (10.9 tok/s), drafts accepted 0 of 1267, 1 checkpoints
strata serve: prompt 27 tokens = 0 reused + 27 read in 380 ms (71.0 tok/s), 256 generated in 20604 ms (12.4 tok/s), drafts accepted 0 of 1275, 1 checkpoints
strata serve: prompt 29 tokens = 0 reused + 29 read in 374 ms (77.5 tok/s), 208 generated in 17997 ms (11.6 tok/s), drafts accepted 0 of 1035, 1 checkpoints
strata serve: prompt 6345 tokens = 0 reused + 6345 read in 3969 ms (1598.8 tok/s), 149 generated in 12429 ms (12.0 tok/s), drafts accepted 0 of 740, 1 checkpoints

Same engine with STRATA_Q6_PLANAR=0:

strata serve: prompt 34 tokens = 0 reused + 34 read in 395 ms (86.1 tok/s), 128 generated in 1763 ms (72.6 tok/s), drafts accepted 97 of 117, 1 checkpoints
strata serve: prompt 27 tokens = 0 reused + 27 read in 376 ms (71.8 tok/s), 128 generated in 2397 ms (53.4 tok/s), drafts accepted 65 of 98, 1 checkpoints
strata serve: prompt 6345 tokens = 0 reused + 6345 read in 3912 ms (1621.8 tok/s), 128 generated in 2535 ms (50.5 tok/s), drafts accepted 66 of 90, 1 checkpoints

Patched engine (planar on):

strata: STRATA_Q6_PLANAR: the DFlash2 draft head (106299 x 2560) is gathered from the planar head
strata serve: prompt 34 tokens = 0 reused + 34 read in 405 ms (84.0 tok/s), 256 generated in 3737 ms (68.5 tok/s), drafts accepted 194 of 224, 1 checkpoints
strata serve: prompt 27 tokens = 0 reused + 27 read in 346 ms (78.1 tok/s), 256 generated in 4672 ms (54.8 tok/s), drafts accepted 134 of 209, 1 checkpoints
strata serve: prompt 29 tokens = 0 reused + 29 read in 355 ms (81.8 tok/s), 208 generated in 3947 ms (52.7 tok/s), drafts accepted 110 of 172, 1 checkpoints
strata serve: prompt 6345 tokens = 0 reused + 6345 read in 3950 ms (1606.7 tok/s), 149 generated in 3028 ms (49.2 tok/s), drafts accepted 76 of 108, 1 checkpoints