From b90120e2f87262319743f0511571055802fa2cf9 Mon Sep 17 00:00:00 2001 From: Daniel Wymark Date: Wed, 23 Sep 2026 22:53:07 -0700 Subject: [PATCH] sycl: explicit-SIMD decode kernels for PTQ1_0 and PQ2_0 On Intel GPUs the MMVQ path for both ternary formats now runs an ESIMD kernel that gives each work-item one weight row, spreads the row's blocks over 16 lanes, and accumulates q8_1 dot products with dp4a. The kernels are compiled under the Intel compiler only and follow GGML_SYCL_ENABLE_ESIMD, so setting it to 0 restores the sub-group kernels for comparison. A one-row GET_ROWS with a single index becomes a vectorized copy, which is the shape the recurrent-state layers issue every token. Co-Authored-By: Claude Fable 5.1 Co-Authored-By: Claude Opus 5.5 (1M context) Claude-Session: https://claude.ai/code/session_01DSrpF2XYi9dav1e3hSAkaS Claude-Session: https://claude.ai/code/session_014smgAQnqKxMyRYQbXnvBgs --- ggml/src/ggml-sycl/getrows.cpp | 27 +++++ ggml/src/ggml-sycl/mmvq.cpp | 196 +++++++++++++++++++++++++++++++++ 2 files changed, 223 insertions(+) diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 36f840e6f5d1..24abe8cf13e7 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -214,6 +214,33 @@ static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tens GGML_TENSOR_BINARY_OP_LOCALS + if constexpr (std::is_same_v && std::is_same_v) { + // A one-row source selected by a single index is a copy; the only valid index is zero, so src1 is + // never read. Recurrent-state layers issue this shape once per layer per token. + if (ne01 == 1 && ne02 == 1 && ne03 == 1 && ggml_nelements(src1) == 1 && + ggml_is_contiguous(src0) && ggml_is_contiguous(dst)) { + if (dst_dd == src0_dd) { + return; + } + const int64_t vectors = (ne00 + 3) / 4; + const sycl::range<1> local(SYCL_GET_ROWS_BLOCK_SIZE); + const sycl::range<1> global(((vectors + local[0] - 1) / local[0]) * local[0]); + stream->parallel_for(sycl::nd_range<1>(global, local), [=](sycl::nd_item<1> item) { + const int64_t i = item.get_global_id(0); + if (4 * i + 3 < ne00) { + sycl::vec values; + values.load(i, src0_dd); + values.store(i, dst_dd); + } else { + for (int64_t j = 4 * i; j < ne00; ++j) { + dst_dd[j] = src0_dd[j]; + } + } + }); + return; + } + } + const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE); const int block_num_x = (ne00 + SYCL_GET_ROWS_BLOCK_SIZE - 1) / SYCL_GET_ROWS_BLOCK_SIZE; const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x); diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 36c4c1778769..1f46bf3fbd1c 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -6,6 +6,19 @@ #include "quants.hpp" #include "vecdotq.hpp" +#if defined(__INTEL_LLVM_COMPILER) + #include + #define GGML_SYCL_MMVQ_HAS_ESIMD +#endif + +#ifdef GGML_SYCL_MMVQ_HAS_ESIMD +// The explicit-SIMD decode kernels give each work-item one weight row and spread that row's blocks +// across the vector lanes; a work-group holds four rows. Both read q8_1 activations and accumulate +// integer dot products per lane, scaling to FP32 once per block. +constexpr int MMVQ_ESIMD_LANES = 16; +constexpr int MMVQ_ESIMD_ROWS_PER_WG = 4; +#endif + template static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __restrict__ vy, float * __restrict__ dst, const int ncols, const int nrows, const sycl::nd_item<3> & nd_item) { @@ -1278,11 +1291,113 @@ static void mul_mat_vec_q1_0_q8_1_sycl_switch_ncols( } } +#ifdef GGML_SYCL_MMVQ_HAS_ESIMD +// Each lane owns one 128-value PTQ1_0 block and its four q8_1 activation blocks. The five trits packed in +// a byte come out through the multiply-by-three recurrence, two bytes per 32-bit half at a time. +static void mul_mat_vec_ptq1_0_q8_1_esimd(const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, dpct::queue_ptr stream) { + constexpr int rows = MMVQ_ESIMD_ROWS_PER_WG; + const sycl::range<1> local(rows); + const sycl::range<1> global(((nrows + rows - 1) / rows) * rows); + stream->parallel_for(sycl::nd_range<1>(global, local), + [=](sycl::nd_item<1> item) [[intel::sycl_explicit_simd]] { + using namespace sycl::ext::intel::esimd; + constexpr int lanes = MMVQ_ESIMD_LANES; + const int row = item.get_global_id(0); + if (row >= nrows) return; + const uint8_t * weights = static_cast(vx) + + static_cast(row) * (ncols / QK_PTQ1_0) * sizeof(block_ptq1_0); + const uint8_t * input = static_cast(vy); + const simd lane(0, 1); + simd acc = 0.0f; + for (int i = 0; i < ncols / QK_PTQ1_0; i += lanes) { + const simd block = lane + i; + const simd_mask valid = block < static_cast(ncols / QK_PTQ1_0); + const simd wb = block * sizeof(block_ptq1_0); + const simd scale_offsets = wb + offsetof(block_ptq1_0, d); + const simd wd = gather( + reinterpret_cast(weights), scale_offsets, valid, simd(sycl::half(0))); + simd input_words[4]; + simd input_scales[4]; + simd sums[4]; +#pragma unroll + for (int chunk = 0; chunk < 4; ++chunk) { + const simd ab = (block * 4 + chunk) * sizeof(block_q8_1); + const simd aq = ab + 4; + input_words[chunk] = gather(reinterpret_cast(input), + aq, valid, simd(0)); + input_scales[chunk] = gather( + reinterpret_cast(input), ab, valid, simd(sycl::half(0))); + sums[chunk] = 0; + } + simd first_words = gather( + reinterpret_cast(weights), wb, valid, simd(0)); + const simd last_offsets = wb + 16; + simd last_words = gather( + reinterpret_cast(weights), last_offsets, valid, simd(0)); +#pragma unroll + for (int group = 0; group < 6; ++group) { + simd packed; + if (group < 4) { + packed = first_words.select(group * lanes); + } else { + packed = last_words.select((group - 4) * lanes); + } + simd v_lo = (packed & 0xff) | ((packed & 0xff00) << 8); + simd v_hi = ((packed >> 16) & 0xff) | ((packed & 0xff000000) >> 8); +#pragma unroll + for (int trit = 0; trit < 5; ++trit) { + const simd w_lo = v_lo * 3; + const simd w_hi = v_hi * 3; + v_lo = w_lo & 0x00ff00ff; v_hi = w_hi & 0x00ff00ff; + simd values = ((w_lo >> 8) & 0xff) | ((w_lo >> 16) & 0xff00) | + ((w_hi << 8) & 0xff0000) | (w_hi & 0xff000000); + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.bit_cast_view(); + const int element = group < 4 ? trit * 16 + 4 * group : 80 + trit * 8 + 4 * (group - 4); + const int chunk = element / 32; + const simd a = input_words[chunk].select((element % 32) / 4 * lanes); + sums[chunk] = dp4a(sums[chunk], w, a); + } + } + const simd tail_offset = wb + offsetof(block_ptq1_0, qh); + const simd tail = gather( + reinterpret_cast(weights), tail_offset, valid, simd(0)); + simd v = (tail & 0xff) | ((tail & 0xff00) << 8); +#pragma unroll + for (int trit = 0; trit < 4; trit += 2) { + const simd w0 = v * 3; v = w0 & 0x00ff00ff; + const simd w1 = v * 3; v = w1 & 0x00ff00ff; + simd values = ((w0 >> 8) & 0xff) | ((w0 >> 16) & 0xff00) | + ((w1 << 8) & 0xff0000) | (w1 & 0xff000000); + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.bit_cast_view(); + const simd a = input_words[3].select((6 + trit / 2) * lanes); + sums[3] = dp4a(sums[3], w, a); + } + simd weighted_sum = 0.0f; +#pragma unroll + for (int chunk = 0; chunk < 4; ++chunk) { + weighted_sum += input_scales[chunk] * simd(sums[chunk]); + } + acc += simd(wd) * weighted_sum; + } + dst[row] = reduce(acc, std::plus<>{}); + }); +} +#endif + static void mul_mat_vec_ptq1_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_PTQ1_0 == 0); +#ifdef GGML_SYCL_MMVQ_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + mul_mat_vec_ptq1_0_q8_1_esimd(vx, vy, dst, ncols, nrows, stream); + return; + } +#endif const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE); @@ -1338,11 +1453,92 @@ static void mul_mat_vec_ptq1_0_q8_1_sycl_switch_ncols( } } +#ifdef GGML_SYCL_MMVQ_HAS_ESIMD +// Each lane owns one 32-value q8_1 activation block and the eight PQ2_0 bytes that pair with it, so four +// lanes share one 128-value weight block. Activations arrive in one gather per lane; weights arrive as +// aligned 32-bit words and are shifted into place when the lane's payload starts mid-word. +static void mul_mat_vec_pq2_0_q8_1_esimd(const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, dpct::queue_ptr stream) { + constexpr int lanes = MMVQ_ESIMD_LANES; + constexpr int rows = MMVQ_ESIMD_ROWS_PER_WG; + const sycl::range<1> local(rows); + const sycl::range<1> global(((nrows + rows - 1) / rows) * rows); + stream->parallel_for(sycl::nd_range<1>(global, local), + [=](sycl::nd_item<1> item) [[intel::sycl_explicit_simd]] { + using namespace sycl::ext::intel::esimd; + const int row = item.get_global_id(0); + if (row >= nrows) { + return; + } + const uint8_t * weights = static_cast(vx) + + static_cast(row) * (ncols / QK_PQ2_0) * sizeof(block_pq2_0); + const uint8_t * input = static_cast(vy); + simd acc = 0.0f; + const simd lane(0, 1); + const simd zero_half = sycl::half(0.0f); + for (int i = 0; i < ncols / QK8_1; i += lanes) { + const simd q = lane + i; + const simd_mask valid = q < static_cast(ncols / QK8_1); + const simd wb = (q / 4) * sizeof(block_pq2_0); + const simd wo = wb + 2 + (q % 4) * 8; + const simd ab = q * sizeof(block_q8_1); + const simd wd = gather( + reinterpret_cast(weights), wb, valid, zero_half); + const simd ad = gather( + reinterpret_cast(input), ab, valid, zero_half); + const simd input_offsets = ab + 4; + simd input_words = gather( + reinterpret_cast(input), input_offsets, valid, simd(0)); + + const simd alignment = (wo + uint32_t(reinterpret_cast(weights) & 3)) & 3; + const simd aligned_offsets = wo - alignment; + simd packed_weights = gather( + reinterpret_cast(weights), aligned_offsets, valid, simd(0)); + const simd_mask shifted = valid & (alignment == 2); + const simd tail_offsets = wo + 6; + const simd tail = gather( + reinterpret_cast(weights), tail_offsets, shifted, simd(0)); + simd first = packed_weights.select(0); + simd second = packed_weights.select(lanes); + first.merge((first >> 16) | (second << 16), shifted); + second.merge((second >> 16) | (tail << 16), shifted); + packed_weights.select(0) = first; + packed_weights.select(lanes) = second; + + simd sum = 0; +#pragma unroll + for (int j = 0; j < 4; ++j) { + const simd pair = + (packed_weights.select((j / 2) * lanes) >> ((j % 2) * 16)) & 0xffff; +#pragma unroll + for (int k = 0; k < 2; ++k) { + simd values = (pair >> (8 * k)) & 0xff; + values = (values | (values << 12)) & 0x000f000f; + values = (values | (values << 6)) & 0x03030303; + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.bit_cast_view(); + const simd a = input_words.select((2 * j + k) * lanes); + sum = dp4a(sum, w, a); + } + } + acc += simd(wd) * simd(ad) * simd(sum); + } + dst[row] = reduce(acc, std::plus<>{}); + }); +} +#endif + static void mul_mat_vec_pq2_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_PQ2_0 == 0); +#ifdef GGML_SYCL_MMVQ_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + mul_mat_vec_pq2_0_q8_1_esimd(vx, vy, dst, ncols, nrows, stream); + return; + } +#endif const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);