Skip to content

sycl: add PTQ1_0 and PQ2_0 support with MMVQ vector dot kernel - #235

Merged
bri-prism merged 7 commits into
PrismML-Eng:prismfrom
kiljoy001:sycl-ptq1-pq2-mmvq
Sep 24, 2026
Merged

bri-prism merged 7 commits into
PrismML-Eng:prismfrom
kiljoy001:sycl-ptq1-pq2-mmvq

Conversation

@kiljoy001

@kiljoy001 kiljoy001 commented Sep 21, 2026 •

Copy link
Copy Markdown

Summary

This PR implements native Intel SYCL / oneAPI support for PTQ1_0 (base-3 trit packed ternary) and PQ2_0 (group-128 2-bit quantization), enabling full GPU offload and execution of ternary models like Ternary-Bonsai-2-27B-PTQ1_0 on Intel Arc GPUs.

Before this PR the SYCL backend had no reference to either format, so these models aborted during load at convert.cpp and could not run on SYCL at all.

Changes

1. Dequantization (dequantize.hpp, convert.cpp)

  • Implemented ptq1_0_trit() base-3 recurrence unpacking and dequantize_ptq1_0() / dequantize_pq2_0() DPC++ kernels.
  • Added PTQ1_0 and PQ2_0 conversion paths in ggml_get_to_fp32_sycl, ggml_get_to_fp16_sycl, and ggml_get_to_fp16_nc_sycl.

2. MMVQ vector dot kernels (vecdotq.hpp, mmvq.cpp) — both formats get a native kernel and full dispatch:

  • vec_dot_ptq1_0_q8_1, VDR_PTQ1_0_Q8_1_MMVQ = 4: evaluates the 128-weight ternary block directly in registers across sub-group lanes, no intermediate FP32 VRAM. Trits are unpacked SIMD-style from 32-bit loads and reduced with DP4A.
  • vec_dot_pq2_0_q8_1, VDR_PQ2_0_Q8_1_MMVQ = 1 with MMVQ_PQ2_0_QI = 16: 16-bit loads through a shared unpack_2bit_to_byte_lanes() permute helper. The dedicated QI spreads each block over 16 lanes rather than the 4 that QI_PQ2_0 would imply, avoiding an 8-deep serial DP4A chain per lane — measured ~2.1x on the mat-vec.
  • Single-column and multi-column (switch_ncols, up to 8) dispatch for both types.

3. Type gating (ggml-sycl.cpp)

  • can_use_mul_mat_vec_q() now gates on the set of types MMVQ actually implements instead of ggml_is_quantized(). Previously the backend claimed every quantized type and mul_mat_vec_q_switch_type() aborted the process on any type without a case — which is why test-backend-ops -o MUL_MAT could not complete on SYCL. Unsupported types now fall back cleanly to CPU. Matching check in supports_op().

4. Operators (getrows.cpp, cpy.cpp)

  • PTQ1_0 / PQ2_0 support in get_rows (embedding gather) and cpy.

Verification

All numbers on Intel Arc Pro B50 (Battlemage, 16 GB), oneAPI 2026.1 / Level Zero, Release.

Correctness — test-backend-ops compares each op against the CPU reference numerically:

PTQ1_0 PQ2_0
MUL_MAT 102/102 102/102
MUL_MAT_ID 75/75 75/75

Coverage includes real Bonsai weight shapes (k = 1024/5120/6144/17408) at odd row counts (m = 67/70) with n = 1..8, added in this PR — the stock generator only emits 16x256 single-column cases, so the multi-column dispatchers had no coverage.

Full -o MUL_MAT sweep: 1098/1098 passed. Before the gating fix this run aborted at mmvq.cpp:2710.

Performance — test-backend-ops perf -o MUL_MAT, m=4096 k=14336:

n PTQ1_0 PQ2_0
1 92.01 us (1.28 TF) 114.66 us (1.02 TF)
4 212.94 us (2.21 TF) 221.49 us (2.12 TF)
8 358.27 us (2.62 TF) 381.36 us (2.46 TF)
512 9224 us (6.52 TF) 8984 us (6.69 TF)

End-to-end — Ternary-Bonsai-2-27B-PTQ1_0.gguf, all 64 layers on SYCL0: 2.20 tok/s, vs 0.89 naive dequant and 0.39 CPU.

Notes and limitations

  • Figures are from one device. Battlemage-class Arc; not generalized to other Arc parts or other SYCL targets.
  • End-to-end verification is output coherence and deterministic evaluation via llama-simple. Per-op numerical agreement is covered by test-backend-ops above; a full perplexity comparison against CPU is not yet done.
  • ggml_sycl_op_fwht supports widths up to 512, so models with hadamard.block_size = 1024 take the dense mat-mul fallback. Numerically correct, and measured at a negligible share of a decode frame at batch 1, but the wide-kernel gap versus CUDA/Metal/Vulkan remains open and is out of scope here.
  • The full test-backend-ops suite still stops at set_rows.cpp:549, a pre-existing abort in an unrelated op.

- Implement base-3 trit unpacking and dequantization (dequantize_ptq1_0, dequantize_pq2_0) in dequantize.hpp and convert.cpp
- Implement vec_dot_ptq1_0_q8_1 and mul_mat_vec_ptq1_0_q8_1_sycl in vecdotq.hpp and mmvq.cpp
- Add get_rows and cpy dispatch for PTQ1_0 and PQ2_0
- Enable PTQ1_0 in can_use_mul_mat_vec_q() for fast in-register Level-Zero vector execution
- Verified on Intel Arc Pro B50 with oneAPI 2026.1 / Level-Zero (2.20 tok/s on 27B)
@kiljoy001

Copy link
Copy Markdown
Author

Performance Update: SIMD Vectorized PTQ1_0 Unpacking + DP4A

Replaced the serial recurrence loop in vec_dot_ptq1_0_q8_1 with SIMD 4-trit unpacking and hardware dpct::dp4a (commit 016be85).

Benchmark results on Intel Arc Pro B50 (16GB, SYCL / Level-Zero):

  • Prompt processing (pp128): 78.66 ± 0.17 t/s
  • Text generation (tg32): 13.64 ± 0.03 t/s (up from 2.20 t/s baseline, 6.2x speedup!)

Exhaustive integer validation passed on all 8,192 byte configurations with zero mismatches against CPU reference.

@kiljoy001

Copy link
Copy Markdown
Author

Performance Update: Native PQ2_0 MMVQ Vector Dot Kernel

Implemented 4-thread cooperative DP4A vector dot kernel for PQ2_0 (vec_dot_pq2_0_q8_1) and multi-column dispatch in mmvq.cpp (commit c5f6c97).

Benchmark results on Intel Arc Pro B50 (16GB, SYCL / Level-Zero):

  • Ternary Bonsai 1.7B (PQ2_0): 83.85 ± 1.75 t/s (up from 13.58 t/s fallback, 6.17x speedup!) | Prompt: 1,515 t/s
  • Ternary Bonsai 4B (PQ2_0): 39.13 ± 0.07 t/s | Prompt: 596 t/s
  • Ternary Bonsai 8B (PQ2_0): 22.73 ± 0.08 t/s | Prompt: 318 t/s

@bri-prism
bri-prism self-requested a review September 21, 2026 16:32
@bri-prism

Copy link
Copy Markdown
Collaborator

Data point: Intel Arc B390 iGPU (Panther Lake Xe3), SYCL / Windows — builds and passes, and it moves a pre-existing abort downstream

Adding an integrated Battlemage-era part alongside the discrete Arc Pro B50 in the description.

SYCL0: Intel(R) Arc(TM) B390 GPU (55789 MiB)
CPU:   Intel(R) Core(TM) Ultra X7 358H
Build with Macros: GGML_SYCL_DNNL=yes  GGML_SYCL_F16=no  GGML_SYCL_GRAPH=yes  GGML_SYCL_SUPPORT_LEVEL_ZERO_API=no

Built with oneAPI 2025.3 + MSVC 2022 Build Tools, Ninja, Release: 360/360 targets, clean. Worth stating explicitly since this PR's only CI job is labeler — there is no build or test coverage on it.

Compared prism @ 422590f5d against this PR @ c5f6c9718, same toolchain and flags for both.

test-backend-ops test -o MUL_MAT -b SYCL0

prism 422590f5d #235 c5f6c9718
cases OK 90 171
FAIL 0 0
PTQ1_0 0 9 OK
PQ2_0 0 9 OK
run ends mmvq.cpp:2562: fatal error: unsupport data type=pq2_0 mmvq.cpp:2710: fatal error: unsupport data type=tq2_0

So on this device stock prism aborts the MUL_MAT suite on PQ2_0; with this PR, PQ2_0 and PTQ1_0 both execute and match the CPU reference, and the abort moves on to the next type that has no mmvq kernel. That is a strict improvement, and the new kernels are correct on every case that ran.

GET_ROWS likewise: 8 cases (4 PTQ1_0 + 4 PQ2_0) go from not supported to OK. The full-suite totals account for exactly that — 658 OK / 78 not-supported on prism, 666 / 70 here.

The abort class is worth fixing separately

Both endpoints above are the same shape: can_use_mul_mat_vec_q() gates on

return ggml_is_quantized(src0->type) && ...

so the backend claims every quantized type, and mul_mat_vec_q_switch_type then GGML_ABORTs on any type it has no case for. That is why prism dies on pq2_0 and this branch dies on tq2_0 — this PR fixes two instances of the bug without changing the pattern, so the next quant type added will land in the same trap.

Gating on the set of types actually implemented (or returning false instead of aborting) would turn a hard crash into a normal fallback. Not a blocker for this PR, but it is the reason a full test-backend-ops run cannot complete on SYCL on this box.

Unrelated and also pre-existing: the full suite aborts identically on both builds at set_rows.cpp:549: Unsupported tensor type!, same mechanism.

Caveat

Because both runs abort partway, the PTQ1_0/PQ2_0 coverage here is only the 9 MUL_MAT cases per type that the stock suite generates before the abort. #188 adds Bonsai-shape MUL_MAT/MUL_MAT_ID cases (k = 1024/5120/6144/17408, odd row counts 67/70, n = 1..8) for exactly these two types on the Vulkan side; porting those test cases over would give this PR's kernels much better coverage than they currently get.

@dwymark

dwymark commented Sep 21, 2026

Copy link
Copy Markdown

Human here. I can help with testing on hardware later tonight if that helps. LG Gram 17 w/ Lunar Lake, 32gb unified memory.

@bri-prism

Copy link
Copy Markdown
Collaborator

Follow-up on the same machine (Arc B390 iGPU, oneAPI 2025.3, Windows) — end-to-end this time, and the before/after is stronger than the test-backend-ops numbers suggested.

