llama.cpp SYCL 反量化改用向量儲存，MXFP4 MoE 預填充最高提速 61%

sycl : vectorize reordered-layout dequantization stores (Q4_0, K-quants, IQ3, MXFP4) ( #30976 ) sycl : widen reordered q4_0 dequantization stores The reordered-layout Q4_0 conversion kernel processed one qs byte per work-item with two scalar stores, strided 16 elements apart. Give it the same treatment the Q8_0 kernel got in #30775 : a flat sycl::range<1> launch with four qs bytes (8 outputs) per work-item inside one 32-element block and two vec<dst_t,4> stores. Same float math and rounding as the scalar kernel; output is bit-identical. The existing k % QK4_0 == 0 assert guarantees whole 32-element blocks, so the k/8 launch covers every row exactly. The now-unreferenced scalar reorder kernel is removed. Arc Pro B70, primed prefill (prefill after decode on a reorder-installed tensor): +6.4% to +7.3% pp512/pp2048 on Q4_0 models (tested up to 27B dense); fresh-layout prefill and decode unchanged; PPL within run noise. sycl : vectorize reordered q4_K dequantization stores The reordered-layout Q4_K conversion kernel already loads its quantized inputs through sycl::vec but writes the dequantized results with scalar stores. Build the results in a vec<dst_t,n> register and emit two vector stores instead. Same per-element math; output is bit-identical. Arc Pro B70, primed prefill (prefill after decode on a reorder-installed tensor): +14.7% to +15.3% pp512/pp2048 on Q4_K-dominant models (tested up to 27B dense); fresh-layout prefill, decode and PPL unchanged. sycl : remap reordered q5_K dequantization to vec4 loads/stores The reordered-layout Q5_K conversion kernel gave each thread only four outputs fed by scalar ql/qh byte loads plus a local-memory scale broadcast, so the stores were never its bottleneck: widening the stores alone measured within noise on mixed-quant models and +0.5% on a q5_K-dominant model. Re-decompose to 32 threads per block, each consuming four consecutive ql and qh bytes and emitting two aligned vec<dst_t,4> stores; a 4-element run shares one scale pair, so the per-element math and the output are identical. Arc Pro B70, primed prefill on a q5_K-dominant 27B model: +3.6% pp512/pp2048 over the scalar kernel; fresh-layout prefill, decode and PPL unchanged. The GET_ROWS reorder path moves to the same 32-thread decomposition. sycl : remap reordered q6_K dequantization to vec4 stores The reordered-layout Q6_K conversion kernel wrote four scalars per thread, strided 32 elements apart, each pair of neighbors sharing neither instruction stream nor store transaction. Re-decompose the same 64-thread work-group so each thread emits one aligned vec<dst_t,4>: a 4-element run shares one 16-element scale group, one ql nibble source and one qh bit pair, so the per-element math and the output are identical to the scalar kernel. Arc Pro B70, primed prefill: +28.7% pp512/pp2048 on a Q6_K 27B model; fresh-layout prefill, decode and PPL unchanged. sycl : vectorize reordered q3_K dequantization stores Each thread of the reordered-layout Q3_K conversion kernel already computes four consecutive outputs; build them in a vec<dst_t,4> register and emit one vector store instead of four scalars. Same per-element math; output is bit-identical. Arc Pro B70, primed prefill: +16.8% pp512/pp2048 on a Q3_K-dominant 27B model; fresh-layout prefill, decode and PPL unchanged. sycl : vectorize reordered iq3_xxs/iq3_s dequantization stores Both reordered-layout IQ3 conversion kernels write eight contiguous outputs per thread as scalar stores (four from each grid lookup). Build two vec<dst_t,4> registers and emit two vector stores. Same per-element math; output is bit-identical. Arc Pro B70, primed prefill: +20.3% to +21.1% pp512/pp2048 on IQ3_XXS- and IQ3_S-dominant 27B models; fresh-layout prefill, decode and PPL unchanged. sycl : remap reordered q2_K dequantization to vec4 stores Re-decompose the reordered-layout Q2_K conversion kernel the same way as Q6_K: one aligned vec<dst_t,4> per thread instead of four scalars strided 32 apart. A 4-element run shares one scale byte and one 2-bit lane, so the per-element math and the output are identical to the scalar kernel. Arc Pro B70, primed prefill: +13.1% pp512/pp2048 on a Q2_K-dominant 27B model; fresh-layout prefill, decode and PPL unchanged. sycl : widen reordered mxfp4 dequantization stores Give the reordered-layout MXFP4 conversion kernel the same flat sycl::range<1> launch as Q4_0: four qs bytes (8 outputs) per work-item inside one 32-element block and two vec<dst_t,4> stores. Same LUT lookups and per-block scale; output is bit-identical. The k % QK_MXFP4 == 0 assert guarantees whole blocks, so the k/8 launch covers every row exactly. The old scalar kernel is removed. The reordered layout for MXFP4 is installed on the MoE expert path (reorder_qw_mxfp4_moe); dense MXFP4 tensors keep the normal layout. Arc Pro B70, MXFP4 MoE 20B model: +57% pp2048 / +61% pp512 prefill (primed and fresh alike), decode unchanged, PPL within the model's run-to-run noise. Website: https://llama.app Attestations: https://github.com/ggml-org/llama.cpp/attestations/61167328 macOS/iOS: macOS Apple Silicon (arm64) macOS Apple Silicon (arm64, KleidiAI enabled) DISABLED macOS Intel (x64) iOS XCFramework Linux: Ubuntu x64 (CPU) Ubuntu arm64 (CPU) Ubuntu s390x (CPU) Ubuntu 26.04 s390x (CPU) Ubuntu 26.04 s390x (zDNN) Ubuntu x64 (Vulkan) Ubuntu arm64 (Vulkan) Ubuntu x64 (CUDA 12) - CUDA 12.8 libraries Ubuntu x64 (CUDA 13) - CUDA 13.4 libraries Ubuntu arm64 (CUDA 13) - CUDA 13.4 libraries Ubuntu x64 (ROCm 10.0) Ubuntu x64 (OpenVINO) Ubuntu x64 (SYCL FP32) Ubuntu x64 (SYCL FP16) Linux arm64 (Snapdragon: CPU, Adreno GPU, Hexagon NPU) - setup guide Android: Android arm64 (CPU) Android arm64 (Snapdragon: CPU, Adreno GPU, Hexagon NPU) - setup guide Windows: Windows x64 (CPU) Windows arm64 (CPU) Windows arm64 (OpenCL Adreno) Windows x64 (CUDA 12) - CUDA 12.4 DLLs Windows x64 (CUDA 13) - CUDA 13.4 DLLs Windows arm64 (CUDA 13) - CUDA 13.4 DLLs Windows x64 (Vulkan) Windows arm64 (Vulkan) Windows x64 (OpenVINO) Windows x64 (SYCL) Windows x64 (ROCm 10.0) openEuler: DISABLED openEuler x86 (310p) openEuler x86 (910b, ACL Graph) openEuler aarch64 (310p) openEuler aarch64 (910b, ACL Graph) UI: UI