ggml-org/llama.cpp b11979

on GitHub

sycl : vectorize reordered-layout dequantization stores (Q4_0, K-quants, IQ3, MXFP4) (#30976)

• sycl : widen reordered q4_0 dequantization stores

The reordered-layout Q4_0 conversion kernel processed one qs byte per
work-item with two scalar stores, strided 16 elements apart. Give it
the same treatment the Q8_0 kernel got in #30775: a flat
sycl::range<1> launch with four qs bytes (8 outputs) per work-item
inside one 32-element block and two vec<dst_t,4> stores.

Same float math and rounding as the scalar kernel; output is
bit-identical. The existing k % QK4_0 == 0 assert guarantees whole
32-element blocks, so the k/8 launch covers every row exactly. The
now-unreferenced scalar reorder kernel is removed.

Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +6.4% to +7.3% pp512/pp2048 on Q4_0
models (tested up to 27B dense); fresh-layout prefill and decode
unchanged; PPL within run noise.

• sycl : vectorize reordered q4_K dequantization stores

The reordered-layout Q4_K conversion kernel already loads its
quantized inputs through sycl::vec but writes the dequantized
results with scalar stores. Build the results in a vec<dst_t,n>
register and emit two vector stores instead. Same per-element
math; output is bit-identical.

Arc Pro B70, primed prefill (prefill after decode on a
reorder-installed tensor): +14.7% to +15.3% pp512/pp2048 on
Q4_K-dominant models (tested up to 27B dense); fresh-layout
prefill, decode and PPL unchanged.

• sycl : remap reordered q5_K dequantization to vec4 loads/stores

The reordered-layout Q5_K conversion kernel gave each thread only
four outputs fed by scalar ql/qh byte loads plus a local-memory
scale broadcast, so the stores were never its bottleneck: widening
the stores alone measured within noise on mixed-quant models and
+0.5% on a q5_K-dominant model. Re-decompose to 32 threads per
block, each consuming four consecutive ql and qh bytes and emitting
two aligned vec<dst_t,4> stores; a 4-element run shares one scale
pair, so the per-element math and the output are identical.

Arc Pro B70, primed prefill on a q5_K-dominant 27B model: +3.6%
pp512/pp2048 over the scalar kernel; fresh-layout prefill, decode
and PPL unchanged. The GET_ROWS reorder path moves to the same
32-thread decomposition.

• sycl : remap reordered q6_K dequantization to vec4 stores

The reordered-layout Q6_K conversion kernel wrote four scalars per
thread, strided 32 elements apart, each pair of neighbors sharing
neither instruction stream nor store transaction. Re-decompose the
same 64-thread work-group so each thread emits one aligned
vec<dst_t,4>: a 4-element run shares one 16-element scale group,
one ql nibble source and one qh bit pair, so the per-element math
and the output are identical to the scalar kernel.

Arc Pro B70, primed prefill: +28.7% pp512/pp2048 on a Q6_K 27B
model; fresh-layout prefill, decode and PPL unchanged.

• sycl : vectorize reordered q3_K dequantization stores

Each thread of the reordered-layout Q3_K conversion kernel already
computes four consecutive outputs; build them in a vec<dst_t,4>
register and emit one vector store instead of four scalars. Same
per-element math; output is bit-identical.

Arc Pro B70, primed prefill: +16.8% pp512/pp2048 on a Q3_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.

• sycl : vectorize reordered iq3_xxs/iq3_s dequantization stores

Both reordered-layout IQ3 conversion kernels write eight contiguous
outputs per thread as scalar stores (four from each grid lookup).
Build two vec<dst_t,4> registers and emit two vector stores. Same
per-element math; output is bit-identical.

Arc Pro B70, primed prefill: +20.3% to +21.1% pp512/pp2048 on
IQ3_XXS- and IQ3_S-dominant 27B models; fresh-layout prefill,
decode and PPL unchanged.

• sycl : remap reordered q2_K dequantization to vec4 stores

Re-decompose the reordered-layout Q2_K conversion kernel the same
way as Q6_K: one aligned vec<dst_t,4> per thread instead of four
scalars strided 32 apart. A 4-element run shares one scale byte and
one 2-bit lane, so the per-element math and the output are
identical to the scalar kernel.

Arc Pro B70, primed prefill: +13.1% pp512/pp2048 on a Q2_K-dominant
27B model; fresh-layout prefill, decode and PPL unchanged.

• sycl : widen reordered mxfp4 dequantization stores

Give the reordered-layout MXFP4 conversion kernel the same flat
sycl::range<1> launch as Q4_0: four qs bytes (8 outputs) per
work-item inside one 32-element block and two vec<dst_t,4> stores.
Same LUT lookups and per-block scale; output is bit-identical. The
k % QK_MXFP4 == 0 assert guarantees whole blocks, so the k/8 launch
covers every row exactly. The old scalar kernel is removed.

The reordered layout for MXFP4 is installed on the MoE expert path
(reorder_qw_mxfp4_moe); dense MXFP4 tensors keep the normal layout.

Arc Pro B70, MXFP4 MoE 20B model: +57% pp2048 / +61% pp512 prefill
(primed and fresh alike), decode unchanged, PPL within the model's
run-to-run noise.