On stock prism @ 422590f5d, SYCL cannot load either Bonsai 2 27B file at all:

llama-bench -m Ternary-Bonsai-2-27B-PTQ1_0.gguf -p 512 -n 128 -ngl 99 -fa 1
  -> ggml/src/ggml-sycl/convert.cpp:832: fatal error: unsupport data type=ptq1_0

llama-bench -m Ternary-Bonsai-2-27B-PQ2_0.gguf  -p 512 -n 128 -ngl 99 -fa 1
  -> ggml/src/ggml-sycl/convert.cpp:832: fatal error: unsupport data type=pq2_0

Hard abort during load, both bands. With this PR @ c5f6c9718 both run:

model pp512 (t/s) tg128 (t/s)
Ternary-Bonsai-2-27B-PTQ1_0 93.67 ± 0.09 11.68 ± 0.38
Ternary-Bonsai-2-27B-PQ2_0 89.63 ± 4.40 5.47 ± 0.03

So the practical before/after here is not a speedup ratio — it is "aborts on load" to "runs correctly", which is a better headline for this PR than the case counts I posted above.

Caveat on those numbers: this build prints warning: asserts enabled, performance may be affected (CMAKE_BUILD_TYPE=Release, icx 2025.3, Ninja), so treat both rows as a floor, not this PR's peak.

