From d81d0d8fb95e52a66992bcfc212b1891b4d2f94a Mon Sep 17 00:00:00 2001 From: budiony Date: Thu, 22 Oct 2026 16:39:01 +0300 Subject: [PATCH] Multi-column on Intel XMX matrix engines refactored (#29864) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Multi-column on Intel XMX matrix engines refactored Multi-column (2 to 80 activation columns) mul_mat_vec_q for reordered K-quant and Q8_0 weights on the Intel XMX matrix engines (int8 DPAS, ESIMD), for the SYCL backend. It targets speculative decoding verification (2 to 8 tokens per step) and short prompt batches (9 to 80 tokens). Three types also use it for a single column (Q2_K, Q3_K, Q8_0). * Fixed regression bug real bug found: the first version produced NaN whenever the row count wasn't a multiple of 4 * Removed GGML_TYPE_Q8_0 activation * rewritten xmx_mul_mat 64-byte alignment and 16-byte pitch check on the activations * Implemented fused gate/up kernel One launch computes up * act(gate) for both FFN weights, reading the activations once and applying SWIGLU or GEGLU in the epilogue. It covers up to 16 columns and works for all six XMX types. * Removed the temporary XMX_NOGLU environment switch * Q2_K min-sum optimization Q2_K weight is d·scl·q − dmin·mc, where scl (the 4-bit scale) and mc (the 4-bit min code) apply to each half of a 32-weight group * Added dispatch check and GRF sweep tg128 on Q4_K_M goes from 26.67 to 28.23 tok/s on the B70 (+5.9%). Q6_K is unchanged. MUL_MAT passes 1727/1727 and MUL_MAT_VEC_FUSION 1361/1361. Perplexity at batch 1 is 5.2064 with XMX and 5.2058 without. * Renamed variables * Added 2 environment variables GGML_SYCL_XMX_MAX_COLS (default 80, any value 0 or higher) and GGML_SYCL_XMX_GLU_MAX_COLS (default 16, range 0 to 16): the largest batch for the fused gate/up kernel. It is clamped to 16 because that is the most the fused kernel is compiled for, so the variable can only lower it. * Update esimd.hpp 1. All five quant structs were repeating the same boilerplate, so I moved it into shared pieces. 2. A new helper xmx_term does one DPAS, converts it to float and scales it per row. The four branches are now one or two lines each, and the shared Cb += z * d8r is written once. The refactor passes all three op tests. MUL_MAT_VEC_FUSION is 1369/1369, MUL_MAT is 2281/2281 (the count includes your uncommitted f32 tests) and MUL_MAT_ID is 937/937. --- docs/backend/SYCL.md | 3 + ggml/src/ggml-sycl/common.hpp | 42 ++ ggml/src/ggml-sycl/esimd.hpp | 810 +++++++++++++++++++++++++++++++ ggml/src/ggml-sycl/fusion.cpp | 12 +- ggml/src/ggml-sycl/ggml-sycl.cpp | 88 +++- ggml/src/ggml-sycl/mmvq.cpp | 163 +++++++ ggml/src/ggml-sycl/mmvq.hpp | 18 + tests/test-backend-ops.cpp | 12 + 8 files changed, 1143 insertions(+), 5 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 81d293e4f..249498450 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -819,6 +819,9 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm | GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. | | GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute. Unsupported types and layouts fall back to the standalone op kernels. See `ggml_sycl_can_fuse()`. | | GGML_SYCL_ENABLE_ESIMD | 0 or 1 (default)| Enable ESIMD kernels when available. | +| GGML_SYCL_ENABLE_XMX_ESIMD | 0 or 1 (default) | Use the XMX (DPAS) multi-column mat-vec kernels for reordered Q2_K to Q6_K and Q8_0 weights on Xe2 GPUs. Speeds up speculative decoding and short prompt batches. | +| GGML_SYCL_XMX_MAX_COLS | non-negative integer, 80 (default) | Most activation columns (batch size) the `GGML_SYCL_ENABLE_XMX_ESIMD` kernels take; wider batches use the other mat-mul kernels. A larger value keeps XMX for bigger batches (as repeated launches), `0` never uses it. | +| GGML_SYCL_XMX_GLU_MAX_COLS | 0 to 16, 16 (default) | Most columns the fused gate and up XMX kernel takes; wider batches run gate and up as separate mat-vecs. 16 is the most the kernel is built for, larger values are clamped to 16 and `0` turns the fused kernel off. | | GGML_SYCL_XMX_GATHER_TYPES | decimal bitmask, all bits set (default) | Weight formats that may use the XMX dequant-GEMM paths, which dequantize weights straight into the XMX tiles. This speeds up prompt processing of MoE models on GPUs with XMX units (Arc A- and B-series, Arc Pro, Data Center GPU Max), for example pp512 of Qwen3-30B-A3B UD-IQ3_XXS by about 50% on an Arc Pro B60. Bits:
* 1: IQ4_NL, 2: IQ3_S, 4: IQ4_XS, 8: IQ3_XXS, 16: IQ2_XXS, 32: IQ2_XS, 64: IQ2_S, 128: IQ1_S, 256: IQ1_M
* 512: Q8_0, 1024: Q4_K, 2048: Q5_K, 4096: Q6_K (MoE `MUL_MAT_ID` only)
Add values to combine them, for example `3` for IQ4_NL and IQ3_S; `0` disables the paths. A set bit does not force the path: batches of more than 64 tokens per expert or row lengths that are not a multiple of 256 (32 for IQ4_NL and Q8_0) use the library GEMM. | | GGML_SYCL_XMX_GATHER_SHAPES | decimal bitmask, 255 (default) | XMX `joint_matrix` combinations the paths of `GGML_SYCL_XMX_GATHER_TYPES` may use; the operand type comes from `GGML_SYCL_DYNAMIC_PRECISION` and the best supported combination is picked automatically (logged as `fg_pick_combo`). Bits:
* Xe2, Xe3, Xe-HPC: 1: f16 8x16x16, 2: f16 16x16x16, 4: f16 32x64x16, 8: f16 32x64x32, 32: tf32 8x16x8, 64: bf16 8x16x16
* Xe-HPG (Arc A770, ARL-H): 16: f16 8x8x16, 128: bf16 8x8x16
Clear a bit to exclude a combination, or set a single bit to force one for testing. | | GGML_SYCL_DYNAMIC_PRECISION | `F16` (default with `GGML_SYCL_F16=ON`), `BF16`, `TF32` or `F32` (default otherwise) | Operand type of the XMX dequant-GEMM paths (`GGML_SYCL_XMX_GATHER_TYPES`); accumulation is always f32. `F16` is the fastest, but activations above 65504 overflow. `BF16` keeps the f32 range at a 7-bit mantissa, `TF32` keeps the range and the f16 mantissa but is about 30% slower and needs Xe2, Xe3 or Xe-HPC, and `F32` turns the XMX paths off. Ops that request a higher src1 precision ([TAG_GGML_PREC]) get it regardless of this setting. | diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index 0bf6facec..c6f285f96 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -68,6 +68,9 @@ extern int g_ggml_sycl_enable_optimize; extern int g_ggml_sycl_enable_fusion; extern int g_ggml_sycl_enable_esimd; extern int g_ggml_sycl_mmvq_wide; +extern int g_ggml_sycl_enable_xmx_esimd; +extern int g_ggml_sycl_xmx_max_cols; +extern int g_ggml_sycl_xmx_glu_max_cols; extern int g_ggml_sycl_prioritize_dmmv; // Which quantized weight formats may take the XMX dequant-GEMM paths. A bitmask rather than one @@ -227,6 +230,41 @@ typedef sycl::float2 dfloat2; #define MMVQ_MAX_BATCH_SIZE 8 +#if defined(__INTEL_LLVM_COMPILER) && __has_include() +# define GGML_SYCL_MMVQ_HAS_XMX +#endif // __INTEL_LLVM_COMPILER + +// default of the most columns the XMX mul_mat_vec_q handles (GGML_SYCL_XMX_MAX_COLS), wider batches use other kernels +#define GGML_SYCL_XMX_MAX_COLS_DEFAULT 80 + +// most columns the fused gate and up XMX kernel is built for, GGML_SYCL_XMX_GLU_MAX_COLS can only lower this limit +#define GGML_SYCL_XMX_GLU_MAX_COLS 16 + +static inline bool ggml_sycl_xmx_supports_type(ggml_type type) { + switch (type) { + case GGML_TYPE_Q2_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: + case GGML_TYPE_Q6_K: + case GGML_TYPE_Q8_0: + return true; + default: + return false; + } +} + +// the DPAS shapes and 2D block loads used by the XMX kernel are those of Xe2 +static inline bool ggml_sycl_xmx_supports_arch(gpu_arch arch) { + return arch == gpu_arch::intel_gpu_bmg_g21 || arch == gpu_arch::intel_gpu_bmg_g31 || + arch == gpu_arch::intel_gpu_lnl_m; +} + +// fewest columns for which the XMX kernel is faster than the other kernels +static inline int ggml_sycl_xmx_min_cols(ggml_type type) { + return type == GGML_TYPE_Q2_K || type == GGML_TYPE_Q3_K || type == GGML_TYPE_Q4_K ? 1 : 2; +} + static int g_all_sycl_device_count = -1; static bool g_ggml_backend_sycl_buffer_type_initialized = false; @@ -310,6 +348,10 @@ struct ggml_sycl_device_info { const ggml_sycl_device_info & ggml_sycl_info(); +static inline bool ggml_sycl_xmx_enabled(int device) { + return g_ggml_sycl_enable_xmx_esimd && ggml_sycl_xmx_supports_arch(ggml_sycl_info().devices[device].hw_info.arch); +} + static constexpr size_t SYCL_BUFFER_ALIGNMENT = 128; struct ggml_sycl_pool { diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 0485ff0ce..85cc93e21 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -2,6 +2,8 @@ #define GGML_SYCL_ESIMD_HPP #include +#include +#include #include "common.hpp" @@ -584,6 +586,814 @@ template <> struct esimd_reorder_q_traits { } }; +// +// Multi-column mul_mat_vec for reordered K-quant weights on the XMX engines (int8 DPAS). +// A thread owns 16 weight rows (B operand, one dword is 4 weights of a row), the q8_1 activation +// columns are the A operand. A super-block is 8 groups of 32, run as 2 parts of 4 groups. A quant +// type supplies its loads and unpack as a traits struct. +// +using sycl::ext::intel::esimd::simd; +constexpr int XMX_ROWS = 16; // weight rows per thread (DPAS N) +constexpr int XMX_CHUNK = 64; // widest single launch, more columns run as repeated launches +constexpr int XMX_FUSED_CHUNK = GGML_SYCL_XMX_GLU_MAX_COLS; // widest launch of the fused gate and up kernel + +// threads that split K for one tile, sized so that the shared memory stays small +template constexpr int xmx_n_split() { + return NC <= 8 ? 32 : (NC == 16 ? 16 : 8); +} + +inline bool xmx_supported(const void * vx, const int ncols) { + // 2D block loads need a 64 byte aligned base + return ncols % QK_K == 0 && (uintptr_t) vx % 64 == 0; +} + +// transposed 2D block load of 16 rows, dword (kd, n) lands at [kd * 16 + n] +// 2D block loads need rows of at least 64 bytes, shorter rows (SR) are read row by row +template +ESIMD_INLINE simd xmx_load_t(const void * base, unsigned pitch, unsigned h, int x, int y) { + using namespace sycl::ext::intel::esimd; + if constexpr (SR) { + simd r; + for (int n = 0; n < XMX_ROWS; ++n) { + const unsigned row = y + n < (int) h ? y + n : h; + const char * p = (const char *) base + (size_t) row * (pitch + 1); + r.template select(n) = block_load((const uint32_t *) p + x); + } + return r; + } else { + return load_2d((const uint32_t *) base, pitch, h, pitch, x, y); + } +} + +// per byte (q - zero) as int8 for q in 0..127, no borrow between bytes +ESIMD_INLINE simd xmx_bytes_sub(simd x, uint32_t zero_x4) { + return ((x | 0x80808080u) - zero_x4) ^ 0x80808080u; +} + +// How the scales of a quant type enter the DPAS epilogue of a group of 32 (see xmx_group) +enum class xmx_scale_kind { + split, // two scales per group, one per 16 wide half: the weights are masked to one half for each DPAS + min, // one scale and one min per group: a second DPAS sums the activations for the min term + min16, // a scale and a min per 16 wide half: the scales are folded into the DPAS operands + plain, // one scale per group +}; + +// What every quant type states about itself. nc_wide is the tile width from which it takes the large register file. +template struct xmx_traits_base { + static constexpr bool short_rows = SR; + static constexpr int nc_wide = NC_WIDE; + static constexpr bool has_min = KIND == xmx_scale_kind::min; + static constexpr bool plain = KIND == xmx_scale_kind::plain; + static constexpr bool min16 = KIND == xmx_scale_kind::min16; +}; + +// The part of a context every quant type shares: the last row (the clamp of 2D loads), this thread's first row and +// the number of super-blocks per row. A type adds its array pointers and 2D load pitches. +struct xmx_ctx_base { + unsigned h; + int row0; + int bpr; +}; + +ESIMD_INLINE void xmx_ctx_init(xmx_ctx_base & c, int nrows, int bpr, int row0) { + c.h = nrows - 1; + c.row0 = row0; + c.bpr = bpr; +} + +// pitch of a 2D load over an array with `bytes` bytes per super-block +constexpr unsigned xmx_pitch(int bpr, int bytes) { + return bpr * bytes - 1; +} + +// per group row scales of the 16 rows, one vector per group +template struct xmx_row_scales { + simd v[N]; + + ESIMD_INLINE simd get(int i) { return v[i]; } +}; + +// the 12 packed scale bytes of a super-block (Q3_K, Q4_K, Q5_K) of the 16 rows as three dwords +ESIMD_INLINE void xmx_load_sc12(const uint8_t * sc, + const simd & bi, + simd & w0, + simd & w1, + simd & w2) { + using namespace sycl::ext::intel::esimd; + w0 = gather((const uint32_t *) sc, bi * 12); + w1 = gather((const uint32_t *) sc, bi * 12 + 4); + w2 = gather((const uint32_t *) sc, bi * 12 + 8); +} + +// d and dmin (a half2 per super-block, Q2_K, Q4_K, Q5_K) of the 16 rows +ESIMD_INLINE void xmx_load_dm(const sycl::half * dm, + const simd & bi, + simd & d, + simd & dmin) { + using namespace sycl::ext::intel::esimd; + d = convert(gather(dm, bi * (2 * (uint32_t) sizeof(sycl::half)))); + dmin = convert( + gather(dm, bi * (2 * (uint32_t) sizeof(sycl::half)) + (uint32_t) sizeof(sycl::half))); +} + +// Q6_K reorder layout: [ql: nb*128] [qh: nb*64] [scales int8: nb*16] [d: nb*2] +template struct xmx_traits_q6_k : xmx_traits_base { + struct ctx : xmx_ctx_base { + const uint8_t * ql; + const uint8_t * qh; + const int8_t * scales; + const sycl::half * dh; + unsigned pl, ph, ps; + }; + + static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) { + const size_t nb = (size_t) nrows * bpr; + ctx c; + xmx_ctx_init(c, nrows, bpr, row0); + c.ql = (const uint8_t *) vx; + c.qh = c.ql + nb * (QK_K / 2); + c.scales = (const int8_t *) (c.qh + nb * (QK_K / 4)); + c.dh = (const sycl::half *) (c.scales + nb * (QK_K / 16)); + c.pl = xmx_pitch(bpr, QK_K / 2); + c.ph = xmx_pitch(bpr, QK_K / 4); + c.ps = xmx_pitch(bpr, QK_K / 16); + return c; + } + + // row scales of the 16 scale groups of a super-block + struct scales_t { + simd sc8; + simd dv; + + ESIMD_INLINE simd get(int gr) { + simd s = sc8.template select<16, 4>(64 * (gr / 4) + (gr % 4)); + return sycl::ext::intel::esimd::convert(s) * dv; + } + }; + + // rows: clamped row index of each of the 16 lanes + static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd & rows) { + using namespace sycl::ext::intel::esimd; + scales_t s; + s.dv = convert(gather(c.dh, (rows * c.bpr + sb) * (uint32_t) sizeof(sycl::half))); + simd raw = xmx_load_t<4, SR>(c.scales, c.ps, c.h, sb * 4, c.row0); + s.sc8 = raw.template bit_cast_view(); + return s; + } + + struct part_t { + simd a, b, h; + }; + + static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) { + part_t p; + p.a = xmx_load_t<8, SR>(c.ql, c.pl, c.h, sb * 32 + part * 16, c.row0); + p.b = xmx_load_t<8, SR>(c.ql, c.pl, c.h, sb * 32 + part * 16 + 8, c.row0); + p.h = xmx_load_t<8, SR>(c.qh, c.ph, c.h, sb * 16 + part * 8, c.row0); + return p; + } + + static ESIMD_INLINE simd unpack(const part_t & p, int g, int /*part*/) { + simd qs = (g & 1) ? p.b : p.a; + simd x; + if (g == 0) { + x = (qs & 0x0F0F0F0Fu) | ((p.h << 4) & 0x30303030u); + } else if (g == 1) { + x = (qs & 0x0F0F0F0Fu) | ((p.h << 2) & 0x30303030u); + } else if (g == 2) { + x = ((qs >> 4) & 0x0F0F0F0Fu) | (p.h & 0x30303030u); + } else { + x = ((qs >> 4) & 0x0F0F0F0Fu) | ((p.h >> 2) & 0x30303030u); + } + return xmx_bytes_sub(x, 0x20202020u); + } +}; + +// Q3_K reorder layout: [qs: nb*64] [hmask: nb*32] [scales: nb*12] [d: nb*2] +// chunk s of 32: 2 bits from qs byte 32*(s/4)+l at shift 2*(s%4), high bit is hmask[l] bit s, value is q - 4 +template struct xmx_traits_q3_k : xmx_traits_base { + struct ctx : xmx_ctx_base { + const uint8_t * qs; + const uint8_t * hm; + const uint8_t * sc; + const sycl::half * dh; + unsigned pq, ph; + }; + + static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) { + const size_t nb = (size_t) nrows * bpr; + ctx c; + xmx_ctx_init(c, nrows, bpr, row0); + c.qs = (const uint8_t *) vx; + c.hm = c.qs + nb * (QK_K / 4); + c.sc = c.hm + nb * (QK_K / 8); + c.dh = (const sycl::half *) (c.sc + nb * 12); + c.pq = xmx_pitch(bpr, QK_K / 4); + c.ph = xmx_pitch(bpr, QK_K / 8); + return c; + } + + using scales_t = xmx_row_scales<16>; + + static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd & rows) { + using namespace sycl::ext::intel::esimd; + scales_t s; + simd bi = rows * c.bpr + sb; + simd w0, w1, w2; + xmx_load_sc12(c.sc, bi, w0, w1, w2); + simd dv = convert(gather(c.dh, bi * (uint32_t) sizeof(sycl::half))); +#pragma unroll + for (int i = 0; i < 16; ++i) { + const simd lo = (i % 8) < 4 ? w0 : w1; + const int bl = i % 4; + simd nib = (lo >> (8 * bl + (i < 8 ? 0 : 4))) & 0xFu; + simd hi = ((w2 >> (8 * bl + 2 * (i / 4))) & 3u) << 4; + simd code = nib | hi; + s.v[i] = convert(code - 32) * dv; + } + return s; + } + + struct part_t { + simd q, h; + }; + + static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) { + part_t p; + p.q = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 16 + part * 8, c.row0); + p.h = xmx_load_t<8, SR>(c.hm, c.ph, c.h, sb * 8, c.row0); + return p; + } + + static ESIMD_INLINE simd unpack(const part_t & p, int g, int part) { + const int s = 4 * part + g; + simd x = ((p.q >> (2 * g)) & 0x03030303u) | (((p.h >> s) & 0x01010101u) << 2); + return xmx_bytes_sub(x, 0x04040404u); + } +}; + +// Q4_K and Q5_K (QH) reorder layout: [qs: nb*128] ([qh: nb*32]) [scales: nb*12] [dm: nb*4] +// group j of 32: nibble of qs byte 32*(j/2)+l (low for even j), Q5_K adds 16 * bit j of qh[l] +// value is dall * sc_j * q - dmin * m_j, one scale and one min per group +template struct xmx_traits_q45_k : xmx_traits_base { + struct ctx : xmx_ctx_base { + const uint8_t * qs; + const uint8_t * qh; + const uint8_t * sc; + const sycl::half * dm; + unsigned pq, ph; + }; + + static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) { + const size_t nb = (size_t) nrows * bpr; + ctx c; + xmx_ctx_init(c, nrows, bpr, row0); + c.qs = (const uint8_t *) vx; + c.qh = c.qs + nb * (QK_K / 2); + c.sc = QH ? c.qh + nb * (QK_K / 8) : c.qh; + c.dm = (const sycl::half *) (c.sc + nb * 12); + c.pq = xmx_pitch(bpr, QK_K / 2); + c.ph = xmx_pitch(bpr, QK_K / 8); + return c; + } + + struct scales_t { + xmx_row_scales<8> d; // dall * scale of each group + xmx_row_scales<8> m; // dmin * min of each group + }; + + static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd & rows) { + using namespace sycl::ext::intel::esimd; + scales_t s; + simd bi = rows * c.bpr + sb; + simd w0, w1, w2; + xmx_load_sc12(c.sc, bi, w0, w1, w2); + simd d, dmin; + xmx_load_dm(c.dm, bi, d, dmin); +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int jj = j % 4; + simd sc, m; + if (j < 4) { + sc = (w0 >> (8 * jj)) & 63u; + m = (w1 >> (8 * jj)) & 63u; + } else { + sc = ((w2 >> (8 * jj)) & 0xFu) | (((w0 >> (8 * jj + 6)) & 3u) << 4); + m = ((w2 >> (8 * jj + 4)) & 0xFu) | (((w1 >> (8 * jj + 6)) & 3u) << 4); + } + s.d.v[j] = convert(simd(sc)) * d; + s.m.v[j] = convert(simd(m)) * dmin; + } + return s; + } + + struct part_t { + simd a, b, h; + }; + + static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) { + part_t p; + p.a = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 32 + part * 16, c.row0); + p.b = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 32 + part * 16 + 8, c.row0); + if constexpr (QH) { + p.h = xmx_load_t<8, SR>(c.qh, c.ph, c.h, sb * 8, c.row0); + } + return p; + } + + static ESIMD_INLINE simd unpack(const part_t & p, int g, int part) { + simd qs = (g < 2) ? p.a : p.b; + simd x = (qs >> (4 * (g & 1))) & 0x0F0F0F0Fu; + if constexpr (QH) { + x |= ((p.h >> (4 * part + g)) & 0x01010101u) << 4; + } + return x; + } +}; + +// Q2_K reorder layout: [qs: nb*64] [scales: nb*16] [dm: nb*4] +// chunk s of 32: 2 bits from qs byte 32*(s/4)+l at shift 2*(s%4), each 16 wide half has a 4 bit scale and min +// value is d * (code & 15) * q - dmin * (code >> 4) +template struct xmx_traits_q2_k : xmx_traits_base { + struct ctx : xmx_ctx_base { + const uint8_t * qs; + const uint8_t * sc; + const sycl::half * dm; + unsigned pq, ps; + }; + + static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) { + const size_t nb = (size_t) nrows * bpr; + ctx c; + xmx_ctx_init(c, nrows, bpr, row0); + c.qs = (const uint8_t *) vx; + c.sc = c.qs + nb * (QK_K / 4); + c.dm = (const sycl::half *) (c.sc + nb * (QK_K / 16)); + c.pq = xmx_pitch(bpr, QK_K / 4); + c.ps = xmx_pitch(bpr, QK_K / 16); + return c; + } + + struct scales_t { + simd code; // 4 bit scale (low) and min (high) of each 16 wide half + simd d, dmin; + + ESIMD_INLINE simd byte(int gr) { return code.template select<16, 4>(64 * (gr / 4) + (gr % 4)); } + }; + + static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd & rows) { + using namespace sycl::ext::intel::esimd; + scales_t s; + simd bi = rows * c.bpr + sb; + simd raw = xmx_load_t<4, SR>(c.sc, c.ps, c.h, sb * 4, c.row0); + s.code = raw.template bit_cast_view(); + xmx_load_dm(c.dm, bi, s.d, s.dmin); + return s; + } + + struct part_t { + simd q; + }; + + static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) { + part_t p; + p.q = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 16 + part * 8, c.row0); + return p; + } + + static ESIMD_INLINE simd unpack(const part_t & p, int g, int /*part*/) { + return (p.q >> (2 * g)) & 0x03030303u; + } +}; + +// Q8_0 reorder layout: [qs: nrows*K bytes] [d: one half per 32 weights] +// the int8 weights feed DPAS as they are, a group is one block with one scale +template struct xmx_traits_q8_0 : xmx_traits_base { + struct ctx : xmx_ctx_base { + const uint8_t * qs; + const sycl::half * dh; + unsigned pq; + }; + + static ESIMD_INLINE ctx make(const void * vx, int nrows, int bpr, int row0) { + ctx c; + xmx_ctx_init(c, nrows, bpr, row0); + c.qs = (const uint8_t *) vx; + c.dh = (const sycl::half *) (c.qs + (size_t) nrows * bpr * QK_K); + c.pq = xmx_pitch(bpr, QK_K); + return c; + } + + struct scales_t { + xmx_row_scales<8> d; // scale of each of the 8 blocks of a super-block + }; + + static ESIMD_INLINE scales_t load_scales(const ctx & c, int sb, const simd & rows) { + using namespace sycl::ext::intel::esimd; + scales_t s; + simd blk = (rows * c.bpr + sb) * (QK_K / QK8_0); +#pragma unroll + for (int j = 0; j < 8; ++j) { + s.d.v[j] = convert(gather(c.dh, (blk + j) * (uint32_t) sizeof(sycl::half))); + } + return s; + } + + struct part_t { + simd q[4]; + }; + + static ESIMD_INLINE part_t load_part(const ctx & c, int sb, int part) { + part_t p; +#pragma unroll + for (int g = 0; g < 4; ++g) { + p.q[g] = xmx_load_t<8, SR>(c.qs, c.pq, c.h, sb * 64 + 8 * (4 * part + g), c.row0); + } + return p; + } + + static ESIMD_INLINE simd unpack(const part_t & p, int g, int /*part*/) { return p.q[g]; } +}; + +template using xmx_traits_q4_k = xmx_traits_q45_k; +template using xmx_traits_q5_k = xmx_traits_q45_k; + +// registers per thread, the large file halves the resident threads but avoids spills in wide tiles, the fused +// kernel keeps two sets of weights and scales and always takes it +template constexpr int xmx_grf() { + return FUSED || NC >= T::nc_wide ? 256 : 128; +} + +// One DPAS of a block of RC columns: the int32 products of the weights B and the activations A, as float and scaled +// per weight row. Every quant type sums one or two of these per group of 32. +template +ESIMD_INLINE simd xmx_term(const simd & B, + const simd & A, + const simd & s) { + using namespace sycl::ext::intel::esimd; + namespace xmx = sycl::ext::intel::esimd::xmx; + + simd C = xmx::dpas<8, RC, int>(B, A); + return convert(C) * s.template replicate(); +} + +// one group of 32 of one weight set, accumulated into Cf for all the columns of the tile +template +ESIMD_INLINE void xmx_group(simd & Cf, + const typename T::part_t & pd, + typename T::scales_t & sc, + A2T & A2, + D8T & d8f, + simd & Bl, + simd & Bh, + const simd & Ol, + const int g, + const int part) { + using namespace sycl::ext::intel::esimd; + + constexpr int NT = XMX_ROWS; + constexpr int RC = NC < 8 ? NC : 8; + constexpr int NB = NC / RC; + constexpr bool split = !T::has_min && !T::plain && !T::min16; // two scales per group of 32 + constexpr bool pair = (RC * 32) % 64 == 0; // a block of RC rows fills whole registers + + const int gi = 4 * part + g; + + simd x = T::unpack(pd, g, part); + simd B = x.template bit_cast_view(); + if constexpr (split) { + Bl.template select<256, 1>(0) = B.template select<256, 1>(0); + Bh.template select<256, 1>(256) = B.template select<256, 1>(256); + } + + simd rd; + simd rm; + simd rs0; + simd rs1; + + // Q2_K: scale and min of each 16 wide half are 4 bit codes, the weight is d * scl * q - dmin * mc. scl * q + // is at most 45, so a dword multiply keeps the bytes apart and the scaled weights go to DPAS as int8. The + // min term sums a * mc over K, the codes mc take the place of the weights in a second DPAS. + simd Bs; + simd Bm; + if constexpr (T::min16) { + const simd c0 = sc.byte(2 * gi); + const simd c1 = sc.byte(2 * gi + 1); + simd sl0 = convert(simd(c0 & (uint8_t) 15)); + simd sl1 = convert(simd(c1 & (uint8_t) 15)); + simd mc0 = convert(simd(c0 >> (uint8_t) 4)) * 0x01010101u; + simd mc1 = convert(simd(c1 >> (uint8_t) 4)) * 0x01010101u; + simd ws; + simd wm; + ws.template select<64, 1>(0) = x.template select<64, 1>(0) * sl0.template replicate<4>(); + ws.template select<64, 1>(64) = x.template select<64, 1>(64) * sl1.template replicate<4>(); + wm.template select<64, 1>(0) = mc0.template replicate<4>(); + wm.template select<64, 1>(64) = mc1.template replicate<4>(); + Bs = ws.template bit_cast_view(); + Bm = wm.template bit_cast_view(); + } + if constexpr (T::plain) { + rd = sc.d.get(gi); + } else if constexpr (T::has_min) { + rd = sc.d.get(gi); + rm = sc.m.get(gi); + } else if constexpr (T::min16) { + // scales are applied through Bs and Bm + } else { + rs0 = sc.get(2 * gi); + rs1 = sc.get(2 * gi + 1); + } + +#pragma unroll + for (int b = 0; b < NB; ++b) { + simd A = + A2[b].template select(pair ? (g % 2) * RC * 8 : 0).template bit_cast_view(); + simd d8r = + d8f[b].template select(gi).template replicate_vs_w_hs(0); + auto Cb = Cf.template select(b * RC * NT); + + simd z; + if constexpr (T::plain) { + z = xmx_term(B, A, rd); + } else if constexpr (T::has_min) { + // the second DPAS (all ones against the activations) is the exact integer sum of the activations + z = xmx_term(B, A, rd) - xmx_term(Ol, A, rm); + } else if constexpr (T::min16) { + z = xmx_term(Bs, A, sc.d) - xmx_term(Bm, A, sc.dmin); + } else { + z = xmx_term(Bl, A, rs0) + xmx_term(Bh, A, rs1); + } + Cb += z * d8r; + } +} + +// activation of the GLU applied to the gate, uniform across the launch +template ESIMD_INLINE simd xmx_glu_act(const simd & gate, const int glu_op) { + using namespace sycl::ext::intel::esimd; + if (glu_op == GGML_GLU_OP_SWIGLU) { + return gate / (1.0f + exp(-gate)); + } + // GEGLU, the tanh is written with exp + simd a = 0.79788456f * gate * (1.0f + 0.044715f * gate * gate); + simd t = 1.0f - 2.0f / (exp(2.0f * a) + 1.0f); + return 0.5f * gate * (1.0f + t); +} + +// activations come in through 2D block loads, one for the (d, s) pairs of a super-block (s is unused) and one per two +// groups for the quants, the DPAS A operand needs no repacking. Groups with two scales (16 wide) mask the +// weights instead, zeros in B drop the other half of K. +// KS is the most threads that split K for one tile, the launch may use fewer. +// FUSED computes up * act(gate) from the weights vx (up) and vg (gate), reading the activations once. +template +ESIMD_INLINE void xmx_mul_mat(const void * vx, + const void * vg, + const char * vy, + float * dst, + const int ncols, + const int nrows, + const int ncv, + const int stride_y_bytes, + const int ldd, + const int glu_op, + const int tile, + const int lid, + const int ks) { + using namespace sycl::ext::intel::esimd; + + constexpr int NT = XMX_ROWS; + constexpr int RC = NC < 8 ? NC : 8; + constexpr int NB = NC / RC; + constexpr int N_ACC = FUSED ? 2 : 1; + constexpr bool pair = (RC * 32) % 64 == 0; // a block of RC rows fills whole registers + + if constexpr (KS > 1) { + slm_init(); + } + + const int bpr = ncols / QK_K; + const int row0 = tile * NT; + const auto c = T::make(vx, nrows, bpr, row0); + const auto cg = T::make(FUSED ? vg : vx, nrows, bpr, row0); + + simd rows; +#pragma unroll + for (int n = 0; n < NT; ++n) { + rows[n] = row0 + n < nrows ? row0 + n : nrows - 1; + } + + const uint32_t * yq = (const uint32_t *) vy; + const uint32_t * yd = (const uint32_t *) (vy + ncols); + const unsigned yqw = ncols - 1; + const unsigned ydw = ncols / QK8_1 * sizeof(sycl::half2) - 1; + const unsigned yh = ncv - 1; + const unsigned yp = stride_y_bytes - 1; + + simd Cf = 0.0f; + simd Cfg = 0.0f; // gate, only used when fused + simd Bl = int8_t(0); + simd Bh = int8_t(0); + simd Ol = int8_t(0); + if constexpr (T::has_min) { + Ol = int8_t(1); + } + + for (int sb = lid; sb < bpr; sb += ks) { + auto sc = T::load_scales(c, sb, rows); + typename T::scales_t scg; + if constexpr (FUSED) { + scg = T::load_scales(cg, sb, rows); + } + + // d of the 8 groups of the super-block, [column * 8 + group] + simd d8f[NB]; +#pragma unroll + for (int b = 0; b < NB; ++b) { + simd ds = load_2d(yd, ydw, yh, yp, sb * 8, b * RC); + simd dh = ds.template bit_cast_view(); + d8f[b] = convert(simd(dh.template select<8 * RC, 2>(0))); + } + +#pragma unroll + for (int part = 0; part < 2; ++part) { + const auto pd = T::load_part(c, sb, part); + typename T::part_t pdg; + if constexpr (FUSED) { + pdg = T::load_part(cg, sb, part); + } + + simd A2[NB]; // quants of two groups + +#pragma unroll + for (int g = 0; g < 4; ++g) { + const int gi = 4 * part + g; + if constexpr (pair) { + if (g % 2 == 0) { +#pragma unroll + for (int b = 0; b < NB; ++b) { + A2[b] = + load_2d(yq, yqw, yh, yp, sb * 64 + gi * 8, b * RC); + } + } + } else { +#pragma unroll + for (int b = 0; b < NB; ++b) { + A2[b].template select(0) = + load_2d(yq, yqw, yh, yp, sb * 64 + gi * 8, b * RC); + } + } + + xmx_group(Cf, pd, sc, A2, d8f, Bl, Bh, Ol, g, part); + if constexpr (FUSED) { + xmx_group(Cfg, pdg, scg, A2, d8f, Bl, Bh, Ol, g, part); + } + } + } + } + + if constexpr (KS > 1) { + slm_block_store(lid * NC * NT * sizeof(float), Cf); + if constexpr (FUSED) { + slm_block_store((KS + lid) * NC * NT * sizeof(float), Cfg); + } + barrier(); + if (lid != 0) { + return; + } + Cf = 0.0f; + Cfg = 0.0f; +#pragma unroll 4 + for (int j = 0; j < ks; ++j) { + Cf += slm_block_load(j * NC * NT * sizeof(float)); + if constexpr (FUSED) { + Cfg += slm_block_load((KS + j) * NC * NT * sizeof(float)); + } + } + } + + if constexpr (FUSED) { + Cf *= xmx_glu_act(Cfg, glu_op); + } + + if (row0 + NT <= nrows && ldd % 4 == 0 && (uintptr_t) dst % 16 == 0) { +#pragma unroll + for (int m = 0; m < NC; ++m) { + if (NC <= 8 || m < ncv) { + block_store(dst + (size_t) m * ldd + row0, Cf.template select(m * NT)); + } + } + } else { + // tail tile or unaligned output + simd lane(0, 1); + simd_mask in_rows = (lane + row0) < (uint32_t) nrows; +#pragma unroll + for (int m = 0; m < NC; ++m) { + if (NC <= 8 || m < ncv) { + scatter(dst + (size_t) m * ldd + row0, lane * (uint32_t) sizeof(float), + Cf.template select(m * NT), in_rows); + } + } + } +} + +// KS_MAX bounds the threads that split K for one tile, the launch uses as few as keep every thread to the same +// number of super-blocks as the busiest one +template +static void xmx_mul_mat_launch_ks(const void * vx, + const void * vg, + const char * vy, + float * dst, + const int ncols, + const int nrows, + const int ncv, + const int stride_y_bytes, + const int ldd, + const int glu_op, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK_K == 0); + const size_t n_tiles = (nrows + XMX_ROWS - 1) / XMX_ROWS; + const int bpr = ncols / QK_K; + const int rounds = (bpr + KS_MAX - 1) / KS_MAX; + const int ks = (bpr + rounds - 1) / rounds; + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(n_tiles * ks), sycl::range<1>(ks)), + sycl::ext::oneapi::experimental::properties{ + sycl::ext::intel::experimental::grf_size()> }, + [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { + xmx_mul_mat(vx, vg, vy, dst, ncols, nrows, ncv, stride_y_bytes, ldd, + glu_op, (int) it.get_group(0), (int) it.get_local_id(0), + ks); + }); + }); +} + +template +static void xmx_mul_mat_launch(const void * vx, + const void * vg, + const char * vy, + float * dst, + const int ncols, + const int nrows, + const int ncv, + const int stride_y_bytes, + const int ldd, + const int glu_op, + dpct::queue_ptr stream) { + // the fused kernel keeps two sets of accumulators in shared memory, half the threads for the same size + constexpr int ks_max = xmx_n_split() / (FUSED ? 2 : 1); + xmx_mul_mat_launch_ks(vx, vg, vy, dst, ncols, nrows, ncv, stride_y_bytes, ldd, glu_op, + stream); +} + +// FUSED: vx and vg are the up and gate weights, dst gets up * act(gate) with glu_op a ggml_glu_op +template +static void xmx_mul_mat_ncols(const void * vx, + const void * vg, + const char * vy, + float * dst, + const int ncols, + const int nrows, + const int ncols_dst, + const int stride_y_bytes, + const int ldd, + const int glu_op, + dpct::queue_ptr stream) { + constexpr int chunk_max = FUSED ? XMX_FUSED_CHUNK : XMX_CHUNK; + for (int c0 = 0; c0 < ncols_dst; c0 += chunk_max) { + const int n = std::min(chunk_max, ncols_dst - c0); + const char * y = vy + (size_t) c0 * stride_y_bytes; + float * d = dst + (size_t) c0 * ldd; + const int nc = n > 32 ? 64 : (n > 16 ? 32 : (n > 8 ? 16 : n)); + switch (nc) { +#define XMX_CASE(N) \ + case N: \ + xmx_mul_mat_launch(vx, vg, y, d, ncols, nrows, n, stride_y_bytes, ldd, glu_op, stream); \ + break + XMX_CASE(1); + XMX_CASE(2); + XMX_CASE(3); + XMX_CASE(4); + XMX_CASE(5); + XMX_CASE(6); + XMX_CASE(7); + XMX_CASE(8); +#undef XMX_CASE + // short row variants are for tests, not built for the wide tiles +#define XMX_CASE_WIDE(N) \ + case N: \ + if constexpr (!T::short_rows && (!FUSED || N <= XMX_FUSED_CHUNK)) { \ + xmx_mul_mat_launch(vx, vg, y, d, ncols, nrows, n, stride_y_bytes, ldd, glu_op, stream); \ + } \ + break + XMX_CASE_WIDE(16); + XMX_CASE_WIDE(32); + XMX_CASE_WIDE(64); +#undef XMX_CASE_WIDE + default: + GGML_ABORT("unsupported ncols_dst=%d for XMX MMV", n); + } + } +} + } // namespace ggml_sycl_esimd #endif // GGML_SYCL_ESIMD_HPP diff --git a/ggml/src/ggml-sycl/fusion.cpp b/ggml/src/ggml-sycl/fusion.cpp index cb4274e79..9b553829b 100644 --- a/ggml/src/ggml-sycl/fusion.cpp +++ b/ggml/src/ggml-sycl/fusion.cpp @@ -37,6 +37,13 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg const bool reorder_pair = wu->type == GGML_TYPE_Q4_K && wg->type == GGML_TYPE_Q4_K; const bool plain_pair = (wu->type == GGML_TYPE_Q5_K || wu->type == GGML_TYPE_IQ4_XS) && (wg->type == GGML_TYPE_Q5_K || wg->type == GGML_TYPE_IQ4_XS); +#ifdef GGML_SYCL_MMVQ_HAS_XMX + // same-type pairs of the types of the XMX kernels, which also take a few more columns + // (the XMX kernels walk whole QK_K super-blocks, so they need rows of a multiple of QK_K) + const bool xmx_pair = wu->type == wg->type && ggml_sycl_xmx_supports_type(wu->type) && wu->ne[0] % QK_K == 0; +#else + const bool xmx_pair = false; +#endif // GGML_SYCL_MMVQ_HAS_XMX // the ESIMD kernel covers same-type q2_K..q6_K and q8_0, single column only (as unfused DMMV) #ifdef GGML_SYCL_DMMV_HAS_ESIMD const bool esimd_pair = g_ggml_sycl_enable_esimd && wu->type == wg->type && act->ne[1] == 1 && @@ -47,7 +54,7 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg #endif // q8_0 rows are walked in 32-wide blocks, so they only need to be a multiple of QK8_0 const int64_t row_align = esimd_pair && wu->type == GGML_TYPE_Q8_0 ? QK8_0 : QK_K; - if ((!reorder_pair && !plain_pair && !esimd_pair) || wu->ne[0] % row_align != 0) { + if ((!reorder_pair && !plain_pair && !xmx_pair && !esimd_pair) || wu->ne[0] % row_align != 0) { return false; } @@ -67,7 +74,8 @@ static bool ggml_sycl_should_fuse_mul_mat_glu(const ggml_tensor * gate, const gg return false; } // mat-vec only: one column per decoded token, up to the batch the reorder kernels cover - if (act->ne[1] > MMVQ_MAX_BATCH_SIZE) { + const int max_cols = xmx_pair ? std::max(MMVQ_MAX_BATCH_SIZE, g_ggml_sycl_xmx_glu_max_cols) : MMVQ_MAX_BATCH_SIZE; + if (act->ne[1] > max_cols) { return false; } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index bfe53c5d7..e07a18bdc 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -103,6 +103,9 @@ int g_ggml_sycl_enable_vmm = 1; int g_ggml_sycl_enable_fusion = 1; int g_ggml_sycl_enable_esimd = 1; int g_ggml_sycl_mmvq_wide = 1; +int g_ggml_sycl_enable_xmx_esimd = 1; +int g_ggml_sycl_xmx_max_cols = GGML_SYCL_XMX_MAX_COLS_DEFAULT; +int g_ggml_sycl_xmx_glu_max_cols = GGML_SYCL_XMX_GLU_MAX_COLS; int g_ggml_sycl_prioritize_dmmv = 0; int g_ggml_sycl_xmx_gather_types = GGML_SYCL_XMX_GATHER_TYPES_DEFAULT; int g_ggml_sycl_xmx_gather_shapes = GGML_SYCL_XMX_GATHER_SHAPES_DEFAULT; @@ -424,6 +427,10 @@ static void ggml_check_sycl() try { g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1); g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1); g_ggml_sycl_mmvq_wide = ggml_sycl_get_env("GGML_SYCL_MMVQ_WIDE", 1); + g_ggml_sycl_enable_xmx_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_XMX_ESIMD", 1); + g_ggml_sycl_xmx_max_cols = std::max(0, ggml_sycl_get_env("GGML_SYCL_XMX_MAX_COLS", GGML_SYCL_XMX_MAX_COLS_DEFAULT)); + g_ggml_sycl_xmx_glu_max_cols = + std::clamp(ggml_sycl_get_env("GGML_SYCL_XMX_GLU_MAX_COLS", GGML_SYCL_XMX_GLU_MAX_COLS), 0, GGML_SYCL_XMX_GLU_MAX_COLS); g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0); g_ggml_sycl_xmx_gather_types = ggml_sycl_get_env("GGML_SYCL_XMX_GATHER_TYPES", GGML_SYCL_XMX_GATHER_TYPES_DEFAULT); g_ggml_sycl_xmx_gather_shapes = ggml_sycl_get_env("GGML_SYCL_XMX_GATHER_SHAPES", GGML_SYCL_XMX_GATHER_SHAPES_DEFAULT); @@ -561,7 +568,13 @@ static void ggml_check_sycl() try { #else GGML_LOG_INFO(" GGML_SYCL_ENABLE_ESIMD: %d disabled by compile flag\n", g_ggml_sycl_enable_esimd); #endif + GGML_LOG_INFO(" GGML_SYCL_MMVQ_WIDE: %d\n", g_ggml_sycl_mmvq_wide); + + GGML_LOG_INFO(" GGML_SYCL_ENABLE_XMX_ESIMD: %d\n", g_ggml_sycl_enable_xmx_esimd); + GGML_LOG_INFO(" GGML_SYCL_XMX_MAX_COLS: %d\n", g_ggml_sycl_xmx_max_cols); + GGML_LOG_INFO(" GGML_SYCL_XMX_GLU_MAX_COLS: %d\n", g_ggml_sycl_xmx_glu_max_cols); + GGML_LOG_INFO(" GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv); g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1); @@ -5165,9 +5178,24 @@ static bool can_use_dequantize_mul_mat_vec(const ggml_tensor * src0, const ggml_ src0->ne[0] % dmmv_x_required == 0 && src1->ne[1] == 1; } -static bool can_use_mul_mat_vec_q(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { +// reordered weights of the types in ggml_sycl_xmx_supports_type() take the XMX kernel for wider batches +static bool can_use_xmx_batch(int device, const ggml_tensor * src0, const ggml_tensor * src1) { +#ifdef GGML_SYCL_MMVQ_HAS_XMX + const auto * extra = static_cast(src0->extra); + return ggml_sycl_xmx_enabled(device) && ggml_sycl_xmx_supports_type(src0->type) && extra && + extra->optimized_feature.reorder && src1->ne[1] <= g_ggml_sycl_xmx_max_cols && src1->ne[2] == 1 && + src1->ne[3] == 1; +#else + GGML_UNUSED(device); + GGML_UNUSED(src0); + GGML_UNUSED(src1); + return false; +#endif // GGML_SYCL_MMVQ_HAS_XMX +} + +static bool can_use_mul_mat_vec_q(int device, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { return ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && - src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; + (src1->ne[1] <= MMVQ_MAX_BATCH_SIZE || can_use_xmx_batch(device, src0, src1)); } static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) { @@ -5207,12 +5235,15 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor // check data types and tensor shapes for custom matrix multiplication kernels: bool use_dequantize_mul_mat_vec = can_use_dequantize_mul_mat_vec(src0, src1, dst); + if (src1->ne[1] == 1 && ggml_sycl_xmx_min_cols(src0->type) == 1 && can_use_xmx_batch(ctx.device, src0, src1)) { + use_dequantize_mul_mat_vec = false; + } #ifdef GGML_SYCL_F16 // dmmv may convert src1 to f16 in this build [TAG_GGML_PREC] use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && ggml_sycl_src1_f16_ok(dst); #endif - bool use_mul_mat_vec_q = can_use_mul_mat_vec_q(src0, src1, dst); + bool use_mul_mat_vec_q = can_use_mul_mat_vec_q(ctx.device, src0, src1, dst); bool use_mul_mat_q = ggml_sycl_supports_mmq(src0->type) && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32; @@ -5321,6 +5352,35 @@ static bool ggml_sycl_mul_mat_glu_mmvq_plain(ggml_backend_sycl_context & ctx, gg GGML_UNUSED(gate); } +// Gate and up weights of one type on the XMX engines: one launch reads the activations once and applies the GLU. +// Both weights must already be in the reorder layout. Returns false if the kernel declined. +static bool ggml_sycl_mul_mat_glu_xmx(ggml_backend_sycl_context & ctx, + ggml_tensor * glu, + ggml_tensor * up, + const ggml_tensor * wu, + const ggml_tensor * wg, + const ggml_tensor * act) { + const int64_t ne00 = wu->ne[0]; + const int64_t ne11 = act->ne[1]; + + const queue_ptr stream = ctx.stream(); + const int src1_padded_cols = GGML_PAD((int) ne00, MATRIX_ROW_PADDING); + const int stride_y_bytes = src1_padded_cols * (int) sizeof(block_q8_1) / QK8_1; + + // log the up mat-mul: glu's own srcs are the two intermediates the fusion never materialises + scope_op_debug_print scope_dbg_print(__func__, up, /*num_src=*/2, " : fused with gate + GLU (XMX)"); + + ggml_sycl_pool_alloc src1_q8_alloc(ctx.pool(), (size_t) ne11 * stride_y_bytes); + char * src1_ddq = src1_q8_alloc.get(); + + quantize_row_q8_1_sycl((const float *) act->data, src1_ddq, (int) ne00, (int) ne11, + src1_padded_cols, stream); + + return ggml_sycl_mul_mat_vec_q_glu_xmx(ctx.device, wu->type, ggml_get_glu_op(glu), wu->data, wg->data, src1_ddq, + (float *) glu->data, (int) ne00, (int) wu->ne[1], (int) ne11, stride_y_bytes, + /*stride_col_dst=*/(int) glu->ne[0], stream); +} + #ifdef GGML_SYCL_DMMV_HAS_ESIMD // {mul_mat(gate), mul_mat(up), GLU} on the reorder layout via the ESIMD DMMV kernel, for // same-type K-quant or Q8_0 pairs and a single activation column. It reads the F32 activation @@ -5389,11 +5449,33 @@ static bool ggml_sycl_mul_mat_glu_fused(ggml_backend_sycl_context & ctx, ggml_cg return false; } + // gate and up of one type the XMX kernel handles, it needs the reorder layout like the unfused mmvq path + if (wg->type == wu->type && gate->src[1] == act && act->ne[1] >= ggml_sycl_xmx_min_cols(wu->type) && + act->ne[1] <= g_ggml_sycl_xmx_glu_max_cols) { + opt_for_reorder(&ctx, wu, act, up, mul_mat_algo::MMVQ); + opt_for_reorder(&ctx, wg, act, gate, mul_mat_algo::MMVQ); + if (can_use_xmx_batch(ctx.device, wu, act) && can_use_xmx_batch(ctx.device, wg, act) && + ggml_sycl_mul_mat_glu_xmx(ctx, glu, up, wu, wg, act)) { + return true; + } + } + + // the kernels below take up to MMVQ_MAX_BATCH_SIZE columns, a wider batch got here for the XMX kernel + if (act->ne[1] > MMVQ_MAX_BATCH_SIZE) { + return false; + } + // quant pairs the reorder kernel does not serve (mixed gate/up types, q5_K off BMG) take the // standard-layout fused path instead; same-type q4_K / q5_K keep the reorder path below const bool reorder_pair = wg->type == wu->type && (wu->type == GGML_TYPE_Q4_K || (wu->type == GGML_TYPE_Q5_K && ggml_sycl_q5_k_mmvq_reuse(ctx.device))); if (!reorder_pair) { + const auto is_plain_type = [](ggml_type type) { + return type == GGML_TYPE_Q5_K || type == GGML_TYPE_IQ4_XS; + }; + if (!is_plain_type(wu->type) || !is_plain_type(wg->type)) { + return false; + } return ggml_sycl_mul_mat_glu_mmvq_plain(ctx, glu, gate, up, wu, wg, act); } diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index e9057777c..ccb0363bd 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -6,6 +6,10 @@ #include "quants.hpp" #include "vecdotq.hpp" +#ifdef GGML_SYCL_MMVQ_HAS_XMX +# include "esimd.hpp" +#endif // GGML_SYCL_MMVQ_HAS_XMX + // vec_dot_q_sycl_t adapters for the IQ vec_dots that take their codebook tables as extra // arguments: bind the constant tables here (as vec_dot_iq2_s_q8_1 / vec_dot_iq1_m_q8_1 already do // internally) so they can be used as template arguments of mul_mat_vec_q_moe. @@ -2557,6 +2561,158 @@ static void mul_mat_vec_iq4_xs_q8_1_sycl_switch_ncols( } } +#ifdef GGML_SYCL_MMVQ_HAS_XMX +// vgate is the gate weight of the fused gate and up kernel, nullptr for a plain mat-vec +template