sycl: add Prism Q1_0 and Q2_0 g128 support - #94
Draft
hancris wants to merge 48 commits into
Draft
Conversation
* Optimized arm NEON(+DOTPROD) q1 dot * Implemented arm I8MM nrc==2 for q1 dot * Applied copilot advice about feature guards for Q1 Arm LUTs
* Implemented ARM NEON DP q1 4x4 repack * Hoisted out scaling by b_d in gemm * Added 4x8 NEON I8MM repack kernels * Cleanup for q1 arm repack * Added missing aliases for arch fallback * Corrected unused var statements * Extended table guard condition to account for i8mm w/o dp build Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --------- Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Q1_0/Q2_0 had no x86 vec_dot path (arch-fallback routed the generic functions to a scalar loop). Add an AVX-512-VNNI/AVX-512VL fast path guarded by __AVX512VNNI__ && __AVX512VL__, scalar fallback otherwise: - helper ggml_hsum_i32_8_vnni to reduce _mm256_dpbusd_epi32 accumulators - Q1_0: build a sign mask from the bit field, blend +qy/-qy, accumulate with dpbusd(ones, sel) - Q2_0: vectorized 2-bit unpack (replicate-4 + 16-bit shift/mask + pack), then dpbusd(codes, qy) - dpbusd(ones, qy) = sum((code-1)*qy) Q2_0 prefill ~3.9x / decode ~3.0x vs scalar on EPYC 9655; Q1_0 ~parity (the +/-1 scalar loop already auto-vectorizes). Bit-exact vs scalar (test-quantize-fns + standalone unit test); KL-divergence vs FP16 unchanged between scalar and VNNI builds. Co-authored-by: Brian <brian@Brians-MacBook-Pro.local>
The Q1_0/Q2_0 VNNI work landed in the generic quants.c, but on x86 the generic Q1_0 path is dead (arch/x86 has an AVX2 ggml_vec_dot_q1_0_q8_0 that wins), and Q2_0 only reached x86 via an arch-fallback alias. - arch/x86/quants.c: add ggml_vec_dot_q2_0_q8_0 (AVX-512-VNNI + scalar fallback), reusing the existing hsum_i32_8 helper. Math is unchanged. - arch-fallback.h: drop the x86 ggml_vec_dot_q2_0_q8_0 alias so x86 uses the arch impl, mirroring how q1_0 is already wired. - quants.c: restore portable scalar for q1_0/q2_0 generic (the VNNI in the generic file was x86-only and is now in arch/x86).
Adds an experimental tensor-core path for Q1_0 mul_mat at batch >= 128: activations quantized to int8 with per-128 absmax scales, weights repacked once per tensor to dense sign-bit words, dequant-in-SMEM via branchless SIMD unpack feeding 64x64x32 int8 wgmma, exact per-block fp32 scaling on the accumulator drain. Hybrid dispatch: persistent stream-K grid for starved shapes, fixed tile grid otherwise. Opt-in at build time (-DGGML_CUDA_HOPPER_Q1=ON -DGGML_CUDA_CUTLASS_DIR=...) and at runtime (env GGML_HOPPER_Q1); falls through to stock MMQ otherwise. Measured on H100 SXM (1-bit test model): pp512 +8.3%, pp2048 +8.6% vs stock MMQ; test-backend-ops MUL_MAT q1_0 43/43; logit-KLD vs stock path 0.0048 mean (noise-level).
Kernels templated on weight width (1- or 2-bit dense fields); Q2_0 adds a per-tensor dense repack of the 2-bit (q-1) fields and a branchless SIMD unpack (per-byte q - 1 via __vsub4, all four field values handled). Same gating, dispatch, and activation-quant path as Q1_0. Measured on H100 SXM (ternary test model): pp512 +7.7%, pp2048 +8.1% vs stock MMQ; test-backend-ops MUL_MAT q1_0+q2_0 86/86; logit-KLD vs stock path 0.0013 mean (noise-level).
- repack cache: key on (device, wdata, N, K, wbits) with a mutex (ggml may dispatch from one host thread per device), and publish entries only after a one-time stream sync so consumers on other streams cannot observe uninitialized dense buffers - per-device attr_set / SM count (a second GPU previously skipped the dynamic-SMEM opt-in and inherited device 0's SM count) - CUDA_CHECK on allocations and attribute calls - defensive int8 clamp in the activation quantizer (unreachable in exact arithmetic; insurance against fp rounding at the boundary) - CMake: target-scoped compile definitions/includes/options instead of directory-wide; option help text covers Q2_0
The path's kernels are sm_90a-only fatbins (wgmma does not exist on Blackwell); cc >= 900 alone admits cc 1200+ where the launch fails. Blackwell support is a separate tcgen05 path.
… version (PrismML-Eng#50) The windows-cpu pack step hardcoded VC\Redist\MSVC\14.44.35112, which broke when the windows-2025 runner image moved to a newer MSVC (arm64 job failed, fail-fast cancelled x64). Glob the VS product and redist version directories and pick the newest match so runner-image updates stop breaking the release. Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
The iOS job configures with LLAMA_BUILD_TOOLS=OFF, but LLAMA_BUILD_APP defaults to ON, so the llama-app target builds without the tools include paths and fails on '#include "build-info.h"'. Upstream's build-apple.yml passes -DLLAMA_BUILD_APP=OFF for the same reason. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ML-Eng#51) * kv-cache: add optional per-channel K-cache mean-centering (Q4_0 only) GGML_TYPE_Q4_0 is a symmetric quantizer, so a K channel with a real, consistent nonzero mean across tokens wastes dynamic range encoding that constant bias. This adds an opt-in mechanism that subtracts a fixed per-(kv-head, channel) bias from K right before it is written into the cache in llama_kv_cache::cpy_k(), gated strictly to k->type == GGML_TYPE_Q4_0. Subtracting the same bias from every cached key is exactly softmax-invariant: it adds the same constant (q . k_bar) to every logit in a query's row, which softmax does not see. Nothing else in attention needs to change, so this is a zero decode-time-cost quantization-fidelity improvement. llama_context_params gets a new path_kv_mean_center field (default NULL); llama_init_from_model() hard-rejects it when the K cache type isn't Q4_0, matching the existing convention for other cache-type-gated mismatches (e.g. "V cache quantization requires flash_attn"). llama_kv_cache::load_kv_mean_center() loads the bias tensors from a GGUF file and applies them; a require_q4_0 escape hatch (used only by tests) exists to validate the underlying math against an unquantized cache without confounding it with real quantization error. Also tags the K tensor right before cpy_k() with the existing cb() graph-build hook ("k_cache_in"), so calibration tooling can capture exactly the tensor that gets written into the cache regardless of what RoPE/rotation preprocessing a given architecture applies upstream. * common: add --kv-mean-center flag and GGUF bias-file writer Adds the CLI-facing side of K-cache mean-centering: common_params gains kv_mean_center_path (plumbed into llama_context_params via common_context_params_to_llama), and --kv-mean-center takes a path to a bias file generated by tools/kv-mean-center. The bias file format is a small GGUF file with one F32 tensor per layer, named "kv_bar.blk.<il>.k", holding n_embd_head_k(il) * n_head_kv(il) values laid out as [n_embd_head_k, n_head_kv]. The writer lives in common/kv-mean-center.* so it can be shared between the calibration tool and the test suite (which needs to synthesize a bias file to check the underlying math). * tools: add llama-kv-mean-center calibration tool New tool, following the tools/imatrix convention: loads a model, runs a plain text calibration corpus through it in chunks, and captures the "k_cache_in" tensor tagged in llm_graph_context::build_attn() via the same backend-scheduler eval-callback mechanism llama-imatrix uses to capture activations (params.cb_eval). The per-(head,channel) mean across all calibration tokens is written out as a bias file consumable by --kv-mean-center. * tests: add K-cache mean-centering regression + invariance tests Uses the same tiny-synthetic-model machinery as test-llama-archs.cpp (llama_model_saver + llama_model_init_from_user with a deterministic random tensor initializer), trimmed to plain LLM_ARCH_LLAMA, to check: - regression safety: two independent contexts with centering disabled produce bit-for-bit identical logits, and a Q4_0 K cache with no bias file loaded still decodes normally (cpy_k()'s new code path is a true no-op when k_bar is empty). - the --kv-mean-center gate: a non-Q4_0 K cache with a bias file is hard-rejected by llama_init_from_model(), while Q4_0 succeeds end-to-end through a real decode. - the softmax-invariance argument itself, against an unquantized F32 K cache with a synthetic nonzero bias applied through the exact same cpy_k() code path (bypassing the Q4_0 gate via load_kv_mean_center()'s require_q4_0=false test seam): output logits match the uncentered baseline to fp32 rounding (nmse ~5e-10 in practice), confirming the math without confounding it with real quantization error. * docs: document K-cache mean-centering Explains the technique, the softmax-invariance argument, usage, the bias file format, and the current scope/limitations (Q4_0-only, plain KV cache only, standard dense/GQA attention path only). * kv-mean-center: address review feedback - kv-mean-center.cpp: the K tensor captured at "k_cache_in" can be F16 or BF16 depending on backend/compute settings, not just F32. The collector was asserting F32-only and reinterpreting raw bytes as float*, which either aborts or silently computes a garbage mean on non-F32 backends. Now accepts F32/F16/BF16 and converts to F32 via ggml_fp16_to_fp32_row/ggml_bf16_to_fp32_row before accumulating. - common/arg.cpp: wire --chunks into the LLAMA_EXAMPLE_KV_MEAN_CENTER example set so the flag the tool's own README documents is actually available, instead of being silently filtered out by the shared arg parser. - docs/kv-mean-center.md: replace non-ASCII "~=" and "." characters (was U+2248 and U+00B7) with ASCII equivalents, matching the project's ASCII-only docs convention.
Generates a calibration corpus from the model itself via a temporary llama-server, removing the need for an external calibration text file. Includes a degenerate-output guard based on gzip compression ratio. Validated end to end: the resulting bias agrees with one calibrated on a standard multi-domain calibration set to within sampling noise.
Print only the header comment as help text instead of grepping every comment line, exit nonzero on unknown options, add gzip to the dependency preflight and the Requires line, and stop forcing -ngl 99: the server's own --n-gpu-layers default now applies unless -g is given.
…r-selfgen-corpus kv-mean-center: add make-calib-corpus.sh self-generated corpus helper
…composes with K rotation) (PrismML-Eng#53) * kv-cache: support mean-centering on hybrid-memory models Hybrid (recurrent + attention) models keep a standard llama_kv_cache for their attention sublayers, but llama_init_from_model only accepted path_kv_mean_center when the whole memory module was that cache, so any hybrid model failed to load a bias file. Route the load through get_mem_attn() for hybrid memory; bias tensors are matched by model layer id and layers absent from the attention cache are skipped, which the loader already handles. * kv-mean-center: document that centering must not be combined with the K-cache rotation Measured end to end, the pre-rotation bias applied post-rotation is worse than either feature alone; strengthen the README note into a warning with the measured numbers. * kv-mean-center: record the calibration basis and reject a rotation mismatch at load The bias lives in the basis the calibration run's K cache used: the collector taps the exact tensor cpy_k() writes, after any Hadamard rotation, and the rotation is gated on the K cache being quantized. A calibration run with the default F16 cache therefore measures the unrotated basis, and applying that bias to a rotated cache measurably degrades quality instead of improving it (KLD vs F16 cache 0.00144 rotation alone vs 0.0020 with the mismatched combination), while a bias calibrated with -ctk q4_0 composes (0.00111, the best of all measured configurations). The tool now detects whether the rotation was active from the captured tensor's ancestry, records it in the output file as kv_mean_center.k_rot, and load_kv_mean_center() rejects a bias whose basis does not match the inference-time rotation state (files predating the flag load with a warning). Docs updated with the calibrate-with-matching-settings rule and the measured numbers. * kv-mean-center: address review feedback Cover SWA memory layouts: the bias load now also routes through llama_kv_cache_iswa and llama_memory_hybrid_iswa (base and SWA sub-caches), so hybrid models with sliding-window attention are no longer rejected. Validate the kv_mean_center.k_rot metadata type before reading it, so a malformed bias file produces a loader error instead of an assertion abort. Run the Q4_0 cache-type validation before the basis check, so an unsupported cache type keeps its actionable error even when the file's basis would also mismatch, and hoist it out of the tensor loop. The gate test now uses a basis-matching file for the F16-cache case so it exercises the cache-type gate specifically.
The Q2_0 2-bit quant type had no entry in test-quantize-fns's per-type error-threshold tables, so it fell through to the default thresholds (MAX_QUANTIZATION_TOTAL_ERROR = 0.002, MAX_DOT_PRODUCT_ERROR = 0.02) that are tuned for 4-bit-and-up formats. A 2-bit format cannot meet those, so the test failed deterministically on every platform (absolute error 0.008678 > 0.002, dot-product error 0.141111 > 0.02). Q2_0 stores one fp16 scale per 128-element block with no zero-point, which puts its error in the same band as the ternary formats (tq1_0/tq2_0 measure 0.008681 / 0.141345 and pass at 0.01 / 0.15). Add matching Q2_0 thresholds (0.01 absolute, 0.15 dot product) rather than loosening the shared 2-bit k-quant constant, so the Q2_K / IQ2_S checks are unaffected.
…fns-q2_0-thresholds tests: add Q2_0 error thresholds to test-quantize-fns
…1_0/Q2_0 dp4a fix (PrismML-Eng#55) * speculative: dspark block-diffusion drafter + CUDA Markov resample; Q1_0/Q2_0 dp4a fix Adds the dspark speculative-decoding drafter and two low-bit/decode improvements. dspark drafter (common/speculative.cpp, src/models/dspark.cpp): - EAGLE-style block-diffusion drafter that reuses a multi-layer target-hidden-state tap (reusable capture path, also useful for EAGLE3-proper) and drafts a block of tokens per round. - Per-round sequential Markov resample: step_logits[k] = base_logits[k] + markov_w2(markov_w1(prev_token)), argmax, chaining the sampled token forward (never batched over the block). Host scalar path by default, optional host BLAS path (LLAMA_DSPARK_MARKOV_BLAS). - GGUF arch scaffolding, converter stub, and forward-graph/loop tests. See docs/dspark-scope.md for scope and the gating rationale. CUDA device-side Markov resample (common/dspark-markov.cu/.h): - Moves the sequential per-position resample onto the GPU: one H2D of the round's base logits, then a fused GEMV + add-base + argmax kernel per position that chains through a device-resident prev token, plus a final reduction. Self-contained (CUDA runtime only). Token-identical to the host scalar/BLAS path. Default when built with CUDA; opt out with LLAMA_DSPARK_MARKOV_CUDA=0. cuda: defer Q1_0/Q2_0 dp4a symbol correction (ggml/src/ggml-cuda/vecdotq.cuh): - vec_dot_q1_0_q8_1 / vec_dot_q2_0_q8_1 (the mul_mat_vec_q decode path) built a signed symbol per element before dp4a. Both now dp4a on the raw unsigned code/bit and apply one deferred affine correction at the end using Q8_1's stored real-valued activation sum (ds.y), matching the pattern vec_dot_q4_0_q8_1_impl already uses: Q1_0 dot = d*(2*sumi*ds.x - ds.y), Q2_0 dot = d*(sumi*ds.x - ds.y). Correctness: test-backend-ops MUL_MAT passes for both types. * cuda: forward-declare ggml_cuda_mul_mat_q1_hopper (fix -Werror=missing-declarations) Pre-existing on prism: the Hopper Q1 entry point is defined in mmq-hopper-q1.cu but only forward-declared locally in ggml-cuda.cu, so the definition's translation unit has no prior declaration and -Werror= missing-declarations breaks the cuda build. The full build matrix only runs on PRs (not prism pushes), so this stayed latent. One-line forward declaration; no behavior change. * tests: fix dspark test cross-platform builds (macos -Werror, windows DLL link) Surfaced by the full CI matrix (only runs on PRs, not prism pushes): - macos clang -Werror,-Wmissing-noreturn: the test-local fail() helpers never return; mark them [[noreturn]] (test-dspark-forward/loop/real-eval). - x64-windows-llvm link error: test-dspark-forward used the llama_model::get_tensor MEMBER function, which is not reliably exported across the Windows DLL boundary. Switch to the exported free function llama_internal_get_tensor_map (same pattern test-quantize-stats already uses cross-platform). No behavior change. Verified locally: all three targets build clean under clang (CPU build). * tests: use public vocab API for dspark n_vocab (windows DLL link) The prior fix swapped llama_model::get_tensor for llama_internal_get_tensor_map, but that free function is also not LLAMA_API-exported (its only other user, test-quantize-stats, is gated NOT WIN32, so it was never windows-linked) and still fails to resolve across the Windows DLL boundary. Use the exported public API instead: llama_vocab_n_tokens(llama_model_get_vocab(model)). The dspark converter's set_vocab() fills the 'none' tokenizer with dummy entries sized to the target's real vocab width, so n_tokens() reports the correct value (the old code comment claiming it is 0 predates that converter behavior). Verified locally: all three dspark test targets build clean under clang. * dspark: widen capture-copy size math to size_t (CodeQL overflow) CodeQL flagged the two dspark capture-copy sites in llama-context.cpp: 'row' was uint32_t, so n_tokens*row and n_outputs*row (feeding the byte size n_*row*sizeof(float)) were evaluated in 32-bit before widening to size_t. No overflow at real capture configs (row=n_capture_layers*n_embd is small), but latent for large captures. Make 'row' size_t so all downstream size math is 64-bit. No behavior change. libllama builds clean. * llama-context: bound capture-layer writes and zero unpopulated capture rows llama_set_capture_layers() wrote one capture_layer_idx[] slot per accepted layer id but only range-checked the id value, so a caller repeating layer ids could advance the write index past the fixed LLAMA_MAX_LAYERS array and corrupt adjacent cparams. Stop once the array is full. On any architecture whose graph does not build a capture tensor (currently every arch except qwen35), the capture copy was skipped while output_reserve() had already allocated embd_capture, so llama_get_embeddings_capture*() returned uninitialized memory. Zero the destination rows and warn once instead. * speculative: harden dspark drafter/target contract and recovery paths Validate the drafter against the target at construction: the drafter consumes the target's hidden states (each capture row is target_hidden wide, copied verbatim) and resamples over the target's vocabulary, so a drafter trained against a different target would over-read capture rows or index the wrong vocab. Fail loudly here rather than corrupt every round. When the per-round drafter cache tail cannot be cropped, do not advance n_cache past a tail that is still physically present; reset the drafter sequence and rebuild context next round. The markov-resample mask_token_id checks aborted via GGML_ASSERT, but a sampled token can legitimately equal the mask sentinel (a real vocab id) -- that only makes a poor draft the target rejects. Warn once instead of aborting a valid run; the sequential chaining is guaranteed structurally by construction. * convert: map dspark per-layer tensors under drafter./dspark. wrappers The per-layer fallthrough passed the original tensor name to map_tensor_name, so a decoder tensor nested under a drafter. or dspark. wrapper kept that unsupported prefix and failed to map (only the model. prefix map_tensor_name understands worked). Strip the dspark-specific wrapper for the fallthrough while preserving any standard model. prefix. * common: build the dspark Markov CUDA TU for ggml's architectures, not SM80 only The resample TU pinned CUDA_ARCHITECTURES to 80 whenever CMAKE_CUDA_ARCHITECTURES was not defined in this scope -- which is the common case, since ggml resolves it inside its own subdirectory and it does not propagate up. That forces a PTX JIT on Hopper/H100 and fails to build for the pre-Ampere GPUs the rest of ggml supports. Inherit the ggml-cuda target's resolved architecture list instead, falling back to native detection. * tests: gate dspark tier-2 on agreement and actually run the rs-ring test Tier 2 only failed on non-finite logits, so it passed on any finite output even with zero argmax matches. Gate on argmax-match-rate and top-5 overlap (both scale-invariant, defaulted high, CLI-overridable), and scan logits from token 0 so a bad first logit is no longer excluded from the diff/non-finite metrics. test-rs-ring-rotation was registered against the non-recurrent stories model, so it always took the self-skip path and never exercised the ring. Generate a tiny recurrent qwen35 fixture with the pure-Python generator when numpy/gguf are importable and require it; otherwise fall back to the stories model unchanged. * docs: correct dspark Markov CUDA default and drop stale scaffolding claims The Markov CUDA path is the default at runtime when the drafter has a Markov head (opt out with LLAMA_DSPARK_MARKOV_CUDA=0), not opt-in via =1 -- fix the root CMake option comment and dspark-markov.h to match the runtime behavior. Note that the device warp tree-reduction changes floating-point accumulation order versus the host scalar/BLAS paths, so the resample is functionally equivalent rather than bit-identical: a near-tie argmax can differ, which only changes the speculative proposal the target verify still arbitrates. Drop the 'no forward graph / scaffolding only' claims in the converter docstring and logging, the GGUF tensor-registry comment, and docs/dspark-scope.md -- the forward graph and block-diffusion draft loop are implemented in this branch. * tests: run rs-ring rotation on CPU (-ngl 0) The ring bit-identity check compares logits from a ring and a no-ring context and requires both to run the identical GDN compute path. On GPU backends without a fused GDN op the op falls back per context and the two diverge (observed on Metal: 'fused Gated Delta Net not supported, set to disabled'). The invariant and this test are defined on CPU, per the test header, so pin the run to CPU. * dspark: load and run drafters with GIDD log-SNR conditioning Some drafters ship a LogSnrEmbed module -- a sinusoidal featurization of a per-position log-SNR value run through a 2-layer SiLU MLP, added to the draft noise embedding before the backbone. Without loader support these drafters fail to load: their four log_snr_fc tensors are unmapped (wrong number of tensors). Add the optional GGUF metadata (dspark.log_snr_conditioning, min/max_log_snr) and the dspark.log_snr_fc1/fc2 tensors, gated on log_snr_conditioning so drafters without it load and run exactly as before. The per-position log-SNR pattern (anchor of each block at max_log_snr, mask positions at min_log_snr) and its featurization are a pure function of n_draft/block_size/min/max_log_snr, all known at graph-build time, so the feature matrix is precomputed host-side and staged via a new llm_graph_input_dspark_logsnr input; only the learned fc1/fc2 weights go through ggml. The featurization divides by (max_log_snr - min_log_snr), so when conditioning is enabled the bounds are required and validated finite and strictly ordered at load time rather than silently producing NaN embeddings. * dspark: drop unused pos from llama_set_dspark_ctx; document draft-dspark contract The pos argument to llama_set_dspark_ctx (and the v_ctx_pos it filled) was never consumed: the drafter graph only uploads the tap features, and each context row's decode position comes from the batch. Shipping a public-API parameter that has no effect is misleading, so drop it (this API is new in this branch, no external callers) along with the dead v_ctx_pos storage. Also document, at the speculative-type registry, that draft-dspark requires the driver to engage multi-layer capture (llama_set_capture_layers plus per-row logits) before drafting: the reference driver is tests/test-dspark-real-eval.cpp, and the generic CLI/server paths do not yet engage capture, so selecting it there fails at the first draft with a clear error. * speculative: validate the target's configured capture-layer count in dspark process() The dspark row copy reads n_embd_cap = n_capture * n_embd floats from each target capture row, but capture layers are engaged by the driver after the impl is constructed, so the ctor's n_embd/n_vocab checks cannot see the configured layer count. A driver that engaged fewer layers than the drafter was trained on would over-read past the end of the capture row; more layers would feed misaligned features. Check llama_get_n_capture(ctx_tgt) against the drafter's n_capture at process() time and fail loudly on mismatch. Also correct a stale ctor comment that still described the CUDA Markov resample as token-identical to the host path; it is functionally equivalent but not bit-identical (see common/dspark-markov.h).
The C++ loader requires dspark.log_snr_fc1/fc2 tensors and the log_snr_conditioning/min_log_snr/max_log_snr KV when a drafter uses log-SNR conditioning, but the Python side never mapped them, so converting any log-SNR-conditioned drafter failed with: Can not map tensor 'log_snr_embed.fc1.bias'.
…rismML-Eng#59) * metal: fix M5 device creation + add Q2_0 multi-column mul_mv kernels Two independent Metal changes: - The AGX_RELAX_CDM_CTXSTORE_TIMEOUT override (added for the long-context command-buffer timeout on M1/M2, ggml-org#20141) prevents MTLCreateSystemDefaultDevice() from returning a device on M5 / current macOS. Keep it on by default and disable it only when sysctl reports an M5 chip; GGML_METAL_RELAX_CDM_CTXSTORE_TIMEOUT=0/1 forces either way. - Add Q1_0-style multi-column mul_mv variants for Q2_0 (nr1 2/3/4): read the streamed weights once for nr1 output columns via a 2-bit weight expansion and an FMA inner loop, instead of re-reading them per column on the mul_mv_ext path. Opt-in via GGML_METAL_Q2_0_NR1 (default routing unchanged). Measured [4096,14336] on M5 Pro: nr1_2 93.2 us at ne11=2 vs 122 for the ext route. 41/41 test-backend-ops MUL_MAT q2_0 on both routings. * speculative: Metal DSpark Markov resample + quantized markov heads Adds a Metal device path for the block Markov resample, alongside the existing CUDA path. It builds one dependency-chain graph for the whole draft block (each step's GPU argmax feeds the next step's get_rows) and submits it once, reading the drafter's still-device-resident logits, so the sequential Markov dependency stays exact with a single sync. Falls back to the host path when the head type or backend is unsupported, or when DSPARK_MARKOV_CPU=1. Also teaches llama_model_dspark_get_markov to dequantize quantized markov head tensors (Q4_0/Q5_0/Q8_0) for the host/CUDA path, instead of rejecting everything but f32/f16/bf16 -- drafters that ship quantized heads previously had the correction silently disabled (has_markov=0). Validated on Metal: accept counts byte-identical to the CPU-forced path; the Metal resample runs ~1.8x the single-thread CPU Markov path.
…rismML-Eng#58) Two independent Metal changes: - The AGX_RELAX_CDM_CTXSTORE_TIMEOUT override (added for the long-context command-buffer timeout on M1/M2, ggml-org#20141) prevents MTLCreateSystemDefaultDevice() from returning a device on M5 / current macOS. Keep it on by default and disable it only when sysctl reports an M5 chip; GGML_METAL_RELAX_CDM_CTXSTORE_TIMEOUT=0/1 forces either way. - Add Q1_0-style multi-column mul_mv variants for Q2_0 (nr1 2/3/4): read the streamed weights once for nr1 output columns via a 2-bit weight expansion and an FMA inner loop, instead of re-reading them per column on the mul_mv_ext path. Opt-in via GGML_METAL_Q2_0_NR1 (default routing unchanged). Measured [4096,14336] on M5 Pro: nr1_2 93.2 us at ne11=2 vs 122 for the ext route. 41/41 test-backend-ops MUL_MAT q2_0 on both routings.
… path) (PrismML-Eng#61) * ggml: rows-indexed state read for the fused GDN op (ring decode path) On the ring-enabled decode path every GDN layer paid two extra dispatches per token just to feed the fused op its input state: a get_rows gather of the per-seq live states into a contiguous scratch, then a cpy of that gather into slot 0 of the (D, K, n_seqs) state input. Both are pure reads of the recurrent cache -- ~786k floats each way per layer on the 27B target -- serialized into a launch-bound decode graph, 96 dispatches and ~300 MB of scratch traffic per token across 48 layers. Add ggml_gated_delta_net_rows: the op takes the 2D cache view plus the per-seq row indices (inp->s_copy_main) as src[6] and reads each sequence's live state directly at cache row rows[seq]. K moves to op_params so both variants share one backend code path. The graph side gains build_rs_cache_view (rs_zero clear + extra-states relocation, no main gather) and qwen35 wires it on the ring path, with GGML_GDN_STATE_GATHER=1 restoring the legacy gathered path for A/B. Implemented on CPU and Metal (function-constant-gated read base, no kargs change). All other backends that support GATED_DELTA_NET reject src[6] in supports_op so rows-mode ops fall back instead of silently reading src[5] as a scratch. test-backend-ops gains rows-mode cases (single/multi-token, multi-seq, snapshot overflow, KDA): 38/38 OK on MTL0, CPU leg green. Real-eval gate: accept counts bit-identical to the gathered path at n_max 1..4 (alpaca x24). Measured on M5 Pro (cont6k Q1_0 x bin6l1 q4_0, ring 4): harness AR 32.0 -> 35.3 tok/s (+10.5%), spec@n3 33.5 -> 35.3 (+5.5%); ring-free llama-bench unchanged (~42), as expected. * ggml: fold recurrent GDN snapshot writes on Metal * metal gdn: always populate snapshot tail on write-fold, handle K==1 rows Two correctness fixes to the folded rows-mode GDN epilogue: - The write-fold followed the SET_ROWS view chain to prove the scatter consumes the GDN result, but not that it is the snapshot tail's sole consumer. The kernel now always writes the op's own documented output tail AND additionally scatters into the cache row, so a second consumer or an output/eval callback never observes an uninitialized region. - WRITE_ROWS scatter existed only in the K>1 branch; a rows-mode graph with K==1 suppressed the SET_ROWS but wrote only the output tail, losing the cache update. The K==1 final-state branch now scatters to the cache row as well. Gate: test-backend-ops GATED_DELTA_NET 39/39 on MTL0; e2e accept invariant 76/116 tau 2.3103 unchanged (default / fold-disabled / gathered). * qwen35: gate GDN rows mode to Metal-only GPU device sets rows mode uses the src[6] GDN variant, implemented on CPU and Metal only; other GPU backends reject it in supports_op, which would move the recurrent op (and its state traffic) to CPU. Select rows mode only when every GPU device in the model is Metal (ACCEL/BLAS devices are skipped). * ggml: disable OpenMP for Emscripten/WASM builds The WASM CI build enables OpenMP (-DGGML_USE_OPENMP -fopenmp=libomp), but Emscripten cannot emit the common symbols libomp's reduction helpers need (.gomp_critical_user_.reduction.var), so ggml-quants.c fails to compile. WASM has no host threads to benefit from OpenMP -- force it off for the Emscripten target instead of failing the build.
…e overflow, write-fold guards, K==1 tests) (PrismML-Eng#62) * address review: CPU workspace sizing, write-fold guards, K==1 rows tests - ggml-cpu: size the GDN scratch from the op-param K (snapshot slots), not src[5]->ne[1] -- in rows mode that dim is the cache row count, so a 1-row cache with K>1 (batch-1 block decode) undersized the scratch and overflowed the work buffer. - metal write-fold: honor ctx->use_fusion (GGML_METAL_FUSION_DISABLE), and verify the SET_ROWS target is exactly the snapshot tail (per-row state width, index count, dest row width) before suppressing it -- descent from the GDN output alone let a mis-sized view be fused, reading row indices out of bounds. - rows-mode state view: document + assert the main/extra row-range disjointness invariant that makes the deferred (read-after-relocate) main read safe. - tests: add rows-mode K==1 cases to exercise the K==1 final-state branch. * review round 2: byte-offset write-fold check, honest rows-mode ordering note, 1-row-cache K>1 test - write-fold: also verify the SET_ROWS view begins at the snapshot-tail byte offset (attn_size + (K-min(T,K))*state_size_per_snap), not just matching size/counts -- a same-sized view at another offset no longer folds. - rows-mode state view: drop the incorrect disjointness assert (s_copy returns idx*size+src0, an arbitrary slot, so it did not establish disjointness). Document the real read-before-relocation hazard (multi-seq; not reachable on the single-seq decode path) as tracked follow-up. - tests: add a rows-mode 1-row-cache K>1 case that reproduces the CPU workspace under-size the planner fix prevents. * write-fold: require compact snapshot-row stride before folding ggml_set_rows only requires contiguous rows (nb[0]); it permits an arbitrary row stride nb[1] that its kernel honors, but the fused GDN epilogue scatters the contiguous snapshot tail. Require the compact [D, n_write] layout (ne[0]==D, nb[0]==type_size, nb[1]==D*type_size) so a strided view falls through to the real SET_ROWS instead of being mis-scattered.
…n every prompt row Gives DSpark's multi-layer hidden-state tap capture its own masked flag, separate from embeddings_nextn_masked which it was previously reusing. Opting into masked=false keeps capture dense (every prompt position) regardless of batch.logits, so callers can request logits=false on context rows (as the plain AR path already does) while still getting a full per-position capture buffer for the drafter. Mirrors the existing embeddings_nextn unmasked path (llama_context.cpp) at every layer: cparams flag, output_reserve sizing, per-decode readback offset/size, and get_embeddings_capture_ith row resolution. test-dspark-real-eval.cpp now engages capture with masked=false and drops the speculative-path prefill's logits back to false on context rows, matching the AR baseline. Fixes the harness-side third of the PP slowdown reported in PrismML-Eng#33: capture previously needed logits=true on every row just to populate a capture row for it, which forced the full-vocab lm_head projection to run on every prompt position instead of one.
- restore the default masked=true on llama_set_capture_layers's public declaration -- it was mandatory there, breaking source compat for any existing 3-arg caller. - set_capture_layers() no longer stomps embeddings_nextn_masked=true as a side effect; that assignment predated the independent capture flag and made the assert below unreachable in the exact case it exists to catch (dense capture silently overriding a caller's masked=false nextn config instead of tripping the guard). - narrow_before_last_layer's capture-side deferral now only applies when the LAST layer is actually one of the requested capture layers; taps at any earlier layer already branched off cur before this point in the loop, so deferring the last layer's own narrowing for them was an unnecessary regression (recovers a little more speed on today's real checkpoints, whose taps never include the last layer). - guard dense (unmasked) capture to single-sequence ubatches: its rows are indexed/reordered assuming raw-token order, which split_equal()'s per-sequence interleaving on a multi-sequence ubatch would violate. No current consumer is multi-sequence; fail loudly instead of silently returning another sequence's capture if that changes.
…d-capture dspark: give tap capture its own unmasked path, avoid full-vocab lm_head on every prompt row
…PrismML-Eng#64) Mirror the Q1_0 default to Q2_0: for ne11>=2 use the nr1=2 multi-column variant that reads each streamed q2_0 weight group once per 2 src1 columns, instead of the mul_mv_ext route. Measured (M5 Pro, in-code microbench): nr1_2 = 93.2us vs 122us ext at ne11=2 (+31%); ne11=4 via 2 passes 171 vs 183. ne11==3 is carved out -- that is the nr1_3 occupancy cliff (195 vs 152 ext, tpb=16) that kept this path opt-in; three columns stay on ext. Output is identical (pure matvec routing); base decode (ne11=1) is unaffected. GGML_METAL_Q2_0_NR1=1 restores the old routing.
* server, speculative-simple: wire dspark tap capture (draft-dspark support) * dspark server/cli: review fixes -- batch headroom, n-max validation, ctx-shift and child-slot gating
Two cpp/wrong-type-format-argument defects (high severity): - the vision_feature_layer/proj_spatial_offsets size-mismatch throw passed size_t values to %d and had no argument for its leading %s; use %zu and pass __func__. - the qwen-flamingo projector-block loop used size_t bid, passed to the %d in the TN_QF_* tensor-name formats; make bid int so the format type matches.
…rismML-Eng#65) The download tests make live HTTP requests to http://ggml.ai/ and assert on the response. Network-restricted CI runners (self-hosted, windows-vulkan) can't reach it, so the good-URL GET fails and takes the whole arg-parser suite down on an unrelated connectivity issue. Probe the endpoint once and assert the download semantics only when it is actually reachable; otherwise print a notice and skip. No behavior change when network is present.
…g#80) common_remote_get_content() returns the HTTP status as long, but the skip-message printf used %d. Format it with %ld so the build does not fail under -Werror=format on the CUDA CI toolchain.
* cuda: speed up Q1_0 extraction with byte permutes * cuda: use unsigned halfword for Q1_0 byte-perm selectors unpack_q1_0_bytes() built the __byte_perm selectors by right-shifting a signed int16_t, so a packed halfword with the top bit set sign-extends through the shift and corrupts the selector nibbles. Take the packed word as uint16_t (widened through uint32_t before the shift) at the helper and at both the MMVQ and MMQ call sites.
…fork note (PrismML-Eng#78) * ci(release): ship full self-contained Windows Vulkan/HIP bundles readme: add Prism fork note (start with Bonsai-demo, main caveats) * readme: fix Q2_0 model-file guidance (fork=Q2_0, mainline=Q2_0_g64, PQ2_0 future) + link demo status; ASCII punctuation
ggml_vec_dot_q2_0_q8_0 gated its VNNI fast path on __AVX512VNNI__ && __AVX512VL__ with no AVX2/AVX-VNNI fallback, so x86 CPUs without AVX-512 silently took the scalar loop. This excludes all Intel 12th-14th gen consumer CPUs (Alder/Raptor Lake), where AVX-512 is fused off for the P/E hybrid design but AVX-VNNI is present. The fast-path body is already entirely 256-bit AVX2; the only AVX-512 dependency is the _mm256_dpbusd_epi32 intrinsic. AVX-VNNI exposes the identical operation as _mm256_dpbusd_avx_epi32, so alias the intrinsic and widen the guard. No algorithmic change, and no behavior change on AVX-512-VNNI hosts. Measured on Intel i5-13400 (Raptor Lake, AVX-VNNI, no AVX-512), 12 threads, CPU-only, same model and prompt (temperature=0): Ternary-Bonsai-8B Q2_0 decode: 2.17 -> 6.92 tok/s (3.2x) Ternary-Bonsai-8B Q2_0 prompt eval: 2.7 -> 8.6 tok/s (3.2x) Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
…-Eng#86) * ggml-cpu: x86 AVX512-VNNI repack GEMV/GEMM for Q1_0 and Q2_0 Q2_0 previously had no repack path at all and Q1_0 only a NEON one, so batched CPU mul_mat for both formats fell back to per-row vec_dot on x86. Add: - block_q2_0x4 (4-row, 8-byte-chunk interleave) with generic repack, gemv and gemm implementations - AVX512-VNNI gemv/gemm kernels for the existing q1_0 4x8 layout and the new q2_0 4x8 layout; sum(qy) is computed once per activation sub-block and the horizontal reduction happens once per output tile - repack type selection on AVX512-VNNI CPUs for both formats * ggml-cpu: fix int-overflow-before-widening in Q2_0 repack (CodeQL) Cast to size_t/int64_t before the nrow*nblocks and i*nblocks multiplications so they cannot overflow int before widening. Resolves CodeQL alerts 630/631.
Collaborator
|
Thanks for adding this, sycl is for intel Gpus? I don't have any to test it ourselves note: main different with llama.cpp's Q2_0 is that its group size is 64 (ours in this fork is 128), jsut recently we merged our q2_0 backens into llama.cpp so doing a big migration soon, and picking up recent chagnes from llama.cpp to our fork. |
Author
|
@khosravipasha Yes, SYCL is for Intel GPUs. I tested this on an Intel Arc B580 (12 GB), and both Prism Q1_0 and Q2_0 g128 run correctly with the SYCL backend. |
|
I can try to test this on a Linux Intel Arch Laptop. |
Author
|
Thanks! That would be great. I only tested it on Windows with an Intel Arc B580 12 GB, so Linux testing would be really valuable. Let us know how it goes! |
Independent test results: Linux / Arc 140V (Lunar Lake iGPU) / oneAPI 2026.0Tested this branch ( Hardware: Intel Core Ultra 7 258V / Arc 140V, 32 GB shared, Arch Linux (kernel 7.1.5), intel-compute-runtime 26.27, level-zero-loader 1.32.0. Built FP32 per Correctness — reproduces your numbers
Two things I checked specifically:
A 4B generation check produced a correct, coherent answer, so the kernels aren't just passing numeric tolerance. Throughput,
|
| Model | PP512 (t/s) | TG128 (t/s) |
|---|---|---|
| Ternary-Bonsai-27B Q2_0 (6.66 GiB) | 45.47 ± 1.03 | 3.28 ± 0.05 |
| Ternary-Bonsai-8B Q2_0 | 134.75 ± 1.43 | 11.50 ± 0.35 |
| Ternary-Bonsai-4B Q2_0 | 327.68 ± 0.41 | 20.79 ± 0.08 |
| Ternary-Bonsai-1.7B Q2_0 | 725.68 ± 22.47 | 38.10 ± 0.44 |
Far below the B580 figures, as expected for an iGPU on shared memory — not a concern, just a different point on the curve. The 27B fits at -ngl 99 without dropping layers.
One observation that may be actionable: prefill is slower than CPU
Same 8B model, same build, -ngl 0 -fa 1 -t 8:
| 8B Q2_0 | PP512 | TG128 |
|---|---|---|
GPU (-ngl 99) |
134.75 | 11.50 |
CPU (-ngl 0) |
169.89 | 7.62 |
Decode is 1.51× the CPU, so the new MMVQ kernel is clearly doing real work. But prefill is ~20% slower than just running on the CPU.
That seems consistent with what the PR covers: mul_mat_vec_q2_0_q8_1_sycl only serves batch-size-1 decode, so prompt processing falls to dequantize_row_q2_0_sycl, which is one work-item per element doing a scalar shift-and-mask:
const int code = (x[ib].qs[local / 4] >> ((local % 4) * 2)) & 0x3;There's no packed/vectorized extract here comparable to what CUDA got in 9ca265a (cuda: extract Q2_0 elements via __byte_perm). If that's the prefill bottleneck, a packed extract in the dequant kernel looks like the natural follow-up for this backend. I haven't profiled to confirm the attribution — offering it as a lead, not a conclusion.
Happy to re-run anything on this hardware if it would help.
Disclosure: test execution and drafting of this report were AI-assisted; all numbers are measured on the hardware described above and are reproducible with the commands listed.
Follow-up: Q1_0 family, KV4, mean-centering, and speculative decoding on SYCLExtending my earlier report on this branch ( My first comment only covered the Ternary (Q2_0) family. Since this PR adds both Q1_0 and Q2_0, here is the 1-bit family end-to-end, plus the KV4 / mean-centering / speculative paths. Correction to my earlier numbersThe 8B figures in my first comment were measured while the machine was under memory pressure and are too low. Re-measured with ~14 GB free and
The CPU baseline was measured in that same degraded batch, so it was re-run too (182.67 vs the 169.89 I quoted). The prefill conclusion is unchanged but the magnitude was overstated — the real gap is 14%, not ~26%. Corrected comparison below. 1-bit Bonsai (Q1_0),
|
| Model | PP512 (t/s) | TG128 (t/s) | TG vs Q2_0 |
|---|---|---|---|
| Bonsai-27B Q1_0 (3.53 GiB) | 53.41 ± 0.10 | 7.92 ± 0.01 | 2.41× |
| Bonsai-8B Q1_0 | 159.29 ± 0.30 | 27.88 ± 0.14 | 2.36× |
| Bonsai-4B Q1_0 | 280.62 ± 12.01 | 39.72 ± 3.60 | 1.91× |
| Bonsai-1.7B Q1_0 | 725.39 ± 49.16 | 79.81 ± 1.54 | 2.09× |
Q1_0 decode is consistently ~2× Q2_0, which is what you would expect from halving the weight bits on a bandwidth-bound path. Practically, the 27B Q1_0 at 7.92 t/s is usable on this iGPU where the Q2_0 27B at 3.28 t/s really is not.
Prefill vs CPU — reproduces on both formats
8B, same build, 5 repetitions, matched memory conditions:
| 8B | GPU PP512 | CPU PP512 | GPU TG128 | CPU TG128 |
|---|---|---|---|---|
| Ternary Q2_0 | 160.51 | 182.67 | 11.83 | 8.46 |
| Bonsai Q1_0 | 159.29 | 186.61 | 27.88 | 19.96 |
Decode is 1.40× the CPU on both. Prefill loses to the CPU by 14% (Q2_0) and 17% (Q1_0). That it reproduces on both formats is the useful part — it points at the shared dequant path rather than anything Q2_0-specific, consistent with the naive one-work-item-per-element dequantize_row_*_sycl kernels this PR adds.
KV4 (--cache-type-k/v q4_0)
Works on SYCL on every size, no errors. Quantized-KV flash attention is well supported here — test-backend-ops shows 336/336 OK for type_K=q4_0,type_V=q4_0.
| Model | FP16 KV (pp/tg) | KV4 (pp/tg) |
|---|---|---|
| Ternary 8B Q2_0 | 160.51 / 11.83 | 174.32 / 11.22 |
| Bonsai 8B Q1_0 | 159.29 / 27.88 | 170.96 / 25.98 |
| Ternary 27B Q2_0 | 45.47 / 3.28 | 47.56 / 3.29 |
| Bonsai 27B Q1_0 | 53.41 / 7.92 | 52.16 / 7.88 |
Decode is flat within noise, matching KV-CACHE.md's framing of KV4 as a memory tool rather than a speed tool.
K-cache mean-centering
test-kv-mean-center passes on SYCL (nmse(baseline, centered) = 4.23e-16).
One methodological note that may be worth documenting: attn_rot_k is enabled only when the K cache is quantized (llama-kv-cache.cpp:337), so llama-kv-mean-center running with the default FP16 K cache always calibrates with rotation inactive, while KV4 inference turns it active. The loader correctly refuses the mismatch. This means rotation and mean-centering are alternative mitigations for the same q4_0 K-cache error, not additive ones — so comparing "KV4+bias" against default "KV4" would conflate two changes. I measured both rotation states separately.
wikitext-2 test, 60 chunks, -fa 1:
| Config | Ternary-8B Q2_0 | Bonsai-8B Q1_0 |
|---|---|---|
| FP16 KV | 10.9908 ± 0.265 | 13.7121 ± 0.347 |
| KV4, rotation off, no bias | 11.1673 ± 0.270 | 14.0034 ± 0.356 |
| KV4 + rotation (default) | 11.0634 ± 0.268 | 13.7756 ± 0.349 |
| KV4 + mean-centering bias | 11.0509 ± 0.266 | 13.7463 ± 0.347 |
The bias recovers 66% (Q2_0) and 88% (Q1_0) of the raw KV4 quality loss, and edges out the rotation path in both cases — though that last margin is well inside the error bars, so treat it as a tie rather than a win. Calibration used wikitext-2 train, evaluation used test.
Speculative decoding (DSpark) — net loss on this backend
SPECULATIVE.md notes the path is "stable and fast on CUDA" and warns about Metal, but says nothing about SYCL. It runs correctly here, with healthy acceptance — but it is consistently slower than plain decoding. Matched prompt, -n 400 --temp 0, llama-speculative-simple vs llama-completion:
| 27B target | Speculative | Baseline | Accept | Net |
|---|---|---|---|---|
| Bonsai Q1_0 + dspark-Q4_1 | 4.99 t/s | 8.04 t/s | 68.7% | 0.62× |
| Ternary Q2_0 + dspark-Q4_1 | 2.58 t/s | 3.37 t/s | 77.6% | 0.77× |
Acceptance is fine (69-78%), so the drafter is doing its job — the cost is in verification. That plausibly ties back to the prefill result above: verifying a drafted block is a batched forward pass, which goes through the same dequant-bound path that already loses to the CPU here. If that reading is right, improving the dequant kernel would help prefill and speculative decoding together. I have not profiled to confirm the attribution, so treat it as a hypothesis.
Might be worth a note in SPECULATIVE.md that SYCL/iGPU currently regresses, alongside the existing Metal caveat.
Summary
Everything in this PR is functionally correct on Linux/SYCL/Arc — 11735/11735 on the full test-backend-ops, both formats, all four sizes, KV4 and mean-centering included. The performance gaps are all downstream of one thing: the scalar dequant kernel.
Happy to re-run anything or test a patch on this hardware.
Disclosure: test execution and drafting of this report were AI-assisted; all numbers are measured on the hardware described and reproducible with the commands given.
There was a problem hiding this comment.
Pull request overview
Adds SYCL support for Prism Q1_0 and Q2_0 group-128 formats.
Changes:
- Adds Q1_0/Q2_0 MMVQ kernels.
- Adds FP16/FP32 and GET_ROWS dequantization.
- Registers backend operation support.
Reviewed changes
Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
vecdotq.hpp |
Implements Q1_0/Q2_0 dot products. |
mmvq.cpp |
Adds and dispatches MMVQ kernels. |
ggml-sycl.cpp |
Enables backend support. |
getrows.cpp |
Dispatches Q2_0 GET_ROWS. |
dequantize.hpp |
Implements Q2_0 dequantization. |
convert.cpp |
Adds contiguous FP16/FP32 conversion. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+691
to
+694
| case GGML_TYPE_Q1_0: | ||
| return dequantize_row_q1_0_sycl; | ||
| case GGML_TYPE_Q2_0: | ||
| return dequantize_row_q2_0_sycl; |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds Intel SYCL support for the Prism Q1_0 and Q2_0 group-128 formats.
Validation
Ternary-Bonsai 1.7B Q2_0:
Ternary-Bonsai 27B Q2_0:
Environment