One thing that may be worth a look: on the same box and the same two files, the Vulkan backend with PQ2_0 support (#188) gets tg128 11.56 on PQ2_0 versus 5.47 here, while PTQ1_0 is much closer (Vulkan 13.16 with mmvq forced, 11.68 here). Different backends and the asserts caveat above both apply, so this is a hint rather than a measurement of your kernel — but it does suggest the PQ2_0 decode path has more headroom than the PTQ1_0 one, if you are looking for somewhere to push next.

@bri-prism

Copy link
Copy Markdown
Collaborator

Correction to the caveat on my previous comment, in the direction of the numbers being more trustworthy, not less.

I flagged those results as a floor because the build printed warning: asserts enabled. Digging into why, the cause was worse than asserts: that build directory had CMAKE_CXX_FLAGS_RELEASE empty and zero occurrences of NDEBUG in build.ninja — so no /DNDEBUG and no /O2 either. (Self-inflicted: two failed cmake -B configures in that tree left a partial CMakeCache.txt that the working configure then reused. The baseline SYCL build I compared against was unaffected — it had /O2 /Ob2 /DNDEBUG correctly.)

I rebuilt from a clean directory with the flags set explicitly (/O2 /Ob2 /DNDEBUG, now 629 NDEBUG occurrences, matching the baseline build, and the asserts warning is gone) and re-ran both bands, two rounds each:

model metric previous (unoptimized) rebuilt (r1 / r2)
PTQ1_0 tg128 11.68 11.45 / 10.90
PTQ1_0 pp512 93.67 94.39 / 90.27
PQ2_0 tg128 5.47 5.37 / 5.19
PQ2_0 pp512 89.63 84.68 / 77.72

No material change — flat to marginally lower, within run-to-run spread. Which is what you would expect for a GPU-bound workload: the device kernels are compiled by the SYCL device compiler regardless of the host-side flags, so unoptimized host code barely registers here.

Consequences for what I wrote earlier:

  • The numbers in my previous comment should no longer be read as a floor. They are representative.
  • The PQ2_0 observation stands, and now rests on a sound build: PQ2_0 decode ~5.2-5.4 t/s against PTQ1_0's ~10.9-11.5 on the same device, model family and flags, while the Vulkan backend on this same box gets ~11.5 on PQ2_0. So the asymmetry is in the SYCL PQ2_0 decode path, not in my toolchain. That remains the most interesting place to look if you want more out of this PR.

Apologies for the noise — I would rather re-measure and confirm than leave a hedge attached to a number that turned out to be fine.

The MMVQ kernel derives its lane split from qi/vdr. PQ2_0 was passing
QI_PQ2_0 (= QK_PQ2_0/32 = 4), which put only 4 lanes on each 128-element
block and left every lane an 8-deep serial dp4a chain over 8 bytes of qs.

Pass a dedicated MMVQ_PQ2_0_QI of 16 instead, so a block is covered by 16
lanes reading 2 bytes / 8 elements each with 2 dp4a per call, and a
sub-group keeps more blocks in flight. Measured on an Arc iGPU with
test-backend-ops perf, MUL_MAT m=4096 n=1 k=14336, alternating the two
builds to cancel out background load: 827/846/837 us before, 397/395/384
us after (~2.1x). On an idle machine the same case goes 247.6 -> 123.0 us.
8 and 32 lanes per block were both tried and were slower; at 32 the
per-call scale setup dominates the single remaining dp4a.

Also fold the 2-bit-to-byte-lane unpack shared with the q2_0 path into
unpack_2bit_to_byte_lanes(), and switch the epilogue to the q4_0-style
sumi*ds.x - zp*ds.y form. The previous PQ2_0 epilogue applied the zero
point with byte_sub_4 and then used only ds[0], dropping the ds.y term.

test-backend-ops test -o MUL_MAT -p type_a=pq2_0: 38/38 on both SYCL
devices, 3/3 backends passed. The unrelated q1_0 failures on SYCL1 are
present with and without this change.

@bri-prism bri-prism left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reviewed at c5f6c97. This fills a real hole: at tag prism-b10709-9a9394a the SYCL backend had zero references to either format, so ternary models had no native path there at all. Having actual Battlemage silicon behind the numbers is worth a lot, since nobody else working on these formats has Intel hardware.

One finding that is not a defect in this PR but that I think changes how you read your own throughput number, then two smaller things.

The Hadamard transform is falling back to a dense mat-mul on SYCL, and this PR cannot see it.

I read the published Ternary-Bonsai-2-27B-PTQ1_0.gguf header directly. It carries:

prism.hadamard.block_size = 1024
prism.hadamard.transform  = normalized-sylvester-walsh-hadamard
prism.hadamard.weight_names = <401 strings>

ggml/src/ggml-sycl/fwht.cpp dispatches only widths 64, 128, 256 and 512, and returns false on anything else. At 1024 it returns false on every call, and the caller falls through to the ordinary mat-mul dispatch. That path is numerically correct, because llama_mul_mat_hadamard in src/llama-impl.h builds the matmul against a materialized rotation tensor rather than a placeholder, so you are getting right answers. You are getting them by running a dense 1024x1024 multiply in place of an O(n log n) transform, roughly a hundred times the arithmetic, on each of 401 folded weights per token.

For comparison, the CUDA, Metal and Vulkan backends all carry the full eight widths from 64 to 8192. SYCL is the only one capped at 512, and this PR does not touch fwht.cpp.

I raise it because 2.20 tok/s is the number a reader will attribute to your MMVQ kernels, and a good part of it is probably this instead. Before you tune the dot products further it is worth checking how much of the frame is going into those dense multiplies. If it is what I expect, the wide SYCL FWHT kernels are a larger win than anything left in the vector dot.

Worth coordinating: someone on our side has started writing SYCL wide-kernel FWHT for exactly this gap. I will make sure they see this PR so the two do not collide, and so they know there is Battlemage hardware that could validate their work, which is the thing they currently lack.

The verification is a coherence check, not a correctness gate.

"Verified output coherence and deterministic evaluation with llama-simple" establishes that the model is not producing garbage, which is a real milestone for a new backend. It does not establish numerical agreement, and coherent output coexists comfortably with meaningful quality loss, particularly on a 2-bit format where a subtly wrong dequant still yields fluent text.

Could you add test-backend-ops -o MUL_MAT filtered to PTQ1_0 and PQ2_0 on SYCL0, with the case counts, and ideally a logit or perplexity comparison of the same prompt against the CPU backend on the same file? The case counts matter as much as the pass, since a backend that silently declines to support a shape reports a green run over a smaller population than you think.

Your summary undersells the PR.

The description says PTQ1_0 was enabled in can_use_mul_mat_vec_q(), which reads as though PQ2_0 got dequant and a kernel but no dispatch. In the actual diff both are dispatched, at mmvq.cpp:2464 and :2478, with single and multi-column paths for each. Worth correcting so a reviewer does not go looking for dead code as I did.

Smaller notes.

The GGML_ABORT on unsupported ncols_dst in both multi-column dispatchers is the right call over a silent fallback, and I would keep it.

The numbers are one configuration on one device. 2.20 against 0.89 for naive dequant is a strong relative result, but there are no prefill figures and no second device. That is fine for landing a first implementation, just worth not generalizing to Arc as a whole in the PR text.

Method note: source reading at c5f6c97, plus GGUF header reads over HTTP range requests against the published artifact. I have no Intel GPU and ran nothing.

@bri-prism

Copy link
Copy Markdown
Collaborator

Correcting my own review above. The FWHT point is real as a fact and wrong as a conclusion, and I would rather retract it clearly than leave you chasing it.

What I got wrong. I said the width-512 cap in ggml-sycl/fwht.cpp was "probably a large part of" your 2.20 tok/s. I never checked what share of a decode frame the Hadamard actually occupies. Doing that now:

At block_size = 1024 over 401 folded weights, the transform is ~4.1 M MAC per token and the dense fallback is ~420 M MAC per token. The 102x ratio I quoted is right. But 420 M MAC is about 0.08 ms at 5 TFLOP/s, against an 85 ms frame at 11.68 tok/s. That is roughly 0.1% of decode. At batch 1 this model is bound by streaming ~7 GB of 2-bit weights, and the extra arithmetic on a 1024-wide activation hides underneath that entirely. A hundred times more work in an operation that costs a tenth of a percent is still a tenth of a percent.

The measurements already in this thread say the same thing, and I should have read them before writing my review rather than after. On the same box and the same files, PTQ1_0 is 13.16 tg128 on Vulkan, which has the full 64-through-8192 FWHT, against 11.68 here with the cap. A ~13% gap across two different backends bounds the FWHT contribution well below what I implied. If the fallback were dominating, that gap would not be 13%.

I also misattributed the number. The 2.20 tok/s in your description is the discrete Arc Pro B50. The 11.68 figure is the B390 iGPU from the follow-up comment, on a build that turned out to have neither /O2 nor NDEBUG. So I reasoned about a throughput figure from one device using a code path measured on another, and the unoptimized-build correction means the real number is higher still.

Where the headroom actually is. The same comparison that clears the FWHT points somewhere much more interesting: PQ2_0 is 11.56 on Vulkan against 5.47 here, a 2.1x gap, while PTQ1_0 sits within 13%. Two backends and an unoptimized build both apply, so treat it as a direction rather than a measurement. But PTQ1_0 looks close to the achievable range on this hardware and PQ2_0 does not, which makes the PQ2_0 decode path the place to push. That is the same conclusion the follow-up comment reached, and it is better grounded than mine was.

What stands from my review. The width cap is still real and still worth closing, just as an ordinary improvement rather than a blocker or a hidden explanation for your throughput. The request for test-backend-ops case counts is also answered better than I asked it: the runs above show the suite aborting at mmvq.cpp:2710 on tq2_0, so the multi-column and MUL_MAT_ID paths this PR adds never execute, and the 9 cases per type that do run are the stock shapes rather than the Bonsai ones. Porting the #188 shape cases over is the coverage that would actually exercise what you wrote.

Apologies for the noise. The finding was sound and the inference on top of it was not.

@bri-prism

Copy link
Copy Markdown
Collaborator

Followed up on where the PQ2_0 headroom actually is, since that is the gap the measurements point at. I think it is in the unpack, and the comparison that makes it clearest is inside this PR rather than against another backend.

Your two kernels are not at the same level of optimization, which matches the numbers: PTQ1_0 lands within about 13% of Vulkan while PQ2_0 sits at roughly half. The commit history says the same thing, since "vectorize PTQ1_0 dot product with SIMD trit unpacking and DP4A" is its own commit and the PQ2_0 kernel arrived after it without an equivalent pass.

PTQ1_0, in this PR, loads 32 bits at a time:

const uint32_t packed = get_int_from_uint8_aligned(bq->qs, g);
uint32_t v_lo = (packed & 0x000000FF) | ((packed & 0x0000FF00) << 8);

PQ2_0, in this PR, loads one byte at a time:

const uint8_t  q  = qs_chunk[j];
const uint32_t vi = ((q >> 0) & 0x3) | (((q >> 2) & 0x3) << 8) |
                    (((q >> 4) & 0x3) << 16) | (((q >> 6) & 0x3) << 24);

Over one 32-element chunk that is eight scalar byte loads and eight shift-mask-or chains, about eleven integer ops each, to produce eight dp4a calls.

For reference, the CUDA kernel for the same format (ggml-cuda/vecdotq.cuh) does the same eight dp4a calls from half the loads and a fraction of the ALU:

const int16_t * qs = (const int16_t *) bq2_0->qs + iqs * 4;
...
const int qe = __byte_perm(0x020100FF, 0x020100FF, q >> 0);
const int qo = __byte_perm(0x020100FF, 0x020100FF, q >> 2);
const int qx = __byte_perm(qe, qo, 0x5140);
const int qy = __byte_perm(qe, qo, 0x7362);

Four 16-bit loads instead of eight 8-bit ones, and the 2-bit code to symbol mapping is a byte-permute lookup against the constant 0x020100FF rather than a shift and mask chain. Same VDR_PQ2_0_Q8_1_MMVQ = 1 and the same per-chunk structure, so this is not a different algorithm, only a different unpack.

So the shape of the fix is the one you already applied to PTQ1_0: widen the load to a 32-bit word through get_int_from_uint8_aligned and unpack four bytes at once, rather than looping over uint8_t. You will not have __byte_perm on SYCL, but you already hand-rolled an equivalent permute for PTQ1_0 in this same file, and the same widening trick applies directly, since PQ2_0's 2-bit codes need less work to separate than base-3 trits do.

Two caveats on this. I have not profiled anything and have no Intel GPU, so this is a hypothesis that fits the measurement rather than a measured attribution. And the cross-backend comparison behind it, Vulkan 11.56 against 5.47 here, came from a build that turned out to lack /O2 and NDEBUG, so the true gap may be smaller than 2.1x. The within-PR asymmetry between your own two kernels does not depend on either caveat, which is why I would trust that part more.

If you want a cheap check before committing to the rewrite: the ratio between your PTQ1_0 and PQ2_0 tg128 on the same box should narrow substantially if the unpack is the binding constraint, and barely move if it is not.

@kiljoy001

kiljoy001 commented Sep 22, 2026 via email

Copy link
Copy Markdown
Author

@bri-prism bri-prism left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reviewed at c5f6c97 against the CPU reference in ggml-quants.c, complementing the review above (which covers the Hadamard fallback and the Battlemage numbers) rather than repeating it.

Coverage first, because it changes what the green run means. The multi-column dispatchers (mul_mat_vec_ptq1_0_q8_1_sycl_switch_ncols, ..._pq2_0_..., ncols_dst 2..8) and the MUL_MAT_ID path have not been executed by any test on this PR as far as I can tell. can_use_mul_mat_vec_q claims every quantized type, so test-backend-ops -o MUL_MAT reaches tq2_0, hits the GGML_ABORT in mul_mat_vec_q_switch_type (mmvq.cpp:2710 on this branch, as the B390 log shows), and never gets to the Bonsai-shape cases that already exist in tests/test-backend-ops.cpp around line 9300: k = 1024/5120/6144/17408, rows 67/70, n = 1..8, plus a PTQ1_0 mul_mat_id case. The nine cases per type that did run are all single-column. Turning that abort into a return-false fallback would let the suite reach them; a run that includes those cases, or -o MUL_MAT_ID, is what I would want before approving.

The kernels themselves check out. The trit ordering in both the scalar ptq1_0_trit and the SIMD path in vec_dot_ptq1_0_q8_1 matches dequantize_row_ptq1_0 exactly, including the 16-wide / 8-wide / qh stage split and the q8 int index for every stage. The 16-bit lane widening is sound (255*3 cannot carry) and the perm assembly puts the four trits in the same byte order as the q8 ints they are dotted against. With QI_PTQ1_0 = 4 and VDR = 4 the mmvq template hands each lane one whole block with iqs = 0, so ignoring iqs is correct here, not an overcount. PQ2_0 matches dequantize_row_pq2_0 (code minus 1, so 11 maps to +2), and reading qs bytewise is right given the 34-byte block puts qs at a 2-byte offset. Dispatch is consistent: DMMV does not claim these types, MMQ is off on SYCL, and the CPY hunk correctly declares quant to F32 unsupported.

Small: the can_use_mul_mat_vec_q hunk is whitespace only. The description says it enables PTQ1_0 there, but ggml_is_quantized already covers it, so the hunk can be dropped.

Not blocking, on the PQ2_0 decode gap noted above: the per-byte unpack does four shift/mask/or steps per byte plus one dp4a. Loading qs as 16-bit pairs (always 2-byte aligned at stride 34) and unpacking two bytes per step would roughly halve the ALU work in that loop.

can_use_mul_mat_vec_q() claimed every quantized type via
ggml_is_quantized(), so any type without a case in
mul_mat_vec_q_switch_type() reached the GGML_ABORT in its default arm
and killed the process. test-backend-ops -o MUL_MAT could not complete
on SYCL for this reason: it died at mmvq.cpp:2710 on tq2_0, which also
meant every case generated after that point never ran.

Gate on an explicit list of the types the switch actually implements.
The list matches the switch exactly, and is also identical to the set
ggml_get_to_fp16_sycl() handles, so a type this declines cannot be
served by the dequantize fallback either - hence the matching check in
supports_op(), which lets the scheduler place those nodes on CPU rather
than abort mid-graph.

Arc Pro B50, oneAPI 2026.1, Level Zero:
  before: test -o MUL_MAT aborts, exit 134
  after:  1098/1098 tests passed, exit 0
tq2_0 goes from "fatal error: unsupport data type" to "not supported".
The stock generator only emits 16 x 256 cases for these types, all of
which take the single-column mmvq path, so the multi-column dispatchers
(ncols_dst 2..8) had no coverage at all.

Add the real Bonsai row lengths (k = 1024/5120/6144/17408) at odd row
counts (m = 67/70, landing a partial row group) with n sweeping 1..8.

Arc Pro B50, oneAPI 2026.1, Level Zero, per type:
  MUL_MAT     9 -> 102/102 passed
  MUL_MAT_ID       75/75 passed
@kiljoy001

Copy link
Copy Markdown
Author

Thanks for the detailed review, and for the two self-corrections — the retraction on the FWHT share of the frame saved me chasing it.

One thing first, since it affects two of your points: you reviewed at c5f6c97, which was the tip at the time but not the latest work. de7a01b is now pushed and it is the PQ2_0 optimization you went on to recommend.

The PQ2_0 unpack is already widened

Your suggestion was to drop the per-byte loop for a wider load. That is de7a01b. The qs_chunk[j] code you quoted no longer exists:

uint32_t val = (uint32_t) *(const uint16_t *) (bq2_0->qs + 2 * iqs);
const uint32_t vi = unpack_2bit_to_byte_lanes(val & 0xFFu);

16-bit load, shared permute helper, same 2-byte alignment reasoning you gave. The commit also spreads a block over 16 lanes instead of 4 (MMVQ_PQ2_0_QI), which was the larger win, and fixes the epilogue to the sumi*ds.x - zp*ds.y form — the old one dropped the ds.y term.

Which changes the headroom conclusion

test-backend-ops perf -o MUL_MAT, m=4096 k=14336, Arc Pro B50:

n PTQ1_0 PQ2_0 ratio
1 92.01 us 114.66 us 1.25x
2 131.33 us ~152 us 1.16x
8 358.27 us 381.36 us 1.06x
512 9224 us 8984 us 0.97x

No 2.1x gap — PQ2_0 is within 25% at n=1 and slightly faster at n=512. Your 5.47 t/s predates de7a01b. Different silicon from your B390, so this does not contradict your measurement, just the inference from it.

(n=2 first read 536 us; that was noise. Three re-runs: 151.6 / 152.7 / 155.8.)

The abort class — fixed

You called it not a blocker; it was worth fixing anyway. can_use_mul_mat_vec_q() now gates on the types the switch implements rather than ggml_is_quantized(). That list is also identical to what ggml_get_to_fp16_sycl() handles, so anything it declines cannot be served by the dequantize fallback either — hence the matching supports_op() check, which lets the scheduler place those nodes on CPU instead of aborting mid-graph.

Before: exit 134 at mmvq.cpp:2710. After: 1098/1098 passed, exit 0, and tq2_0 reports not supported.

The suite now runs to set_rows.cpp:549 — the unrelated pre-existing abort you identified, same pattern, different op.

Coverage

Added the Bonsai shapes (k = 1024/5120/6144/17408, m = 67/70, n = 1..8). Per type, on B50:

  • MUL_MAT: 9 -> 102/102
  • MUL_MAT_ID: 75/75

So the multi-column dispatchers and the MUL_MAT_ID path both execute now. Note test-backend-ops compares each op against the CPU reference numerically, so these are agreement checks, not just non-crash checks.

Two corrections

The can_use_mul_mat_vec_q hunk was not whitespace-only at c5f6c97 — it added src0->type != GGML_TYPE_PQ2_0, a real exclusion while the PQ2_0 kernel was being brought up. Dropping it as suggested would have regressed. It is now replaced properly.

You are right on the description. It is updated: both types are shown as fully dispatched with their own kernels and switch_ncols chains, the gating fix is called out as its own item, and the numbers are scoped to the specific device rather than Arc generally.

Still open

  • Perplexity/logit vs CPU — fair, and not done. The per-op comparisons above are stronger evidence than the coherence check I originally cited, but they are not whole-model quality.
  • FWHT width cap — real, unchanged here. I read the CUDA strategy; fwht_cuda_block alone covers 512–8192 and ports cleanly (__shfl_xor_sync -> permute_sub_group_by_xor, __shared__ -> local_accessor). One thing for whoever picks it up: SYCL builds with GGML_SYCL_WARP_SIZE=16 on some targets, so the register-path math is tighter than CUDA's at a given N, and the SYCL op currently has no signs or F16-input variant. Given the ~0.1% figure and that you said someone is already on it, I would rather not duplicate — I have B50 silicon and am happy to validate their work, which sounds like the thing that is missing.

@dwymark

dwymark commented Sep 23, 2026

Copy link
Copy Markdown

Data point: Intel Arc 140V (Lunar Lake laptop), Windows, oneAPI 2025.3, at head 450c5c3eb, with a whole-model logit check

Daniel (dwymark) instructed a series of coding agents to reproduce this PR's measurements on his laptop, and worked with agents on further optimizations of the same kernels. Claude Code wrote this comment in Claude's voice on Daniel's instruction. The first section is the reproduction; the second section describes the optimization work.

Reproduction

Two commits are compared throughout: the base, c5f6c9718, which is the first native PQ2_0 kernel in this PR, and the head, 450c5c3eb. The only kernel change between them is de7a01b.

Environment and method:

  • Intel graphics driver 32.0.101.9030
  • Fresh Ninja Release builds with /O2 /Ob2 /DNDEBUG
  • llama-bench -p 512 -n 128 -r 2 -ngl 99 -fa 1, with GGML_SYCL_ENABLE_DNN=0
  • One process per build, alternating between builds

Values are median tokens per second. Repetitions agree within 6 percent.

Model Build Prompt 512 Decode 128
Bonsai 2 27B PQ2_0 base c5f6c9718 46.9 3.27
Bonsai 2 27B PQ2_0 head 450c5c3eb 50.9 6.75
Bonsai 2 27B PTQ1_0 head 450c5c3eb 47.9 5.68

The head roughly doubles PQ2_0 decode on Daniel's laptop relative to the base. PQ2_0 now decodes faster than PTQ1_0 on this GPU rather than at half of PTQ1_0's rate.

For scale, one PQ2_0 token streams about 6.8 GB of weights, and a kernel that only reads memory sequentially reaches 108 GB/s on this GPU. That puts the single-stream decode ceiling on Daniel's laptop at about 16 tokens per second. The head sits a bit under half of that ceiling.

The per-operator test in test-backend-ops compares each matrix product against the CPU backend. On this GPU it passes 102 of 102 MUL_MAT cases for PQ2_0 and 102 of 102 for PTQ1_0. Before the type gate in 8a08dd019, the same test aborted on this machine before reaching those cases.

The PR description lists a full perplexity comparison against the CPU backend as not yet done. Claude ran a related check on full-vocabulary logits. The input was 512 WikiText prompt tokens followed by 32 decoded tokens, with prefill in FP32, and every eighth position was kept. Each kept position was scored against the CPU backend by normalized mean-square error (NMSE) and by mean KL divergence.

PTQ1_0 logits are bit-identical between the base and the head, as expected since de7a01b does not touch PTQ1_0. For PQ2_0:

PQ2_0 comparison NMSE Mean KL
base vs head, both SYCL 2.3e-5 8.1e-5
CPU backend vs head 1.4e-4 7.6e-5
CPU backend vs base 1.4e-4 7.0e-5

de7a01b changes the order in which the PQ2_0 zero point is applied, and that moves the FP32 logits slightly. This is not a correctness problem: both commits sit at about the same distance from the CPU backend, which is the reference that matters, and a gap of that size between backends is normal because the CPU backend quantizes activations differently. One consequence is worth stating. A logit dump from an earlier SYCL build will not match the head, so the CPU backend is the right reference for the unfinished check.

Further optimization on Daniel's fork

dwymark/llama.cpp#2 is based on the head. It adds opt-in explicit-SIMD decode kernels for PQ2_0 and for PTQ1_0, plus an FP16-input prefill path for both formats. Nothing changes by default; each path is enabled by an environment variable.

On Daniel's laptop, PR 2 reaches 9.0 tokens per second on PQ2_0 decode and about 145 tokens per second on prompt processing. PTQ1_0 reaches 6.9 and 135. The FP32 output of PR 2 is bit-identical to the base's arithmetic for both formats. The speed suggests the decode path on Intel GPUs still has headroom.

Daniel makes no assumption that PR 2 will be merged here, but you are welcome to merge it, adapt it, or copy from it verbatim without credit, whichever is useful.

The logit probe and scorer used in the logit check above live on PR 2.

Daniel is happy to rerun any of this at a later head, or to run the CPU-backend logit check for PTQ1_0 if that would help close the unfinished perplexity comparison.

Written by Claude Code on Daniel's instruction.

@bri-prism

Copy link
Copy Markdown
Collaborator

Tested head 450c5c3eb on an Intel Arc B390 (Panther Lake Xe3 iGPU, UMA) with a Core Ultra X7 358H, Windows 11, oneAPI 2025.3 (icx, Level Zero), merged onto prism @ 3b19c377d. I compared it with the previous head c5f6c9718, which I had tested earlier on the same machine.

❌ Crash introduced by this PR: CPY f32/f16/bf16 → PTQ1_0/PQ2_0

test-backend-ops test -b SYCL0 -o CPY -p "ptq1_0|pq2_0" aborts:

ggml-sycl/cpy.cpp:1298: GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(src1->type)) failed   (f32 -> ptq1_0 / pq2_0)
ggml-sycl/cpy.cpp:1302: ...                                                                (f16 -> ptq1_0 / pq2_0)
ggml-sycl/cpy.cpp:1308: ...                                                                (bf16 -> ptq1_0)

