public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-13 01:35:40 +02:00
Code Issues Packages Projects Releases Wiki Activity

SYCL: Q8_0 DMMV ESIMD and MMVQ wide load (llama/29186)

Browse Source 
* Adding wide-load mmvq for Q8_0 and esimd dmmv for q8_0

Assisted-by: Codex

* remove guard for q8_0

* remove docs

* Simplify by committing to clean code without fallback

* Add feature flag as requested

Assisted-by: Claude Opus 5

---------

Co-authored-by: cwriter <cwriter@localhost>
This commit is contained in:
cwriter
2026-10-06 10:38:17 +03:00
committed by Georgi Gerganov
co-authored by cwriter
parent 95a05ed4da
commit cd1bee5203
5 changed files with 192 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/common.hpp
+1
View File
@@ -64,6 +64,7 @@ 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_mmvq_wide;
extern int g_ggml_sycl_prioritize_dmmv;
extern int g_ggml_sycl_enable_flash_attention;
extern int g_ggml_sycl_dev2dev_memcpy;
ggml/src/ggml-sycl/dmmv.cpp
+147 -2
View File
@@ -2006,6 +2006,135 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const
});
}
// Q8_0 SOA reorder layout: [qs: nb*QK8_0] [d: nb*sizeof(half)].
// Process eight blocks per stripe and process remaining blocks one at a time.
template <int NBLK>
ESIMD_INLINE void q8_0_mac_stripe(
const int8_t * qs_a, const int8_t * qs_b,
const sycl::half * d_a, const sycl::half * d_b, bool has_b,
sycl::ext::intel::esimd::simd<float, 32 * NBLK> & 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<int8_t, 32 * NBLK> qa = block_load<int8_t, 32 * NBLK>(qs_a);
simd<int8_t, 32 * NBLK> qb = 0;
// Scale rows can be only 2-byte aligned when nblk_row is odd.
simd<sycl::half, NBLK> da = block_load<sycl::half, NBLK>(d_a, element_aligned_tag{});
simd<sycl::half, NBLK> db = 0;
if (has_b) {
qb = block_load<int8_t, 32 * NBLK>(qs_b);
db = block_load<sycl::half, NBLK>(d_b, element_aligned_tag{});
}
simd<float, NBLK> da_f = convert<float>(da);
simd<float, NBLK> db_f = convert<float>(db);
#pragma unroll
for (int s = 0; s < NBLK; ++s) {
simd<float, 32> y_s = y_vec.template select<32, 1>(s * 32);
simd<int8_t, 32> qa_s = qa.template select<32, 1>(s * 32);
simd<int8_t, 32> qb_s = qb.template select<32, 1>(s * 32);
const float sa = da_f[s];
const float sb = db_f[s];
acc_a += y_s * (convert<float>(qa_s) * sa);
acc_b += y_s * (convert<float>(qb_s) * sb);
}
}
template <int WG>
ESIMD_INLINE void dequantize_mul_mat_vec_q8_0_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;
constexpr int STRIPE = 8;
const int nblk_row = ncols / QK8_0;
const size_t nb = (size_t) nrows * nblk_row;
const int8_t * qs = (const int8_t *) vx;
const sycl::half * d = (const sycl::half *) (qs + nb * QK8_0);
const int tid = it.get_local_id(0);
const int row_pair = it.get_group(0);
const int row0 = row_pair * 2;
const bool has_row1 = row0 + 1 < nrows;
const size_t base0 = (size_t) row0 * nblk_row;
const size_t base1 = has_row1 ? (size_t) (row0 + 1) * nblk_row : base0;
simd<float, 32> acc0 = 0.0f;
simd<float, 32> acc1 = 0.0f;
// Each thread processes one contiguous stripe.
int ib = 0;
for (; ib + WG * STRIPE <= nblk_row; ib += WG * STRIPE) {
const int b = ib + tid * STRIPE;
simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) b * QK8_0);
q8_0_mac_stripe<STRIPE>(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0,
d + base0 + b, d + base1 + b, has_row1, y_vec, acc0, acc1);
}
// Distribute remaining blocks across the work-group.
for (int b = ib + tid; b < nblk_row; b += WG) {
simd<float, 32> y_vec = block_load<float, 32>(y + (size_t) b * QK8_0);
q8_0_mac_stripe<1>(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0,
d + base0 + b, d + base1 + b, 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 < WG; ++p) {
sum0 += lmem[p * 2 + 0];
sum1 += lmem[p * 2 + 1];
}
dst[row0 + 0] = sum0;
if (has_row1) {
dst[row0 + 1] = sum1;
}
}
}
template <int WG>
static void q8_0_esimd_launch(const void * vx, const float * y, float * dst, const int ncols,
const int nrows, dpct::queue_ptr stream) {
const int workgroups = (nrows + 1) / 2;
stream->submit([&](sycl::handler & h) {
sycl::local_accessor<float, 1> lmem(sycl::range<1>(WG * 2), h);
h.parallel_for(
sycl::nd_range<1>(sycl::range<1>((size_t) workgroups * WG), sycl::range<1>(WG)),
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
dequantize_mul_mat_vec_q8_0_reorder_esimd<WG>(vx, y, dst, ncols, nrows, lmem, it);
});
});
}
static void dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(const void *vx, const float *y,
float *dst, const int ncols,
const int nrows,
dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK8_0 == 0);
// Scale the work-group with the number of blocks per row.
const int nblk_row = ncols / QK8_0;
if (nblk_row >= 64) {
q8_0_esimd_launch<8>(vx, y, dst, ncols, nrows, stream);
} else if (nblk_row >= 32) {
q8_0_esimd_launch<4>(vx, y, dst, ncols, nrows, stream);
} else if (nblk_row >= 16) {
q8_0_esimd_launch<2>(vx, y, dst, ncols, nrows, stream);
} else {
q8_0_esimd_launch<1>(vx, y, dst, ncols, nrows, stream);
}
}
#endif // GGML_SYCL_DMMV_HAS_ESIMD
static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
@@ -2072,11 +2201,20 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
ggml_sycl_pool_alloc<sycl::half> src1_dfloat_a(ctx.pool());
sycl::half *src1_dfloat = nullptr; // dfloat == half
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
// The ESIMD Q8_0 kernel reads F32 activations.
const bool q8_0_esimd = src0->type == GGML_TYPE_Q8_0 && g_ggml_sycl_enable_esimd &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra) &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder;
#else
const bool q8_0_esimd = false;
#endif
bool src1_convert_f16 =
src0->type == GGML_TYPE_Q1_0 ||
src0->type == GGML_TYPE_Q4_0 || src0->type == GGML_TYPE_Q4_1 ||
src0->type == GGML_TYPE_Q5_0 || src0->type == GGML_TYPE_Q5_1 ||
src0->type == GGML_TYPE_Q8_0 || src0->type == GGML_TYPE_F16 ||
(src0->type == GGML_TYPE_Q8_0 && !q8_0_esimd) || src0->type == GGML_TYPE_F16 ||
src0->type == GGML_TYPE_BF16;
if (src1_convert_f16) {
@@ -2120,7 +2258,14 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
case GGML_TYPE_Q8_0:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
if (g_ggml_sycl_enable_esimd) {
dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
} else
#endif
{
dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
} else {
dequantize_mul_mat_vec_q8_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
ggml/src/ggml-sycl/ggml-sycl.cpp
+4 -1
View File
@@ -103,6 +103,7 @@ int g_ggml_sycl_memtrace_step = 64;
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_prioritize_dmmv = 0;
int g_ggml_sycl_use_async_mem_op = 0;
int g_ggml_sycl_use_async_mem_op_requested = 1;
@@ -398,6 +399,7 @@ static void ggml_check_sycl() try {
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_mmvq_wide = ggml_sycl_get_env("GGML_SYCL_MMVQ_WIDE", 1);
g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
@@ -521,7 +523,7 @@ 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_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);
@@ -4142,6 +4144,7 @@ static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) {
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;
ggml/src/ggml-sycl/mmvq.cpp
+10
View File
@@ -1192,6 +1192,16 @@ static void reorder_mul_mat_vec_q8_0_q8_1_sycl(const void * vx, const void * vy,
const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
if (g_ggml_sycl_mmvq_wide) {
stream->submit([&](sycl::handler & cgh) {
cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
mul_mat_vec_q_reorder<reorder_vec_dot_q8_0_wide>(vx, vy, dst, ncols, nrows, nd_item);
});
});
return;
}
stream->submit([&](sycl::handler & cgh) {
cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
ggml/src/ggml-sycl/vecdotq.hpp
+30
View File
@@ -415,6 +415,36 @@ template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0> {
};
};
// Load four contiguous dwords per operand instead of loading each value separately.
struct reorder_vec_dot_q8_0_wide {
static constexpr ggml_type gtype = GGML_TYPE_Q8_0;
using q8_0_block = ggml_sycl_reordered::block_q_t<GGML_TYPE_Q8_0>;
using q8_0_traits = typename q8_0_block::traits;
__dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair<int, int> ibx_offset,
const std::pair<int, int> d_offset, const int8_t * q8_1_quant_ptr,
const sycl::half2 * q8_1_ds, const int & iqs) {
static_assert(q8_0_traits::vdr_mmvq == 4, "the wide load moves exactly four dwords");
const uint8_t * base = static_cast<const uint8_t *>(vbq);
const int8_t * qs = reinterpret_cast<const int8_t *>(base + ibx_offset.first);
const ggml_half d = *reinterpret_cast<const ggml_half *>(base + d_offset.first);
const sycl::int4 v = *reinterpret_cast<const sycl::int4 *>(qs + sizeof(int) * iqs);
const sycl::int4 u = *reinterpret_cast<const sycl::int4 *>(q8_1_quant_ptr + sizeof(int) * iqs);
int sumi = 0;
#pragma unroll
for (int i = 0; i < 4; ++i) {
sumi = dpct::dp4a(v[i], u[i], sumi);
}
const sycl::half2 ds_values = *q8_1_ds;
return static_cast<float>(d) * static_cast<float>(ds_values[0]) * sumi;
}
};
template <> struct reorder_vec_dot_q_sycl<GGML_TYPE_Q8_0> {
static constexpr ggml_type gtype = GGML_TYPE_Q8_0;