sycl : Add DMMV ESIMD Q3_K kernel (llama/26251)
* Add DMMV Q4_K and Q6_K ESIMD kernels Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable. Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Refactor ESIMD kernels to share common code Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Move control of ESIMD from compile to runtime Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Use ESIMD by default when available Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Fix possible error when using ESIMD by default While not an issue in the current version, this will become an issue when additional QK ESIMD kernels are added (such as Q2_K). Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Add explicit unroll to ESIMD kernels Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Tidy up ESIMD kernels a bit Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Add DMMV Q3_K ESIMD kernel Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
This commit is contained in:
Todd Malsbary
Georgi Gerganov
parent
f7011606ad
commit
90d4ed1859
|
|
@ -61,6 +61,7 @@ void ggml_sycl_host_free(void* ptr);
|
|||
extern int g_ggml_sycl_debug;
|
||||
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_prioritize_dmmv;
|
||||
extern int g_ggml_sycl_enable_flash_attention;
|
||||
extern int g_ggml_sycl_dev2dev_memcpy;
|
||||
|
|
|
|||
|
|
@ -8,6 +8,9 @@
|
|||
#include <sycl/ext/oneapi/bfloat16.hpp>
|
||||
#define GGML_SYCL_DMMV_HAS_BF16
|
||||
#endif
|
||||
#include <sycl/ext/intel/esimd.hpp>
|
||||
#include "esimd.hpp"
|
||||
#define GGML_SYCL_DMMV_HAS_ESIMD
|
||||
#endif
|
||||
|
||||
static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){
|
||||
|
|
@ -1864,6 +1867,113 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
|
|||
});
|
||||
}
|
||||
|
||||
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
|
||||
using ggml_sycl_esimd::GGML_SYCL_DMMV_ESIMD_WG_SIZE;
|
||||
|
||||
// generic reordered dequantize-matvec: each work-group owns a pair of
|
||||
// consecutive output rows and updates one 32-wide accumulator per row
|
||||
template <ggml_type T>
|
||||
ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd(
|
||||
const void * vx, const float * y, float * dst,
|
||||
const int ncols, const int nrows,
|
||||
sycl::local_accessor<float, 1> lmem,
|
||||
const sycl::nd_item<1> & it) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
using traits = ggml_sycl_esimd::esimd_reorder_q_traits<T>;
|
||||
|
||||
const int num_blocks_per_row = ncols / QK_K;
|
||||
const size_t nb = (size_t) nrows * num_blocks_per_row;
|
||||
const auto ps = traits::make_ptrs(vx, nb);
|
||||
|
||||
const int tid = it.get_local_id(0);
|
||||
const int row_pair = it.get_group(0);
|
||||
const int row0 = row_pair * 2; // two consecutive output rows
|
||||
const bool has_row1 = row0 + 1 < nrows;
|
||||
|
||||
// one 32-wide accumulator per output row (small footprint, no spill)
|
||||
simd<float, 32> acc0 = 0.0f;
|
||||
simd<float, 32> acc1 = 0.0f;
|
||||
|
||||
for (int ib = tid; ib < num_blocks_per_row; ib += GGML_SYCL_DMMV_ESIMD_WG_SIZE) {
|
||||
simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) ib * QK_K);
|
||||
|
||||
const size_t bi0 = (size_t) (row0 + 0) * num_blocks_per_row + ib;
|
||||
const size_t bi1 = (size_t) (row0 + 1) * num_blocks_per_row + ib;
|
||||
|
||||
traits::mac_pair(ps, bi0, ps, bi1, has_row1, y_vec, acc0, acc1);
|
||||
}
|
||||
|
||||
lmem[tid * 2 + 0] = reduce<float>(acc0, std::plus<>{});
|
||||
lmem[tid * 2 + 1] = reduce<float>(acc1, std::plus<>{});
|
||||
it.barrier(sycl::access::fence_space::local_space);
|
||||
|
||||
if (tid == 0) {
|
||||
float sum0 = 0.0f;
|
||||
float sum1 = 0.0f;
|
||||
for (int p = 0; p < GGML_SYCL_DMMV_ESIMD_WG_SIZE; ++p) {
|
||||
sum0 += lmem[p * 2 + 0];
|
||||
sum1 += lmem[p * 2 + 1];
|
||||
}
|
||||
dst[row0 + 0] = sum0;
|
||||
if (has_row1) {
|
||||
dst[row0 + 1] = sum1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const float *y,
|
||||
float *dst, const int ncols,
|
||||
const int nrows,
|
||||
dpct::queue_ptr stream) {
|
||||
GGML_ASSERT(ncols % QK_K == 0);
|
||||
const int workgroups = (nrows + 1) / 2;
|
||||
stream->submit([&](sycl::handler &h) {
|
||||
sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
|
||||
h.parallel_for(
|
||||
sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
|
||||
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
|
||||
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q3_K>(
|
||||
vx, y, dst, ncols, nrows, lmem, it);
|
||||
});
|
||||
});
|
||||
}
|
||||
|
||||
static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const float *y,
|
||||
float *dst, const int ncols,
|
||||
const int nrows,
|
||||
dpct::queue_ptr stream) {
|
||||
GGML_ASSERT(ncols % QK_K == 0);
|
||||
const int workgroups = (nrows + 1) / 2;
|
||||
stream->submit([&](sycl::handler &h) {
|
||||
sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
|
||||
h.parallel_for(
|
||||
sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
|
||||
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
|
||||
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q4_K>(
|
||||
vx, y, dst, ncols, nrows, lmem, it);
|
||||
});
|
||||
});
|
||||
}
|
||||
|
||||
static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const float *y,
|
||||
float *dst, const int ncols,
|
||||
const int nrows,
|
||||
dpct::queue_ptr stream) {
|
||||
GGML_ASSERT(ncols % QK_K == 0);
|
||||
const int workgroups = (nrows + 1) / 2;
|
||||
stream->submit([&](sycl::handler &h) {
|
||||
sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
|
||||
h.parallel_for(
|
||||
sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)),
|
||||
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
|
||||
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q6_K>(
|
||||
vx, y, dst, ncols, nrows, lmem, it);
|
||||
});
|
||||
});
|
||||
}
|
||||
|
||||
#endif // GGML_SYCL_DMMV_HAS_ESIMD
|
||||
|
||||
static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
|
||||
float *dst, const int ncols,
|
||||
const int nrows,
|
||||
|
|
@ -1992,7 +2102,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
|
|||
case GGML_TYPE_Q3_K:
|
||||
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
|
||||
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
|
||||
dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
|
||||
if (g_ggml_sycl_enable_esimd) {
|
||||
dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
else
|
||||
#endif
|
||||
{
|
||||
dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
} else {
|
||||
dequantize_mul_mat_vec_q3_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
|
|
@ -2000,7 +2118,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
|
|||
case GGML_TYPE_Q4_K:
|
||||
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
|
||||
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
|
||||
dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
|
||||
if (g_ggml_sycl_enable_esimd) {
|
||||
dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
else
|
||||
#endif
|
||||
{
|
||||
dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
} else {
|
||||
dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
|
|
@ -2016,7 +2142,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
|
|||
case GGML_TYPE_Q6_K:
|
||||
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
|
||||
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
|
||||
dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
|
||||
if (g_ggml_sycl_enable_esimd) {
|
||||
dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
else
|
||||
#endif
|
||||
{
|
||||
dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
} else {
|
||||
dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -0,0 +1,392 @@
|
|||
//
|
||||
// MIT license
|
||||
// Copyright (C) 2026 Intel Corporation
|
||||
// SPDX-License-Identifier: MIT
|
||||
//
|
||||
|
||||
//
|
||||
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
|
||||
// See https://llvm.org/LICENSE.txt for license information.
|
||||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
|
||||
//
|
||||
|
||||
#ifndef GGML_SYCL_ESIMD_HPP
|
||||
#define GGML_SYCL_ESIMD_HPP
|
||||
|
||||
#include <sycl/ext/intel/esimd.hpp>
|
||||
|
||||
#include "common.hpp"
|
||||
|
||||
namespace ggml_sycl_esimd {
|
||||
|
||||
constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4;
|
||||
|
||||
//
|
||||
// Shared ESIMD building blocks for the reordered K-quant dequantize-matvec
|
||||
// kernels.
|
||||
//
|
||||
// The reordered K-quant ESIMD matvec kernels share one skeleton: per super-block,
|
||||
// load a 256-float activation slice, load one weight block, dequantize it into 8
|
||||
// chunks of 32 and MAC each chunk against the matching activation slice, then
|
||||
// reduce and run a lane-0 epilogue.
|
||||
//
|
||||
// Each K-quant kernel emits exactly 8 chunks of 32 mapping to activation slices
|
||||
// 0..7, so the per-block work is captured by esimd_reorder_q_traits<T>::mac_pair,
|
||||
// which dequantizes two weight blocks and MACs both against a shared activation
|
||||
// vector with the two FMA chains interleaved (co-scheduled to hide FMA latency).
|
||||
// The "pair" is the (row0,row1) row pair owned by one work-group, so the
|
||||
// layout+dequant is written once per quant type here.
|
||||
//
|
||||
|
||||
template <ggml_type T> struct esimd_reorder_q_traits;
|
||||
|
||||
// build a 32-lane vector whose low 16 lanes are `lo` and high 16 are `hi`
|
||||
// (a super-chunk splits into two 16-wide halves with distinct scale/min codes).
|
||||
static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 32> splat_lo_hi(float lo, float hi) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
simd<float, 32> v;
|
||||
v.select<16, 1>(0) = lo;
|
||||
v.select<16, 1>(16) = hi;
|
||||
return v;
|
||||
}
|
||||
|
||||
// unpack one block of Q4_K/Q5_K scale/min codes (get_scale_min_k4 layout) into 8
|
||||
// float scales (dall * sc) and 8 float mins (-dmin * m); the min carries the
|
||||
// negation so the dequant epilogue adds.
|
||||
static ESIMD_INLINE void unpack_scale_min_k4(
|
||||
sycl::ext::intel::esimd::simd<uint8_t, 12> scales, float dall, float dmin,
|
||||
sycl::ext::intel::esimd::simd<float, 8> & scale_f,
|
||||
sycl::ext::intel::esimd::simd<float, 8> & min_f) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
simd<uint8_t, 8> sc = 0;
|
||||
simd<uint8_t, 8> m = 0;
|
||||
simd<uint8_t, 4> scale_lo = scales.select<4, 1>(0);
|
||||
simd<uint8_t, 4> min_lo = scales.select<4, 1>(4);
|
||||
simd<uint8_t, 4> hi_bits = scales.select<4, 1>(8);
|
||||
sc.select<4, 1>(0) = scale_lo & simd<uint8_t, 4>(0x3F);
|
||||
sc.select<4, 1>(4) = (hi_bits & simd<uint8_t, 4>(0x0F)) |
|
||||
((scale_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
|
||||
m.select<4, 1>(0) = min_lo & simd<uint8_t, 4>(0x3F);
|
||||
m.select<4, 1>(4) = (hi_bits >> simd<uint8_t, 4>(4)) |
|
||||
((min_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
|
||||
scale_f = convert<float>(sc) * dall;
|
||||
min_f = convert<float>(m) * (-dmin);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Q3_K, SOA reorder layout produced by reorder_qw_q3_k:
|
||||
// [qs: nb*(QK_K/4)] [hmask: nb*(QK_K/8)] [scales: nb*12] [d: nb*sizeof(half)]
|
||||
// with nb = nrows*num_blocks_per_row. Single super-block scale d, no dmin.
|
||||
//
|
||||
// 3 bits per weight: 2 low bits in qs, 1 high bit in hmask. The 8 output chunks
|
||||
// of 32 (matching dequantize_row_q3_K) map to super-chunk s (0..7): byte base
|
||||
// 32*(s/4) into the 64-byte qs array, bit shift 2*(s%4); the low 16 lanes use
|
||||
// scale code 2s, the high 16 use 2s+1. hmask is a 32-byte array (like Q5_K's
|
||||
// qh) where chunk s uses bit s of the same 32 bytes, but INVERTED: the value is
|
||||
// (q & 3) - (hmask_bit_set ? 0 : 4), i.e. (q & 3) + 4*bit - 4.
|
||||
//
|
||||
// The 16 6-bit scale codes are packed into 12 bytes (get_scale_min layout for
|
||||
// Q3_K): low nibbles from bytes 0..7, high 2 bits from bytes 8..11 shifted by
|
||||
// 0/2/4/6; the dequant scale is d * (code - 32).
|
||||
// ---------------------------------------------------------------------------
|
||||
template <> struct esimd_reorder_q_traits<GGML_TYPE_Q3_K> {
|
||||
struct ptrs {
|
||||
const uint8_t * qs;
|
||||
const uint8_t * hmask;
|
||||
const uint8_t * scales;
|
||||
const sycl::half * d;
|
||||
};
|
||||
|
||||
static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
|
||||
const uint8_t * qs = (const uint8_t *) vx;
|
||||
const uint8_t * hmask = qs + nb * (QK_K / 4);
|
||||
const uint8_t * scales = hmask + nb * (QK_K / 8);
|
||||
const sycl::half * d = (const sycl::half *) (scales + nb * 12);
|
||||
return { qs, hmask, scales, d };
|
||||
}
|
||||
|
||||
// unpack the 12 packed bytes into 16 6-bit scale codes (dequantize_row_q3_K
|
||||
// aux layout), returned as float scale = d * (code - 32).
|
||||
// done with wide (8/16-lane) ops rather than four 4-lane groups.
|
||||
static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 16> unpack_scales(
|
||||
sycl::ext::intel::esimd::simd<uint8_t, 12> in, float d) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
|
||||
// low 6-bit part: codes 0..7 = low nibble of bytes 0..7,
|
||||
// codes 8..15 = high nibble of bytes 0..7
|
||||
simd<uint8_t, 8> lo8 = in.select<8, 1>(0);
|
||||
simd<uint8_t, 16> code;
|
||||
code.select<8, 1>(0) = lo8 & simd<uint8_t, 8>(0x0F);
|
||||
code.select<8, 1>(8) = lo8 >> simd<uint8_t, 8>(4);
|
||||
|
||||
// high 2-bit part: bytes 8..11 replicated 4x, group g (0..3) shifted 2*g
|
||||
simd<uint8_t, 16> hib;
|
||||
hib.select<4, 1>(0) = in.select<4, 1>(8);
|
||||
hib.select<4, 1>(4) = in.select<4, 1>(8);
|
||||
hib.select<4, 1>(8) = in.select<4, 1>(8);
|
||||
hib.select<4, 1>(12) = in.select<4, 1>(8);
|
||||
simd<uint8_t, 16> hshift;
|
||||
hshift.select<4, 1>(0) = 0;
|
||||
hshift.select<4, 1>(4) = 2;
|
||||
hshift.select<4, 1>(8) = 4;
|
||||
hshift.select<4, 1>(12) = 6;
|
||||
hib = (hib >> hshift) & simd<uint8_t, 16>(0x03);
|
||||
|
||||
code = code | (hib << simd<uint8_t, 16>(4));
|
||||
return (convert<float>(code) - 32.0f) * d;
|
||||
}
|
||||
|
||||
static ESIMD_INLINE void mac_pair(
|
||||
const ptrs & pa, size_t bia,
|
||||
const ptrs & pb, size_t bib, bool has_b,
|
||||
sycl::ext::intel::esimd::simd<float, 256> & y_vec,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_a,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
|
||||
simd<uint8_t, 64> qs_a = block_load<uint8_t, 64>(pa.qs + bia * (QK_K / 4));
|
||||
simd<uint8_t, 64> qs_b = 0;
|
||||
simd<uint8_t, 32> hmask_a = block_load<uint8_t, 32>(pa.hmask + bia * (QK_K / 8));
|
||||
simd<uint8_t, 32> hmask_b = 0;
|
||||
simd<uint8_t, 12> scales_a = block_load<uint8_t, 12>(pa.scales + bia * 12);
|
||||
simd<uint8_t, 12> scales_b = 0;
|
||||
|
||||
const float d_a = (float) pa.d[bia];
|
||||
float d_b = 0.0f;
|
||||
if (has_b) {
|
||||
qs_b = block_load<uint8_t, 64>(pb.qs + bib * (QK_K / 4));
|
||||
hmask_b = block_load<uint8_t, 32>(pb.hmask + bib * (QK_K / 8));
|
||||
scales_b = block_load<uint8_t, 12>(pb.scales + bib * 12);
|
||||
d_b = (float) pb.d[bib];
|
||||
}
|
||||
|
||||
simd<float, 16> scale_f_a = unpack_scales(scales_a, d_a);
|
||||
simd<float, 16> scale_f_b = unpack_scales(scales_b, d_b);
|
||||
|
||||
#pragma unroll
|
||||
for (int s = 0; s < 8; ++s) {
|
||||
const int byte_base = 32 * (s / 4);
|
||||
const uint8_t shift = (uint8_t) (2 * (s % 4));
|
||||
simd<float, 32> y_s = y_vec.select<32, 1>(s * 32);
|
||||
|
||||
// 2 low bits from qs, high bit from hmask (bit s of the same 32 bytes);
|
||||
// value = (q & 3) + 4*bit - 4 (inverted hmask: subtract 4 when bit clear).
|
||||
// merge in the integer domain: q3 = (q & 3) | (bit << 2) in {0..7},
|
||||
// then a single convert + subtract yields q3 - 4 (one convert, not two)
|
||||
simd<uint16_t, 32> q3_a = convert<uint16_t>(
|
||||
(qs_a.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
|
||||
q3_a |= convert<uint16_t>(
|
||||
((hmask_a >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
|
||||
simd<uint16_t, 32> q3_b = convert<uint16_t>(
|
||||
(qs_b.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
|
||||
q3_b |= convert<uint16_t>(
|
||||
((hmask_b >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
|
||||
|
||||
simd<float, 32> qf_a = convert<float>(q3_a) - 4.0f;
|
||||
simd<float, 32> qf_b = convert<float>(q3_b) - 4.0f;
|
||||
|
||||
const float scale_a_lo = scale_f_a[2 * s + 0];
|
||||
const float scale_a_hi = scale_f_a[2 * s + 1];
|
||||
const float scale_b_lo = scale_f_b[2 * s + 0];
|
||||
const float scale_b_hi = scale_f_b[2 * s + 1];
|
||||
|
||||
simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
|
||||
simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
|
||||
|
||||
simd<float, 32> deq_a = qf_a * scale_vec_a;
|
||||
simd<float, 32> deq_b = qf_b * scale_vec_b;
|
||||
|
||||
acc_a += y_s * deq_a;
|
||||
acc_b += y_s * deq_b;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Q4_K, SOA reorder layout produced by reorder_qw_q4_k:
|
||||
// [qs: nb*(QK_K/2)] [scales: nb*K_SCALE_SIZE] [dm: nb*sizeof(half2)]
|
||||
// with nb = nrows*num_blocks_per_row.
|
||||
// ---------------------------------------------------------------------------
|
||||
template <> struct esimd_reorder_q_traits<GGML_TYPE_Q4_K> {
|
||||
struct ptrs {
|
||||
const uint8_t * qs;
|
||||
const uint8_t * scales;
|
||||
const sycl::half * dm;
|
||||
};
|
||||
|
||||
static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
|
||||
const uint8_t * qs = (const uint8_t *) vx;
|
||||
const uint8_t * scales = qs + nb * (QK_K / 2);
|
||||
const sycl::half * dm = (const sycl::half *) (scales + nb * K_SCALE_SIZE);
|
||||
return { qs, scales, dm };
|
||||
}
|
||||
|
||||
static ESIMD_INLINE void mac_pair(
|
||||
const ptrs & pa, size_t bia,
|
||||
const ptrs & pb, size_t bib, bool has_b,
|
||||
sycl::ext::intel::esimd::simd<float, 256> & y_vec,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_a,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
|
||||
simd<uint8_t, 128> qs_a = block_load<uint8_t, 128>(pa.qs + bia * (QK_K / 2));
|
||||
simd<uint8_t, 128> qs_b = 0;
|
||||
simd<uint8_t, 12> scales_a = block_load<uint8_t, 12>(pa.scales + bia * K_SCALE_SIZE);
|
||||
simd<uint8_t, 12> scales_b = 0;
|
||||
|
||||
const float dall_a = (float) pa.dm[bia * 2 + 0];
|
||||
const float dmin_a = (float) pa.dm[bia * 2 + 1];
|
||||
float dall_b = 0.0f;
|
||||
float dmin_b = 0.0f;
|
||||
if (has_b) {
|
||||
qs_b = block_load<uint8_t, 128>(pb.qs + bib * (QK_K / 2));
|
||||
scales_b = block_load<uint8_t, 12>(pb.scales + bib * K_SCALE_SIZE);
|
||||
dall_b = (float) pb.dm[bib * 2 + 0];
|
||||
dmin_b = (float) pb.dm[bib * 2 + 1];
|
||||
}
|
||||
|
||||
simd<float, 8> scale_f_a, min_f_a, scale_f_b, min_f_b;
|
||||
unpack_scale_min_k4(scales_a, dall_a, dmin_a, scale_f_a, min_f_a);
|
||||
unpack_scale_min_k4(scales_b, dall_b, dmin_b, scale_f_b, min_f_b);
|
||||
|
||||
simd<uint8_t, 128> qs_lo_a = qs_a & simd<uint8_t, 128>(0x0F);
|
||||
simd<uint8_t, 128> qs_hi_a = qs_a >> simd<uint8_t, 128>(4);
|
||||
simd<uint8_t, 128> qs_lo_b = qs_b & simd<uint8_t, 128>(0x0F);
|
||||
simd<uint8_t, 128> qs_hi_b = qs_b >> simd<uint8_t, 128>(4);
|
||||
|
||||
#pragma unroll
|
||||
for (int sb = 0; sb < 8; sb += 2) {
|
||||
const int q_offset = sb * 16;
|
||||
simd<float, 32> y_lo = y_vec.select<32, 1>(sb * 32);
|
||||
simd<float, 32> y_hi = y_vec.select<32, 1>((sb + 1) * 32);
|
||||
|
||||
const float scale_a_lo = scale_f_a[sb];
|
||||
const float scale_a_hi = scale_f_a[sb + 1];
|
||||
const float min_a_lo = min_f_a[sb];
|
||||
const float min_a_hi = min_f_a[sb + 1];
|
||||
const float scale_b_lo = scale_f_b[sb];
|
||||
const float scale_b_hi = scale_f_b[sb + 1];
|
||||
const float min_b_lo = min_f_b[sb];
|
||||
const float min_b_hi = min_f_b[sb + 1];
|
||||
|
||||
simd<uint8_t, 32> qa_lo = qs_lo_a.select<32, 1>(q_offset);
|
||||
simd<uint8_t, 32> qa_hi = qs_hi_a.select<32, 1>(q_offset);
|
||||
simd<uint8_t, 32> qb_lo = qs_lo_b.select<32, 1>(q_offset);
|
||||
simd<uint8_t, 32> qb_hi = qs_hi_b.select<32, 1>(q_offset);
|
||||
|
||||
simd<float, 32> deq_a_lo = convert<float>(qa_lo) * scale_a_lo + min_a_lo;
|
||||
simd<float, 32> deq_a_hi = convert<float>(qa_hi) * scale_a_hi + min_a_hi;
|
||||
simd<float, 32> deq_b_lo = convert<float>(qb_lo) * scale_b_lo + min_b_lo;
|
||||
simd<float, 32> deq_b_hi = convert<float>(qb_hi) * scale_b_hi + min_b_hi;
|
||||
|
||||
acc_a += y_lo * deq_a_lo;
|
||||
acc_b += y_lo * deq_b_lo;
|
||||
acc_a += y_hi * deq_a_hi;
|
||||
acc_b += y_hi * deq_b_hi;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Q6_K, SOA reorder layout:
|
||||
// [ql: nb*(QK_K/2)] [qh: nb*(QK_K/4)] [scales(int8): nb*(QK_K/16)] [d: nb*half]
|
||||
// ---------------------------------------------------------------------------
|
||||
template <> struct esimd_reorder_q_traits<GGML_TYPE_Q6_K> {
|
||||
struct ptrs {
|
||||
const uint8_t * ql;
|
||||
const uint8_t * qh;
|
||||
const int8_t * scales;
|
||||
const sycl::half * d;
|
||||
};
|
||||
|
||||
static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
|
||||
const uint8_t * ql = (const uint8_t *) vx;
|
||||
const uint8_t * qh = ql + nb * (QK_K / 2);
|
||||
const int8_t * scales = (const int8_t *) (qh + nb * (QK_K / 4));
|
||||
const sycl::half * d = (const sycl::half *) (scales + nb * (QK_K / 16));
|
||||
return { ql, qh, scales, d };
|
||||
}
|
||||
|
||||
static ESIMD_INLINE void mac_pair(
|
||||
const ptrs & pa, size_t bia,
|
||||
const ptrs & pb, size_t bib, bool has_b,
|
||||
sycl::ext::intel::esimd::simd<float, 256> & y_vec,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_a,
|
||||
sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
|
||||
using namespace sycl::ext::intel::esimd;
|
||||
|
||||
simd<uint8_t, 128> ql_a = block_load<uint8_t, 128>(pa.ql + bia * (QK_K / 2));
|
||||
simd<uint8_t, 128> ql_b = 0;
|
||||
simd<uint8_t, 64> qh_a = block_load<uint8_t, 64>(pa.qh + bia * (QK_K / 4));
|
||||
simd<uint8_t, 64> qh_b = 0;
|
||||
simd<int8_t, 16> scales_a = block_load<int8_t, 16>(pa.scales + bia * (QK_K / 16));
|
||||
simd<int8_t, 16> scales_b = 0;
|
||||
|
||||
const float d_a = (float) pa.d[bia];
|
||||
float d_b = 0.0f;
|
||||
if (has_b) {
|
||||
ql_b = block_load<uint8_t, 128>(pb.ql + bib * (QK_K / 2));
|
||||
qh_b = block_load<uint8_t, 64>(pb.qh + bib * (QK_K / 4));
|
||||
scales_b = block_load<int8_t, 16>(pb.scales + bib * (QK_K / 16));
|
||||
d_b = (float) pb.d[bib];
|
||||
}
|
||||
|
||||
simd<float, 16> sc_a = convert<float>(scales_a);
|
||||
simd<float, 16> sc_b = convert<float>(scales_b);
|
||||
|
||||
#pragma unroll
|
||||
for (int im = 0; im < 2; ++im) {
|
||||
simd<uint8_t, 32> ql_lo_a = ql_a.select<32, 1>(64 * im);
|
||||
simd<uint8_t, 32> ql_hi_a = ql_a.select<32, 1>(64 * im + 32);
|
||||
simd<uint8_t, 32> qh_bits_a = qh_a.select<32, 1>(32 * im);
|
||||
simd<uint8_t, 32> ql_lo_b = ql_b.select<32, 1>(64 * im);
|
||||
simd<uint8_t, 32> ql_hi_b = ql_b.select<32, 1>(64 * im + 32);
|
||||
simd<uint8_t, 32> qh_bits_b = qh_b.select<32, 1>(32 * im);
|
||||
|
||||
// reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K)
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 4; ++g) {
|
||||
simd<float, 32> y_g = y_vec.select<32, 1>(32 * (4 * im + g));
|
||||
|
||||
const float scale_a_lo = sc_a[8 * im + 2 * g + 0] * d_a;
|
||||
const float scale_a_hi = sc_a[8 * im + 2 * g + 1] * d_a;
|
||||
const float scale_b_lo = sc_b[8 * im + 2 * g + 0] * d_b;
|
||||
const float scale_b_hi = sc_b[8 * im + 2 * g + 1] * d_b;
|
||||
|
||||
simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
|
||||
simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
|
||||
|
||||
simd<uint8_t, 32> qa;
|
||||
simd<uint8_t, 32> qb;
|
||||
switch (g) {
|
||||
case 0:
|
||||
qa = (ql_lo_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
|
||||
qb = (ql_lo_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
|
||||
break;
|
||||
case 1:
|
||||
qa = (ql_hi_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
|
||||
qb = (ql_hi_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
|
||||
break;
|
||||
case 2:
|
||||
qa = (ql_lo_a >> simd<uint8_t, 32>(4)) | (qh_bits_a & simd<uint8_t, 32>(0x30));
|
||||
qb = (ql_lo_b >> simd<uint8_t, 32>(4)) | (qh_bits_b & simd<uint8_t, 32>(0x30));
|
||||
break;
|
||||
default:
|
||||
qa = (ql_hi_a >> simd<uint8_t, 32>(4)) | ((qh_bits_a & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
|
||||
qb = (ql_hi_b >> simd<uint8_t, 32>(4)) | ((qh_bits_b & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
|
||||
break;
|
||||
}
|
||||
|
||||
simd<float, 32> deq_a = (convert<float>(qa) - 32.0f) * scale_vec_a;
|
||||
simd<float, 32> deq_b = (convert<float>(qb) - 32.0f) * scale_vec_b;
|
||||
|
||||
acc_a += y_g * deq_a;
|
||||
acc_b += y_g * deq_b;
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace ggml_sycl_esimd
|
||||
|
||||
#endif // GGML_SYCL_ESIMD_HPP
|
||||
|
|
@ -43,6 +43,9 @@
|
|||
# include <sycl/ext/oneapi/virtual_mem/virtual_mem.hpp>
|
||||
# define GGML_SYCL_SUPPORT_VMM
|
||||
#endif
|
||||
#if defined(__INTEL_LLVM_COMPILER)
|
||||
#define GGML_SYCL_DMMV_HAS_ESIMD
|
||||
#endif
|
||||
#include <sycl/half_type.hpp>
|
||||
|
||||
#include "ggml.h"
|
||||
|
|
@ -90,6 +93,7 @@ int g_ggml_sycl_fa_onednn = 1;
|
|||
int g_ggml_sycl_fa_onednn_max_kv = 0;
|
||||
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_prioritize_dmmv = 0;
|
||||
int g_ggml_sycl_use_async_mem_op = 0;
|
||||
int g_ggml_sycl_use_async_mem_op_requested = 1;
|
||||
|
|
@ -298,6 +302,7 @@ static void ggml_check_sycl() try {
|
|||
g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0);
|
||||
g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
|
||||
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_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
|
||||
|
||||
g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL);
|
||||
|
|
@ -392,6 +397,12 @@ static void ggml_check_sycl() try {
|
|||
|
||||
GGML_LOG_INFO(" GGML_SYCL_ENABLE_FUSION: %d\n", g_ggml_sycl_enable_fusion);
|
||||
|
||||
#if defined(__INTEL_LLVM_COMPILER)
|
||||
GGML_LOG_INFO(" GGML_SYCL_ENABLE_ESIMD: %d\n", g_ggml_sycl_enable_esimd);
|
||||
#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_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);
|
||||
|
|
@ -3740,6 +3751,22 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) {
|
|||
}
|
||||
}
|
||||
|
||||
static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) {
|
||||
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
|
||||
switch (type) {
|
||||
case GGML_TYPE_Q3_K:
|
||||
case GGML_TYPE_Q4_K:
|
||||
case GGML_TYPE_Q6_K:
|
||||
return true;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
#else
|
||||
GGML_UNUSED(type);
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
static bool ggml_sycl_supports_dmmv(enum ggml_type type) {
|
||||
switch (type) {
|
||||
case GGML_TYPE_Q1_0:
|
||||
|
|
@ -4443,19 +4470,22 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
|
|||
use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
|
||||
#endif // SYCL_USE_XMX
|
||||
|
||||
// Dispatch becomes obscure with the reorder, MMVQ when the reorder optimization
|
||||
// is enabled takes precedence over DMMV, the current if-else implementation
|
||||
// requires disabling DMMV if both conditions are met
|
||||
// When reorder is enabled, both ESIMD, MMVQ and DMMV kernels may be used. For
|
||||
// best performance use ESIMD when supported, followed by MMVQ, and finally DMMV.
|
||||
// But the reordered ESIMD path cannot be used without reordered MMVQ. A later
|
||||
// multi-token call (ne[1] in 2..8) will take the MMVQ path and it would read the
|
||||
// reordered bytes as if they were still the unreordered layout.
|
||||
|
||||
if (!g_ggml_sycl_prioritize_dmmv && ((should_reorder_tensor(ctx, dst) &&
|
||||
ggml_sycl_supports_reorder_mmvq(src0->type)))) {
|
||||
// Arc770 get benefit with Q4_0 by skipping it.
|
||||
if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
|
||||
gpu_arch::intel_gpu_acm_g10 &&
|
||||
src0->type == GGML_TYPE_Q4_0)) {
|
||||
use_dequantize_mul_mat_vec =
|
||||
use_dequantize_mul_mat_vec && !use_mul_mat_vec_q;
|
||||
}
|
||||
bool use = g_ggml_sycl_enable_esimd && ggml_sycl_supports_reorder_esimd(src0->type);
|
||||
// Arc770 get benefit with Q4_0 by skipping MMVQ path
|
||||
if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
|
||||
gpu_arch::intel_gpu_acm_g10 &&
|
||||
src0->type == GGML_TYPE_Q4_0)) {
|
||||
use = use || !use_mul_mat_vec_q;
|
||||
}
|
||||
use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && use;
|
||||
}
|
||||
|
||||
if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {
|
||||
|
|
|
|||
Loading…
Reference in New Issue