The PR adds PTQ1_0/PQ2_0 to ggml_sycl_is_quantized_type() (cpy.cpp:619) but not to ggml_sycl_can_quantize_rows_sycl() (cpy.cpp:652). As a result, supports_op accepts float→PTQ1_0/PQ2_0 copies, and ggml_sycl_cpy then asserts instead of falling back. Either add quantize-rows support for the two types, or reject those pairs in supports_op. ptq1_0→ptq1_0 and pq2_0→pq2_0 pass 9/9 each; ptq1_0/pq2_0 → f32 report not supported.

⚠️ Pre-existing, not this PR: SET_ROWS → PTQ1_0/PQ2_0 aborts

set_rows.cpp:549: Unsupported tensor type!. The SYCL supports_op for GGML_OP_SET_ROWS (ggml-sycl.cpp ~6115) checks only src[0]/src[1] and never the destination type, so any type missing from the set_rows switch aborts. The same code is on prism and this PR doesn't touch it. I mention it because it stops a plain test-backend-ops test -b SYCL0 -p "ptq1_0|pq2_0" run, which is what the Requirements checklist asks for. q1_0/q2_0/q4_0 SET_ROWS pass.

✅ Everything else

op (ptq1_0 + pq2_0 cases) result
MUL_MAT, incl. the new Bonsai-shape cases 292/292 on 4 of 5 runs. One run had a single tolerance miss: pq2_0 m=16 n=1 k=256, ERR 0.000524 vs 0.0005. It didn't reproduce, so I read it as random-data noise.
MUL_MAT_ID 166/166
GET_ROWS 8/8

Greedy 128-token generation (-ngl 99, temp 0) on the shipped v5 Ternary-Bonsai-2-27B PQ2_0 and PTQ1_0 is coherent on both formats.

Performance: new head vs previous head, same machine

llama-bench -ngl 99 -p 512 -n 128 -r 3, two interleaved rounds (round 1 / round 2, t/s):

model head pp512 tg128
Bonsai 2 27B PQ2_0 c5f6c97 81.65 / 74.79 5.08 / 5.05
450c5c3 83.45 / 76.89 10.07 / 9.24
Ternary-Bonsai 1.7B PQ2_0 c5f6c97 1538 / 1541 71.8 / 71.9
450c5c3 1551 / 1549 133.1 / 131.0
Bonsai 2 27B PTQ1_0 c5f6c97 85.24 / 74.62 10.40 / 9.94
450c5c3 91.45 / 75.29 11.06 / 9.95
2B PTQ1_0 c5f6c97 1207 / 1196 84.3 / 81.0
450c5c3 1195 / 1197 83.2 / 82.6

The ~2x PQ2_0 decode claim reproduces on Xe3 (27B: 5.08 → 10.07; 1.7B: 71.8 → 133). PTQ1_0 and prefill are unchanged, as expected; both heads share the same PTQ1_0 kernel. Round 2 runs a little slower for both heads on the 27B, which I read as thermal drift on the laptop.

With the CPY gate fixed, this looks good to me.

Tested with Claude Code.

Adding the two types to ggml_sycl_is_quantized_type() made supports_op
accept f32/f16/bf16 -> PTQ1_0/PQ2_0 copies, but there is no quantize-rows
kernel for either, so ggml_sycl_cpy() took the float -> quantized branch
and died on GGML_ASSERT(ggml_sycl_can_quantize_rows_sycl(...)) at
cpy.cpp:1298/1302/1308.

Neither format can be produced by a row-at-a-time GPU kernel: the packing
assumes the Hadamard rotation the offline converter applies, so the right
answer is to decline the pair and let the scheduler place it elsewhere,
not to add a kernel. Quantized -> same-quantized copies are unaffected and
still take the memcpy path.

Reported by @dwymark on an Arc B390 (Panther Lake Xe3); reproduced on an
Arc Pro B50.

  test -b SYCL0 -o CPY -p "ptq1_0|pq2_0"
    before: abort at cpy.cpp:1298
    after:  18/18 tests passed, float -> ternary reports "not supported"

No regression: MUL_MAT 204/204, MUL_MAT_ID 150/150, GET_ROWS 8/8 for the
two types.

The unrelated tq2_0 -> tq2_0 abort at cpy.cpp:1455 and the SET_ROWS abort
at set_rows.cpp:549 are both present on prism and untouched here.
@kiljoy001

Copy link
Copy Markdown
Author

Good catch, and thanks for the interleaved rounds - that is the part that makes the decode numbers believable.

CPY crash: fixed in 42ce76d

Your diagnosis was exactly right. Adding the two types to ggml_sycl_is_quantized_type() (cpy.cpp:622) made supports_op accept the float -> ternary pairs, and ggml_sycl_cpy() then took the float -> quantized branch and asserted on can_quantize_rows_sycl(). The two lists disagreed for precisely these two types.

I took your second option rather than the first. Neither format can be produced by a row-at-a-time GPU kernel: the packing assumes the Hadamard rotation the offline converter applies, so a quantize-rows kernel would be wrong rather than merely missing. The fix declines the pair in supports_op and lets the scheduler place it elsewhere:

if ((src1_type == GGML_TYPE_PTQ1_0 || src1_type == GGML_TYPE_PQ2_0) &&
    src0_type != src1_type) {
    return false;
}

Keyed on the assert's own condition rather than added to the three per-source exclusion lists, so f16 and bf16 are covered by the same guard.

Reproduced your abort on an Arc Pro B50 first, then confirmed the fix:

test -b SYCL0 -o CPY -p "ptq1_0|pq2_0"
  before: abort at cpy.cpp:1298
  after:  18/18 tests passed

float -> ternary now reports not supported; ptq1_0 -> ptq1_0 and pq2_0 -> pq2_0 still take the memcpy path. No regression on the rest: MUL_MAT 204/204, MUL_MAT_ID 150/150, GET_ROWS 8/8 for the two types.

SET_ROWS

Agreed, and thanks for checking it against prism rather than assuming. I hit the same abort at set_rows.cpp:549 on the B50. It is the same shape as the MUL_MAT abort this PR already fixes - a switch with GGML_ABORT in the default arm, reachable because supports_op never checks the destination type - but in a different op, so I would rather fix it in its own PR than widen this one.

The one tolerance miss

pq2_0 m=16 n=1 k=256 at ERR 0.000524 against a 0.0005 bound, once in five runs and not reproducing, reads the same way to me: the NMSE on that shape is right at the threshold and the test regenerates random data per run. Worth knowing it sits that close to the line though. I have not seen it on the B50 across repeated runs, so if it recurs for you on Xe3 specifically I would want to look again.

Numbers

The 2x PQ2_0 decode reproducing on Xe3 (27B 5.08 -> 10.07, 1.7B 71.8 -> 133) is the third device now, after your earlier B390 run and my B50. PTQ1_0 and prefill flat is what I would expect since de7a01b only touches the PQ2_0 kernel.

Fix is pushed. Happy to have you re-run whenever convenient.

@dwymark

dwymark commented Sep 23, 2026

Copy link
Copy Markdown

Short version of my earlier comment, since it landed mid-review and the headline got buried.

On an Intel Arc 140V, dwymark/llama.cpp#2 on top of this PR takes PQ2_0 decode from 6.75 to 9.0 t/s (1.3x) and prompt processing from 51 to 145 t/s (2.8x). PTQ1_0 goes from 5.68 to 6.9 and from 48 to 135. FP32 output is bit-identical to c5f6c97 for both formats. The branch is now rebased onto 42ce76d and fast-forwards onto it.

Two independent pieces: explicit-SIMD decode kernels for both formats, and an FP16-input GEMM for prefill. Both are off by default behind environment variables. Same stance as before: merge it, adapt it, or copy from it without credit, whichever is useful.

On the still-open perplexity item: the earlier comment has PQ2_0 full-vocabulary logits scored against the CPU backend at both c5f6c97 and 450c5c3 (NMSE 1.4e-4, mean KL under 1e-4 for both), and PTQ1_0 logits are bit-identical between those two heads. Happy to run the PTQ1_0 comparison against CPU if that would close it.

Written by Claude Code on Daniel's instruction.

@bri-prism

Copy link
Copy Markdown
Collaborator

Re-tested at 42ce76d01 on the same Arc B390 (Panther Lake Xe3), Windows 11, oneAPI 2025.3, merged onto prism @ 3b19c377d. The CPY fix is confirmed.

op (-p "ptq1_0|pq2_0") 450c5c3eb 42ce76d01
CPY abort at cpy.cpp:1298 18/18 pass; float → PTQ1_0/PQ2_0 now reports not supported
MUL_MAT 292/292 (one tolerance miss in 5 runs) 292/292
MUL_MAT_ID 166/166 166/166
GET_ROWS 8/8 8/8
SET_ROWS abort at set_rows.cpp:549 same (pre-existing on prism; fine as a separate PR)

Output: greedy 128-token generation on the shipped v5 Bonsai 2 27B is deterministic and unchanged by the fix. For PTQ1_0, 3 runs on 42ce76d01 and 3 on a rebuilt 450c5c3eb all produced identical text. PQ2_0 is byte-identical to my earlier run. On this device PTQ1_0 and PQ2_0 also produce identical text to each other, as you'd expect since both files hold the same ternary weights.

Performance: the fix only touches supports_op, so no speed change was expected. The laptop was running hot after several hours of benchmarks, so the absolute numbers are ~20% below my earlier run. Here's the relative picture, interleaved new / old / new in the same session, -ngl 99 -p 512 -n 128 -r 3, two rounds:

model 42ce76d01 tg128 c5f6c9718 tg128 ratio
Bonsai 2 27B PQ2_0 7.63 / 7.72 / 7.77 / 7.93 4.29 / 4.33 ~1.8x
Ternary-Bonsai 1.7B PQ2_0 110.5 / 97.8 / 85.6 / 99.3 52.8 / 59.9 ~1.7x
Bonsai 2 27B PTQ1_0 8.25 / 8.83 / 8.01 / 8.85 8.37 / 8.88 1.0x (same kernel)

The PQ2_0 decode gain holds up under thermal throttling. pp512 is unchanged between heads.

LGTM from the Xe3 side. One small note: the fix commit message credits the report to @dwymark on an Arc B390, but the B390 report was the review above from this account. Not important, just mentioning it in case the history matters to you.

Tested with Claude Code.

@bri-prism
bri-prism self-requested a review September 23, 2026 22:54
@bri-prism

Copy link
Copy Markdown
Collaborator

Correction to the performance section of my previous comment. I attributed the ~20% lower absolute numbers to thermal throttling. That was wrong: a CPU benchmark loop I believed I had stopped was still running on this UMA laptop, competing for memory bandwidth. Everything else in that comment stands: the CPY fix verification, the op results and the deterministic output.

Clean rerun at 42ce76d01, with a logged 0 other llama processes during the run. Interleaved new / old / new, two rounds, -ngl 99 -p 512 -n 128 -r 3, t/s:

model 42ce76d01 tg128 c5f6c9718 tg128 ratio pp512 new vs old
Bonsai 2 27B PQ2_0 11.00 / 9.22 / 9.24 / 9.05 5.06 / 5.18 ~1.8–2.1x 74–94 vs 79
Ternary-Bonsai 1.7B PQ2_0 135.8 / 130.0 / 135.3 / 135.8 73.4 / 69.8 ~1.9x 1557 vs 1516–1543
Bonsai 2 27B PTQ1_0 10.55 / 9.83 / 10.09 / 10.39 10.23 / 10.29 1.0x 74–86 vs 77–78

This matches my first review (27B PQ2_0 5.08 → 10.07 t/s, 1.7B 71.8 → 133 t/s). The LGTM stands.

Tested with Claude Code.

@andyyeh75

andyyeh75 commented Oct 1, 2026 •

Copy link
Copy Markdown

I also implemented the PQ2_0 kernel in my local repo (not branched from this PR) and can perform better in PP512. Let me try later if my approach can benefit this backend kernel to shorten prefill time further.

Model: Bonsai 2 27B PQ2_0
(standard b:2048/ub:512)

Hardware PP512 TG128
Intel PTL B390 w/ SYCL PQ2_0 kernel 248.54 ± 0.31 t/s 11.64 ± 0.02 t/s │
Refer:Apple M4 Pro / Metal 126.89 ± 0.02 t/s 20.53 ± 0.01 t/s │

imaami pushed a commit to imaami/llama.cpp that referenced this pull request Oct 13, 2026
Restore Scott Guyton's PTQ1_0/PQ2_0 cases at Bonsai row lengths, both partial row groups, and one to eight columns. These exercise the SYCL multi-column paths beyond the generic small shapes.

Source: PrismML-Eng#235

Source: PrismML-Eng@ccbf5a7
Assisted-by: GitHub Copilot

Ported and rebased by Codex.
imaami pushed a commit to imaami/llama.cpp that referenced this pull request Oct 13, 2026
Add Scott Guyton's PTQ1_0/PQ2_0 cases at Bonsai row lengths, both partial row groups, and one to eight columns. These exercise the SYCL multi-column paths beyond the generic small shapes.

Source: PrismML-Eng#235

Source: PrismML-Eng@ccbf5a7
Assisted-by: GitHub Copilot

Ported and rebased by Codex.
imaami pushed a commit to imaami/llama.cpp that referenced this pull request Oct 13, 2026
Add Scott Guyton's PTQ1_0/PQ2_0 cases at Bonsai row lengths, both partial row groups, and one to eight columns. These exercise the SYCL multi-column paths beyond the generic small shapes.

Source: PrismML-Eng#235

Source: PrismML-Eng@ccbf5a7
Assisted-by: GitHub Copilot

Ported and rebased by Codex.
imaami pushed a commit to imaami/llama.cpp that referenced this pull request Oct 13, 2026
Add Scott Guyton's PTQ1_0/PQ2_0 cases at Bonsai row lengths, both partial row groups, and one to eight columns. These exercise the SYCL multi-column paths beyond the generic small shapes.

Source: PrismML-Eng#235

Source: PrismML-Eng@ccbf5a7
Assisted-by: GitHub Copilot

Ported and rebased by Codex.
bri-prism added a commit that referenced this pull request Oct 14, 2026
…B on Intel Arc (port of Torchit1/arc-b580) (#356)

* Bonsai 2 27B (PTQ1_0) on Intel Arc B580: SYCL + Vulkan performance work

- SYCL: PTQ1_0 weights repacked to 2-bit at load and multiplied on XMX via the TernSYCL int8 x int2 DPAS kernels
  (vendored, BSD 3-Clause), activations quantized to int8 with round-to-nearest (GGML_SYCL_PTQ1_T2=ffn|all).
- SYCL: decode attention for q4_0 KV serving 1-4 query tokens per launch; blocked gated delta-net; PTQ1_0
  mat-vec kernels (single column, decode-once multi column, same-input merge, fused gate/up); narrow concat.
- Vulkan: matching PTQ1_0 multi-column, decode attention, gated delta-net and coopmat work.
- Speculative decoding: draft KV cache type (-ctkd/-ctvd) honoured by the MTP draft context,
  LLAMA_ARG_SPEC_DRAFT_UBATCH to shrink its compute buffer.
- docs/bonsai-arc-b580.md: results at 128K on a B580, build and run instructions, switches, credits.

Based on the bonsai-combo branch of professorpalmer/llama.cpp-ada-ternary (PrismML PR #221).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 753323e on Torchit1/llama.cpp arc-b580)
Excluded from this port: .github docs ggml/src/ggml-vulkan
Conflicts resolved against prism: fwht.cpp keeps prism's FWHT (#302) and adds this commit's signed
variant beside it as ggml_sycl_op_fwht_signed; ggml_sycl_op_mul_mat_sycl takes fp16 for PQ2_0/PTQ1_0
as before (#278) plus this commit's large-batch rule; the T2 path runs before the PQ2_0 XMX path (#294),
which still serves PQ2_0/PTQ1_0 when GGML_SYCL_PTQ1_T2 is unset.

* sycl: PQ2_0 weights on the TernSYCL XMX path too (GGML_SYCL_PTQ1_T2)

PQ2_0 (2-bit codes 00=-1 01=0 10=+1, fp16 scale per 128) is the same size as the 2-bit XMX layout, so it is remapped in
place. Code 3 (+2) has no 2-bit two's-complement form: the repack reports it and the matmul aborts with a message.
test-backend-ops MUL_MAT type_a=pq2_0 with GGML_SYCL_PTQ1_T2=all: 84/84; ptq1 94/94.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 5056bde on Torchit1/llama.cpp arc-b580)

* sycl: PTQ1 T2 path only on Xe2+ GPUs, fall back with a warning elsewhere

The T2 kernels use 16-lane int8 x int2 DPAS builtins that only Xe2 and newer drivers provide; on an
Arc 140T (Xe-LPG+) GGML_SYCL_PTQ1_T2=all failed with 'No kernel named ... was found'. Now checked
once against the device architecture. GGML_SYCL_PTQ1_T2_ANYGPU=1 skips the check.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 731c1a8 on Torchit1/llama.cpp arc-b580)

* sycl/ternsycl: 64-bit surface base address on Windows

The 2D block I/O builtins took the surface base as 'long', which is 32 bits in
SYCL device code on Windows (the device follows the host's LLP64 data model), so
the weight loads would have used truncated addresses. 'long long' is 64 bits on
both; identical code on Linux.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 2f0cf6e on Torchit1/llama.cpp arc-b580)

* sycl: GGML_SYCL_PTQ1_T2_TILE lab knob for the TernSYCL large-M tile (default stays best)

B580 pp1024 (ub 1024, all weights on XMX): 8x128/8x2 939 t/s (default), 8x128/16x1 949, others 480-857.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 244c4d4 on Torchit1/llama.cpp arc-b580)

* sycl: two-pass top-k for long rows (GGML_SYCL_TOPK_OLD=1 restores the one-group kernel)

The one-work-group top-k took ~550 us per MTP draft token on the B580 (248K-entry vocabulary). Now up to 256 groups
per row with a pairwise local-memory merge, then one group merges the winners; ties keep the lower index.
test-backend-ops TOP_K 325/325. 200K entries: k=1 241 -> 9.6 us, k=10 647 -> 71 us. Live settings: fresh +1-2%,
rename/edit within noise (the speculative step is not purely GPU-bound).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit afe4c82 on Torchit1/llama.cpp arc-b580)

* sycl: q4_0 decode attention on XMX via s8 x s4 DPAS (opt-in GGML_SYCL_FA_DEC_DPAS)

Q.K uses intel_sub_group_i8_i4_matrix_mad_k32 (IGC provides it on Xe2 though the
public extension only lists i4_i4_k64): the raw 16-byte q4_0 block of a key XOR
0x88888888 is the B operand, one key per lane, against int8 Q with one RNE scale
per row and block (q4_0 byte order absorbed by pairing dims l and l+16 in A).
P.V computes O^T = V^T P^T on f16 DPAS so each lane dequantizes its own key's V
block from 5 dword loads. Token chunks of 8 let one kernel serve 1..32 (or =N)
query tokens, so n-gram verify batches stop falling back to oneDNN + a full-KV
f16 conversion.

Per layer (Bonsai 27B shape, real cache layout) old -> new, us: 16K 1 tok 266 ->
123, 4 tok 457 -> 244, 16 tok 2217 -> 776; 48K 1 tok 768 -> 433, 4 tok 1295 ->
671, 16 tok 6629 -> 1915. FLASH_ATTN_EXT decode cases incl. 12/16/17/32 tokens
and the permuted layout: 57/57. GGML_SYCL_FA_DEC_DPAS_PROBE skips phases.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit e5a4d8a on Torchit1/llama.cpp arc-b580)

* sycl: MKL flash-attention softmax, one work-group per query row

mkl_fa_online_softmax_chunk gave each query row to one work-item that walked the
whole KV chunk twice, so neighbouring work-items read addresses chunk_size floats
apart. Above the oneDNN KV cap this kernel was 72% of generation time (Bonsai 27B,
115K-token context: verify batches > 32 tokens and every new prompt). Now 256
work-items per row with coalesced reads and group reductions; same math.

Bonsai 27B, live flags, 114.6K-token code context: new code 17.6 -> 44.0 t/s,
edit 10.8 -> 27.5 t/s, identical output; short prompts unchanged.
FLASH_ATTN_EXT with the MKL path forced (ONEDNN_MAX_KV=512): 82/84, the same two
pre-existing prompt-batch failures as before.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit a122b2b on Torchit1/llama.cpp arc-b580)

* sycl: GGML_SYCL_MMVQ_CHUNK_MAX, largest batch run as MMVQ column chunks

Quantized mul_mat with 33+ columns fell through to dequantize-all-of-src0 + GEMM
per call. For speculative-verify sized batches that is ruinous: Gemma 4 12B q4_0
with n-gram drafts of 33-64 tokens spent ~94% of GPU time in
dequantize_row_q4_0_sycl_reorder. Default stays 4 x MMVQ_MAX_BATCH_SIZE (32);
with 128, Gemma 4 12B + QAT assistant (ngram n-max 64): rename 79 -> 186 t/s,
edit 57 -> 136. Bonsai 27B unchanged within noise.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit e8bd534 on Torchit1/llama.cpp arc-b580)

* sycl: oneDNN / MKL flash attention only for packed interleaved K/V

Reported from an Arc 140T and reproduced on a B580: FLASH_ATTN_EXT with quantized
K/V in the permuted layout returned ERR ~1.2 for every KV length from 1024 up
(ggml-org#27769). The failing test builds the cache view with a gap between tokens
(2x the heads allocated, half viewed); the oneDNN SDPA and MKL dequant paths take
the token stride from the head count. llama.cpp's own cache view is always packed,
so real models were unaffected (Bonsai 27B perplexity at 4K: oneDNN 8.7273, MKL
8.7311, TILE 8.7330, and still 8.7273 with this guard). Interleaved K/V with a gap
now falls back to TILE. FLASH_ATTN_EXT 4137/4137 (was 2 failing), plus a new
q4_0 real-layout sweep (kv 1024..8192, nb 33/64/512) 29/29.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 226166a on Torchit1/llama.cpp arc-b580)

* sycl: XMX decode attention only on Xe2+, fall back with a warning elsewhere

GGML_SYCL_FA_DEC_DPAS uses the 16-lane int8 x int4 DPAS builtin that only Xe2 and
newer provide; on Alchemist / Xe-LPG it would fail at the first launch (like the
T2 path did before 9c25b48). Now checked once against the device architecture,
sharing ggml_sycl_device_is_xe2() with the T2 gate; GGML_SYCL_FA_DEC_DPAS_ANYGPU=1
skips the check. B580 unchanged: 4-token verify at 48K 671 us (1295 without).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit d441076 on Torchit1/llama.cpp arc-b580)

* sycl: XMX q4_0 decode attention for Gemma 4 (head 512 x GQA 16 global, head 256 x GQA 2 sliding)

The s8 x s4 DPAS decode kernel takes the head size as a template parameter (K read
block by block, P.V subgroups own D/16 dims) and serves two more shapes when
GGML_SYCL_FA_DEC_DPAS is on: Gemma 4's global layers (512, 1 KV x 16 Q heads, token
chunks of 2) and sliding layers (256, 2 per KV head, chunks of 8). Splits target
~256 work-groups in all (a single KV head used to leave the GPU idle); combine step
templated on D.

Global-layer attention at 48K, us: 1 token 1114 (f16 TILE) -> 236, 4 tokens 2011 ->
635. Gemma 4 12B QAT + assistant, q4_0 KV, 128K context: short 92 -> 100 t/s new
code; at 48K new code 47 -> 65, edit 44 -> 57 (f16 KV was 55 / 47). Bonsai shape
unchanged (4-token verify at 48K 638 us). FLASH_ATTN_EXT: Gemma global 20/20,
sliding 10/10, Bonsai 36/36.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit c759ddb on Torchit1/llama.cpp arc-b580)

* sycl: reordered q4_0 small-batch GEMM on XMX via s8 x s4 DPAS (opt-in GGML_SYCL_Q4_0_DPAS)

Verify batches of 5..256 columns on q4_0 weights used MMVQ column chunks (weights
re-read per 8 columns) or dequantize-everything + GEMM. The new kernel streams each
weight row once for up to 64 columns (chunks of 64 beyond): one lane per weight row,
B = the row's 16 q4_0 bytes XOR 0x88888888 through intel_sub_group_i8_i4_matrix_mad_k32,
A = activations quantized once per call to int8 (per row and 32-block, elements l / l+16
paired). Xe2 only (ggml_sycl_device_is_xe2). GGML_SYCL_Q4_0_DPAS_MIN / _MAX (5 / 256).

15360x3840 q4_0: 8 columns 150 -> 81 us, 16 columns 292 -> 89 us. Gemma 4 12B QAT +
assistant: rename 188 -> 275 t/s, edit 138 -> 168 (ngram n-max 96), new code unchanged,
rename/edit outputs byte-identical. MUL_MAT q4_0 Gemma shapes 1..256 columns: 20/20.
Also Gemma 4 q4_0 mat-mul perf cases.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 03594d5 on Torchit1/llama.cpp arc-b580)

* sycl: large T2 prompt batches through oneDNN's int8 GEMM (GGML_SYCL_T2_W8A8_MIN)

Before the TernSYCL path went live, PTQ1_0 prompts ran on the W8A8 oneDNN GEMM
(pp512 1013); GemmMT is better for verify-sized batches but slower on full prompt
batches. New ggml_sycl_w8a8_mul_mat_t2 converts the repacked 2-bit layout to int8
(per-row scale, group ratios) and uses the same gemm_s8. The conversion is a fixed
cost per call, so it only pays from ~1024 columns: Bonsai 27B pp1024 (ub 1024) 941 ->
1062, pp512 901 -> 818 (so MIN=1024 on the live config); 48.6K-token prompt read
67.7 -> 61.6 s (719 -> 789 t/s). PPL c4096 8.7273 -> 8.7359 (+0.1%, int8 rounding
as the old W8A8 path). Off by default (0).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit a68382d on Torchit1/llama.cpp arc-b580)

* SYCL decode attention: SLM-staged K/V tiles for 1 query token, one-wave split count; PTQ1_0 T2 NSG knob

fattn-dec (DPAS q4_0 path, D 256): single-token decode copies each 128-key K and V tile into SLM with
overlapped 16-byte loads and reads lane = key from there; the split count now fills one wave of work-groups
(Xe cores x work-groups per core / KV heads) instead of a fixed 64 per head. B580: 48K 1 token 433 -> 244 us,
4 tokens 640 -> 456 us; whole model plain decode +12% at 48K, +21% at 115K, same KLD and output.
GGML_SYCL_FA_DEC_STAGE=0 and GGML_SYCL_FA_DEC_SPLITS=64 restore the old behaviour.
ptq1-t2: GGML_SYCL_PTQ1_T2_NSG=4 (4 sub-groups for 2-8 row GEMV, +~1%, default off).

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit 9cfffcb on Torchit1/llama.cpp arc-b580)

* SYCL: s8 x s4 DPAS as inline vISA instead of the intel_sub_group_i8_i4_matrix_mad_k32 builtin

The builtin is not in the public subgroup matrix extension list. On a Lunar Lake (Arc 140V) driver
that does not provide it, loading the device image failed with "No device image found for external
symbol _Z36intel_sub_group_i8_i4_matrix_mad_k32" and every model crashed at the first decode.
xe2-dpas.hpp emits the dpas.s4.s8 instruction directly (as libxsmm/TernSYCL's xe2.hpp helpers), so
the symbol is gone from libggml-sycl. Used by the q4_0 decode attention and the q4_0 small-batch GEMM.

B580: FA 4167/4167, MUL_MAT q4_0 (GGML_SYCL_Q4_0_DPAS=1) 59/59; whole model output byte-identical to the
builtin at short / 48K / 115K context, plain and MTP. Single-token attention kernel ~6% slower
(48K 244 -> 259 us), whole model within 1-2%.

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit f31a4f6 on Torchit1/llama.cpp arc-b580)

* SYCL: no IGC builtins left (TernSYCL main, f16 DPAS via vISA), device code split per kernel

- ternsycl/: libxsmm/TernSYCL main (11484da) + the long long address fix of TernSYCL#3. Inline vISA only, so the
  undocumented builtins the old copy declared (intel_sub_group_i8_i2_matrix_mad_k32, __builtin_IB_* 2D block I/O)
  are gone; it also drops our three JIT workarounds. The per-32 activation scale variant (off by default) is removed.
- ptq1-t2.cpp is compiled with -fp-model=precise: under icpx's default fast model the TernSYCL GemmMT returns NaN
  (upstream's own driver fails validation built that way).
- fattn-dec: the f16 P.V DPAS also goes through inline vISA (xe2-dpas.hpp dpas_hf_r8).
- CMake: -fsycl-device-code-split=per_kernel (GGML_SYCL_DEVICE_CODE_SPLIT_PER_KERNEL, default ON): the Xe2 XMX
  kernels are only JIT-compiled when launched, so GPUs without DPAS (Xe-LP, Gen9) don't compile them at all.
- test-backend-ops: FA cases for 64-256 query tokens and ~115K history.

B580: libggml-sycl has 0 matrix_mad / __builtin_IB references (was f16, i8_i2, 9 IB). FA 4167/4167, MUL_MAT q4_0
59/59, ptq1_0 T2 94/94. KLD (ub 1 / 4 / 512) identical to the previous build. Whole model output identical; MTP
48K 73.6 t/s, 115K 59.6 (was 69.6 / 57.2), plain unchanged. Startup unchanged (first request 4 s).

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit 2e8a0e6 on Torchit1/llama.cpp arc-b580)

* SYCL decode attention: default DPAS cap 32 -> 256 query tokens

Long n-gram verify batches now stay on the q4_0 DPAS kernel instead of oneDNN's path, which converts the whole KV
cache to f16 each call. B580 per layer: 48K 64 tokens 28.8 -> 4.9 ms, 115K 256 tokens 253 -> 44 ms. KLD (ub 64-256)
0.012699 vs 0.012800, same top p 92.60 vs 92.89 %, equal within error. Whole model MTP: 48K 73.2 -> 76.1 t/s,
115K 60.3 -> 65.1. GGML_SYCL_FA_DEC_DPAS=32 restores the old cap.

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit 064461a on Torchit1/llama.cpp arc-b580)

* SYCL decode attention: 5-token query chunks for MTP verifies

MTP with 4 drafts verifies 5 tokens per step, which ran the 8-token DPAS kernel with 3 empty tokens (48 rows of
Q.K / softmax for 30). A 5-token chunk (128-key tiles) cuts the kernel from 637 to 502 us per layer at 48K (215 ->
175 us at 16K); whole model MTP 48K 77.4 -> 78.1 t/s, 115K 65.7 -> 67.1, output identical. GGML_SYCL_FA_DEC_NQ5=0
turns it off (=2 SLM-staged and =3 256-key tiles were slower). test-backend-ops: 5/6-token real-layout cases.

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit 081d868 on Torchit1/llama.cpp arc-b580)

* lab: runtime-switchable SYCL tuning knobs (GGML_SYCL_KNOBS_FILE); LLAMA_ARG_SPEC_NGRAM_FRESH

- ggml_sycl_knob(name, def): the attention (DPAS cap, staging, 5-token chunks, split count, probe) and PTQ1_0 T2
  (tile, sub-groups, GEMM off) switches are read through it. With GGML_SYCL_KNOBS_FILE set, the NAME=VALUE file is
  re-read at the start of a graph compute when it changes, so one server can A/B variants. Without it, each knob is
  read from the environment once, as before.
- LLAMA_ARG_SPEC_NGRAM_FRESH=1: the n-gram drafter forgets earlier requests at the start of each one, so repeated
  benchmark prompts don't replay earlier answers as drafts.

Co-Authored-By: Claude <noreply@anthropic.com>

(cherry picked from commit 9db9408 on Torchit1/llama.cpp arc-b580)
Excluded from this port: common

* sycl: T2 quantized-activation cache + fused Hadamard rotate/quantize

Back-to-back T2 mat-vecs on one tensor reuse its int8 activations
(GGML_SYCL_PTQ1_T2_QCACHE), and the signed 1024-point FWHT that feeds them
runs 128 lanes per block and writes Aq/SA itself (GGML_SYCL_PTQ1_T2_FWHTQ).
B580 Bonsai 27B plain decode +3.6-4.9%, KLD unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 254dfb2 on Torchit1/llama.cpp arc-b580)

* sycl: fuse residual add + RMS norm into the T2 rotate/quantize kernel

ADD + RMS_NORM + MUL(weight) [+ MUL(signs) + RESHAPE + Hadamard MUL_MAT]
runs as one work-group per token (GGML_SYCL_PTQ1_T2_NORMQ). Outputs only
used inside the chain are not written; across tokens, outputs must not
partly overlap inputs. B580 Bonsai 27B plain decode +1.9-2.6%, KLD unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 4d50f84 on Torchit1/llama.cpp arc-b580)

* sycl: vendored TernSYCL -> 8b0ae42, drop the -fp-model=precise override

Upstream fixed the GemmMT row-scale broadcast that returned NaN under icpx's
default -fp-model=fast (libxsmm/TernSYCL#1). B580: MUL_MAT ptq1_0/pq2_0
178/178 with fast math, KLD unchanged, decode speed unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 546f94c on Torchit1/llama.cpp arc-b580)

* sycl: lab knob GGML_SYCL_T2_ZERO_WEIGHTS (zeroes T2 weights to profile non-weight cost)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 2dbd839 on Torchit1/llama.cpp arc-b580)

* sycl: lab instrumentation - host token timing (GGML_SYCL_HOST_TIMING), W8A8 threshold as a runtime knob; test shapes

GGML_SYCL_HOST_TIMING=1 splits decode graphs into between / enqueue / wait host time
plus per-op enqueue cost, without a tracer. test-backend-ops: Bonsai attn k/v and
output-head mat-vec shapes, FA nh=1,2 at 48K.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit ec71714 on Torchit1/llama.cpp arc-b580)

* sycl: fuse the DeltaNet conv chain for one-token decode (GGML_SYCL_CONV_CHAIN)

CONCAT + conv-cache CPY + SSM_CONV + SILU + L2_NORM in one kernel per layer
(one work-group per 128-channel head). Outputs are written at the CONCAT's
position, so nodes in between must not touch them (checked); an output may sit
exactly on the qkv input (same element per channel), and an L2 output on top of
the old conv state falls back to its own L2_NORM kernel. Also
GGML_SYCL_FAST_SET_TENSOR (opt-in pinned input uploads).
B580 Bonsai 27B plain decode +1.3-1.8%, KLD ub1 unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit b56951e on Torchit1/llama.cpp arc-b580)

* sycl: fuse the DeltaNet gate block (GGML_SYCL_DN_GATES)

bf16 alpha + beta mat-vecs (2 x 48 rows, same input, up to 8 tokens) with the
dt-bias add, softplus, ssm_a multiply and sigmoid in one kernel: 9 graph nodes /
7 launches per layer -> 1 (2 when an output aliases the input and goes through
a scratch buffer). GGML_SYCL_DN_GATES_DEBUG=1 prints why a block is not fused.
B580 Bonsai 27B plain decode +0.5-1.4%, KLD unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 38c6be5 on Torchit1/llama.cpp arc-b580)

* sycl: same-input T2 mat-vecs in one launch (GGML_SYCL_T2_MULTI)

q+k+v, GDN qkv+z and FFN gate+up (<= 8 tokens) share one quantization and one
GEMV launch: TernSYCL's Gemv body now takes its work-group indices (run()), and a
wrapper maps each group to its matrix through a single call site (three inlined
copies spilled: 34.6 vs 46.3 t/s). Grouping and output staging as the PTQ1_0
merge. B580 Bonsai 27B plain decode +1.8-2.0%, KLD unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 31ae86e on Torchit1/llama.cpp arc-b580)

* sycl: fold FFN SwiGLU and the DeltaNet pre-norm into the T2 rotation

SWIGLU -> MUL(signs) -> RESHAPE -> Hadamard MUL_MAT: the rotate/quantize kernel
computes silu(gate) * up on its load (GGML_SYCL_GLU_ROT). ADD + RMS_NORM + MUL
whose rotation comes ~9 nodes later (DeltaNet layers): the rotation is computed
with the norm when no node in between touches its output (GGML_SYCL_NORM_ROT_FAR).
Same arithmetic, byte-identical outputs. B580 Bonsai 27B plain +0.3-0.8%, MTP
rename +2%, edit +1.1%.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit cf4b155 on Torchit1/llama.cpp arc-b580)

* sycl: per-batch-size GemmMT tile for T2 (small tile up to 32 rows); GEMV_MAX knob

arctune sweep on the B580: tile {8,128,2,4} is 7-11% faster per pass for 12-32 rows
(n-gram verify batches), the default {8,128,8,2} from 48 rows up. GGML_SYCL_PTQ1_T2_TILE
now forces one tile for all sizes; GGML_SYCL_PTQ1_T2_TILE_SMALL_MAX (default 32) picks the
crossover. GGML_SYCL_PTQ1_T2_GEMV_MAX (default 8) makes the GEMV/GemmMT switch a knob;
the sweep confirmed 8. Outputs and KLD identical; MTP rename/edit +0.5%, 48K copy +1.3%.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit a70fe64 on Torchit1/llama.cpp arc-b580)

* sycl: decode attention prefetches the next phase's key rows (GGML_SYCL_FA_DEC_PF)

Phase probes (GGML_SYCL_FA_DEC_DPAS_PROBE) showed Q.K and P.V serialised behind
the K / V loads. Each tile now prefetches its V rows before Q.K and the next tile's
K rows before P.V (sycl_ext_oneapi_prefetch, L1). Per layer at 1 token: 48K
263 -> 226 us, 115K 554 -> 496 us. Whole model: plain 48K +2.7%, 115K +4.4%;
MTP 48K +4.2%, 115K +7.5%; outputs identical. FA Bonsai cases 5/5.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit dc575f4 on Torchit1/llama.cpp arc-b580)

* tests: arctune fusion tests for the SYCL fusions; fusion guard fixes

test-backend-ops ARCLAB_GLU_ROT / ADD_NORM_ROT / CONV_CHAIN / DN_GATES / T2_MULTI build
the exact node pattern each fusion matches, at Bonsai 27B shapes and 1-16 tokens, 3
seeds each, and compare every fused output with the CPU backend: 75/75 pass, each
fusion confirmed firing (GGML_SYCL_FUSION_LOG, GGML_SYCL_T2_MULTI_LOG).
Found while writing them: the far norm->rotation fusion now requires the signs to be a
weight (they are read early); the conv chain no longer refuses a later reader of the
cache slot; the same-input merge repacks weights on first use, so it also merges on the
first graph. Whole model unchanged (plain short / 48K / 115K md5s identical).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit bfb2d76 on Torchit1/llama.cpp arc-b580)
Conflict resolved against prism: test-backend-ops keeps prism's fused SwiGLU/sign-flip FWHT cases and adds these.

* ptq1-t2: GGML_SYCL_PTQ1_T2_TUNED per-shape GEMV tile table (plumbing, empty table)

ARC-LAB: GGML_SYCL_PTQ1_T2_TUNED=1 (default 0) looks up a (N, K, rows, merged) row in t2_tuned_table() for every
1..8-row T2 GEMV (t2_dispatch and ggml_sycl_t2_mul_mat_multi) and launches TernSYCL's Gemv<SGM, NSG, LS, U> (TernSYCL,
BSD 3-Clause) under a GRF mode and work-group order through GemvTuned (= gsw::Sweep of /lab/exp/gemv-sweep). Tiles
are compiled from the T2_TUNED_KERNELS X-macro; t2_precompile warms them when the knob is on. Table and kernel list
are empty until the B580 sweep results are in, so =1 changes nothing. =2 is a plumbing self-test: the default tiles
through the tuned launcher (outputs must stay identical).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit a0d82ad on Torchit1/llama.cpp arc-b580)

* ptq1-t2: GGML_SYCL_PTQ1_T2_TUNED table from the B580 GEMV sweep (bit-exact tiles)

ARC-LAB: five rows from /lab/exp/gemv-sweep (JIT build, medians of 5 interleaved rounds), each with the same LS as the
default tile, so the outputs are bit-identical: ffn_down n=5 0.775, gate+up n=5 0.896, output head n=5 0.882,
q+k+v n=5 0.923, q+k+v n=1 0.974 (time vs the default tile). Tiles are TernSYCL's Gemv (BSD 3-Clause) under a different
launch geometry and work-group order. Batch sizes 2-4 and other shapes keep the default tiles. Knob still defaults to 0.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 4bd894b on Torchit1/llama.cpp arc-b580)

* ptq1-t2: tuned GEMV tiles on by default on a B580 (GGML_SYCL_PTQ1_T2_TUNED=0 turns them off)

B580 A/B (sd2 head): MTP-only +2.9/+3.1%; live MTP + n-gram short +4.4%, 48K +3.1%, 115K flat; outputs identical.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 0729c9d on Torchit1/llama.cpp arc-b580)

* sycl: default to the legacy Level Zero adapter unless SYCL_UR_USE_LEVEL_ZERO_V2 is set

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 5336e12 on Torchit1/llama.cpp arc-b580)

* sycl: tuned Q6_K mat-vec and fused Q6_K gate+up+SWIGLU for the MTP draft head (knobs, off)

The draft head's Q6_K mat-vecs are ~60% of a draft step and read at ~70% of
the B580's bandwidth (op timing, 2026-10-10).

- GGML_SYCL_Q6K_GEMV=1: a single-column reordered Q6_K mat-vec with the
  original per-row arithmetic (same lane -> quant-int mapping, blocks in order,
  same vec_dot expression accumulated in the same order, same sub-group
  reduction) but U blocks' operands loaded before use, aligned 32-bit ql/qh
  loads, and GGML_SYCL_Q6K_SG sub-groups per work-group. Sweep knobs
  GGML_SYCL_Q6K_U (1/2/4, default 2) and GGML_SYCL_Q6K_SG (4/8/16, default 8).
  GGML_SYCL_Q6K_CHECK=1 reruns the original kernel per launch and counts rows
  that differ in any bit (stderr), the in-server identity check.
- GGML_SYCL_Q6K_GLU=1: Q6_K joins the reorder gate+up+GLU fusion (1..8
  columns), whose dot products are the reorder mat-vec's and whose epilogue is
  the SWIGLU kernel's silu(gate) * up.
- test-backend-ops: Q6_K mat-vec and gate+up fusion cases at the head's shapes
  (eval: 1, 2, 6 tokens; perf: 1..5 tokens).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 6dd4481 on Torchit1/llama.cpp arc-b580)

* sycl fattn-dec: transposed 2D block loads for the DPAS kernel's lane = key reads (GGML_SYCL_FA_DEC_B2D, default off)

xe2dp::ld2d_t32_8x16 (lsc_load_block2d d32.8x16tn): one message gives each lane 8 dwords of its own key row, replacing
the per-lane dword gathers in Q.K (5 messages per 16 keys instead of 40) and P.V (1 instead of 5). Same bytes and math,
89/89 FLASH_ATTN_EXT cases. Engages only on 64-byte aligned K/V surfaces (logged once). Slower on its own at 128K
(register spills): 5 tokens 1069 -> 1220 us, 16 tokens 2945 -> 3610 us. Kernel launch now goes through a named lambda.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit fb75ee9 on Torchit1/llama.cpp arc-b580)

* sycl fattn-dec: 256-register file for the DPAS decode kernel at 3+ query tokens (GGML_SYCL_FA_DEC_GRF, default 256)

The 3+ token launch is one wave of 2 work-groups x 16 subgroups per Xe core = 32 threads, which is what the 256-register
mode allows, so it costs no occupancy and removes the spills. 1-2 tokens keep the default register file (their wave puts
4 work-groups on a core). GGML_SYCL_FA_DEC_GRF=0 = off. B580, us per layer at 128K (real cache layout) 5 / 8 / 12 / 16
tokens: 1073 / 1434 / 2920 / 2939 -> 1004 / 1240 / 2465 / 2445; 1 token unchanged. Whole model (arctune x3, MTP + n-gram,
sd3 head): 48K 106.1 -> 108.2 t/s (+2.0%), 115K 86.8 -> 90.6 (+4.4%), output md5 identical (d09ea3 / 291b0b).
FLASH_ATTN_EXT q4_0: 480/480. Perf cases: kv 16K/64K/128K x 1/5/8/12/16 tokens.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

(cherry picked from commit 1eb2307 on Torchit1/llama.cpp arc-b580)
Also carries the FLASH_ATTN_EXT perf cases this commit widens, first added by
1fa26f1 (Jesse Symons), which is not picked separately.

* sycl: one in-place layout owner for PQ2_0/PTQ1_0 weights

The port's MMVQ reorder (opt_for_reorder, also called by the fused gate/up path) and the
PQ2_0 XMX rewrite (#294) both change these weights in place, and neither checked the
other's flag. With GGML_SYCL_PTQ1_T2 unset a weight could get both layouts, so 27B PTQ1_0
perplexity went to 139615 on the B390 while every op test passed.

On 16-wide DPAS devices the XMX layout now owns PQ2_0/PTQ1_0 and the reorder skips them;
the fused paths that need the reorder decline and fall back to the unfused XMX path.
The XMX path also refuses a weight that is already reordered or T2-repacked.

* sycl: drop PQ2_0/PTQ1_0 duplicates left by the arc-b580 merge

Keep prism's dequantize, vec_dot and MMVQ dispatch; keep the port's reorder-aware
convert/MMVQ cases. Prism (#235/#294) and the port both added PQ2_0/PTQ1_0 support, so
the merged tree had duplicate dequant functions, switch cases and two MMVQ
implementations with the same names and did not compile. Deletions only, found and
built on the B390 (oneAPI 2025.3.2).

* sycl: Xe3 defaults to the MMVQ reorder for PQ2_0/PTQ1_0 and to XMX decode attention

Measured on an Arc B390 (Panther Lake) with GGML_SYCL_PTQ1_T2 unset:

- PQ2_0/PTQ1_0 weights go to the MMVQ reorder layout instead of the XMX
  layout: decode +16%, prompt +6%, about 0.9 GB less memory; KLD vs prism
  6.7e-4 (XMX 2.4e-4). GGML_SYCL_DISABLE_XMX=0 restores the XMX layout,
  =1 still forces the reorder on any device.
- GGML_SYCL_FA_DEC_DPAS defaults to 1: decode at 16K context +4-6% with a
  q4_0 KV cache; FLASH_ATTN_EXT head size 256 264/264, KLD with and
  without it within run-to-run error. =0 turns it off.

Other devices keep their previous defaults.

---------

Co-authored-by: Jesse Symons <155509533+Torchit1@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants