Skip to content

Commit 547208d

Browse files
malsbatmalsbat
authored
Restore ESIMD kernels on FFN path (ggml-org#30239)
* Reenable ESIMD on the FFN path Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Extend ESIMD fusion on FFN path to include Q8_0 Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> * Remove comment explicitly referencing ESIMD in test-backend-ops Signed-off-by: Todd Malsbary <todd.malsbary@intel.com> --------- Signed-off-by: Todd Malsbary <todd.malsbary@intel.com>
1 parent 848bbda commit 547208d

6 files changed

Lines changed: 263 additions & 79 deletions

File tree

‎ggml/src/ggml-sycl/common.hpp‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,10 @@ namespace syclexp = sycl::ext::oneapi::experimental;
3737
#endif
3838
#endif
3939

40+
#if defined(__INTEL_LLVM_COMPILER)
41+
#define GGML_SYCL_DMMV_HAS_ESIMD
42+
#endif
43+
4044
#if GGML_SYCL_DNNL
4145
#include "dnnl.hpp"
4246
#include "dnnl_sycl.hpp"

‎ggml/src/ggml-sycl/dmmv.cpp‎

Lines changed: 179 additions & 71 deletions
Original file line numberDiff line numberDiff line change
@@ -10,7 +10,6 @@
1010
#endif
1111
#include <sycl/ext/intel/esimd.hpp>
1212
#include "esimd.hpp"
13-
#define GGML_SYCL_DMMV_HAS_ESIMD
1413
#endif
1514

1615
static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){
@@ -1873,11 +1872,67 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
18731872
#ifdef GGML_SYCL_DMMV_HAS_ESIMD
18741873
using ggml_sycl_esimd::GGML_SYCL_DMMV_ESIMD_WG_SIZE;
18751874

1876-
// generic reordered dequantize-matvec: each work-group owns a pair of
1877-
// consecutive output rows and updates one 32-wide accumulator per row
1878-
template <ggml_type T>
1875+
// policies for dequantize_mul_mat_vec_reorder_esimd: row mapping and what to do with the two sums
1876+
// plain mat-vec: a work-group owns two consecutive rows of one matrix
1877+
struct esimd_dmmv_row_pair_policy {
1878+
const void * vx;
1879+
float * dst;
1880+
1881+
ESIMD_INLINE const void * w_a() const { return vx; }
1882+
ESIMD_INLINE const void * w_b() const { return vx; }
1883+
1884+
ESIMD_INLINE int row_a(int g) const { return g * 2; }
1885+
ESIMD_INLINE int row_b(int g) const { return g * 2 + 1; }
1886+
ESIMD_INLINE bool has_b(int g, int nrows) const { return row_b(g) < nrows; }
1887+
1888+
ESIMD_INLINE void epilogue(int g, int nrows, float sum_a, float sum_b) const {
1889+
dst[row_a(g)] = sum_a;
1890+
if (has_b(g, nrows)) {
1891+
dst[row_b(g)] = sum_b;
1892+
}
1893+
}
1894+
};
1895+
1896+
// fused gate+up+GLU: a work-group owns one row of both matrices; a is the gate, b the up
1897+
struct esimd_dmmv_glu_policy {
1898+
const void * vgate;
1899+
const void * vup;
1900+
float * dst;
1901+
ggml_glu_op glu_op;
1902+
1903+
ESIMD_INLINE const void * w_a() const { return vgate; }
1904+
ESIMD_INLINE const void * w_b() const { return vup; }
1905+
1906+
ESIMD_INLINE int row_a(int g) const { return g; }
1907+
ESIMD_INLINE int row_b(int g) const { return g; }
1908+
ESIMD_INLINE bool has_b(int, int) const { return true; } // up always contributes
1909+
1910+
ESIMD_INLINE void epilogue(int g, int, float gate, float up) const {
1911+
using namespace sycl::ext::intel::esimd;
1912+
1913+
// SwiGLU: silu(gate)*up = gate*sigmoid(gate)*up. GEGLU's tanh form reduces
1914+
// to the same shape via tanh(y) = 2*sigmoid(2y) - 1, so only the sigmoid
1915+
// argument differs. sigmoid via exp2 (width-1 simd is a degenerate case for
1916+
// the EM pipe, so broadcast to 16 lanes); lane 0 holds the result.
1917+
const float GELU_COEF_A = 0.044715f;
1918+
const float SQRT_2_OVER_PI = 0.79788456080286535587989211986876f;
1919+
const float NEG_LOG2E = -1.44269504088896341f;
1920+
1921+
const float arg = (glu_op == GGML_GLU_OP_GEGLU) ?
1922+
2.0f * SQRT_2_OVER_PI * gate * (1.0f + GELU_COEF_A * gate * gate) :
1923+
gate;
1924+
1925+
simd<float, 16> t = arg * NEG_LOG2E;
1926+
simd<float, 16> e = exp2(t);
1927+
dst[g] = gate / (1.0f + e[0]) * up;
1928+
}
1929+
};
1930+
1931+
// generic reordered dequantize-matvec: each work-group owns a pair of rows chosen by the policy
1932+
// and updates one 32-wide accumulator per row
1933+
template <ggml_type T, typename Policy>
18791934
ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd(
1880-
const void * vx, const float * y, float * dst,
1935+
const Policy & policy, const float * y,
18811936
const int ncols, const int nrows,
18821937
sycl::local_accessor<float, 1> lmem,
18831938
const sycl::nd_item<1> & it) {
@@ -1886,41 +1941,41 @@ ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd(
18861941

18871942
const int num_blocks_per_row = ncols / QK_K;
18881943
const size_t nb = (size_t) nrows * num_blocks_per_row;
1889-
const auto ps = traits::make_ptrs(vx, nb);
1944+
const auto pa = traits::make_ptrs(policy.w_a(), nb);
1945+
const auto pb = traits::make_ptrs(policy.w_b(), nb);
18901946

1891-
const int tid = it.get_local_id(0);
1892-
const int row_pair = it.get_group(0);
1893-
const int row0 = row_pair * 2; // two consecutive output rows
1894-
const bool has_row1 = row0 + 1 < nrows;
1947+
const int tid = it.get_local_id(0);
1948+
const int g = it.get_group(0);
1949+
const int row_a = policy.row_a(g);
1950+
const int row_b = policy.row_b(g);
1951+
const bool has_b = policy.has_b(g, nrows);
18951952

1896-
// one 32-wide accumulator per output row (small footprint, no spill)
1897-
simd<float, 32> acc0 = 0.0f;
1898-
simd<float, 32> acc1 = 0.0f;
1953+
// one 32-wide accumulator per row (small footprint, no spill)
1954+
simd<float, 32> acc_a = 0.0f;
1955+
simd<float, 32> acc_b = 0.0f;
18991956

19001957
for (int ib = tid; ib < num_blocks_per_row; ib += GGML_SYCL_DMMV_ESIMD_WG_SIZE) {
19011958
simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) ib * QK_K);
19021959

1903-
const size_t bi0 = (size_t) (row0 + 0) * num_blocks_per_row + ib;
1904-
const size_t bi1 = (size_t) (row0 + 1) * num_blocks_per_row + ib;
1960+
const size_t bi_a = (size_t) row_a * num_blocks_per_row + ib;
1961+
const size_t bi_b = (size_t) row_b * num_blocks_per_row + ib;
19051962

1906-
traits::mac_pair(ps, bi0, ps, bi1, has_row1, y_vec, acc0, acc1);
1963+
traits::mac_pair(pa, bi_a, pb, bi_b, has_b, y_vec, acc_a, acc_b);
19071964
}
19081965

1909-
lmem[tid * 2 + 0] = reduce<float>(acc0, std::plus<>{});
1910-
lmem[tid * 2 + 1] = reduce<float>(acc1, std::plus<>{});
1966+
lmem[tid * 2 + 0] = reduce<float>(acc_a, std::plus<>{});
1967+
lmem[tid * 2 + 1] = reduce<float>(acc_b, std::plus<>{});
19111968
it.barrier(sycl::access::fence_space::local_space);
19121969

19131970
if (tid == 0) {
1914-
float sum0 = 0.0f;
1915-
float sum1 = 0.0f;
1971+
float sum_a = 0.0f;
1972+
float sum_b = 0.0f;
19161973
for (int p = 0; p < GGML_SYCL_DMMV_ESIMD_WG_SIZE; ++p) {
1917-
sum0 += lmem[p * 2 + 0];
1918-
sum1 += lmem[p * 2 + 1];
1919-
}
1920-
dst[row0 + 0] = sum0;
1921-
if (has_row1) {
1922-
dst[row0 + 1] = sum1;
1974+
sum_a += lmem[p * 2 + 0];
1975+
sum_b += lmem[p * 2 + 1];
19231976
}
1977+
1978+
policy.epilogue(g, nrows, sum_a, sum_b);
19241979
}
19251980
}
19261981

@@ -1936,7 +1991,7 @@ static void dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(const void *vx, const
19361991
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)),
19371992
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
19381993
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q2_K>(
1939-
vx, y, dst, ncols, nrows, lmem, it);
1994+
esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, lmem, it);
19401995
});
19411996
});
19421997
}
@@ -1953,7 +2008,7 @@ static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const
19532008
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)),
19542009
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
19552010
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q3_K>(
1956-
vx, y, dst, ncols, nrows, lmem, it);
2011+
esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, lmem, it);
19572012
});
19582013
});
19592014
}
@@ -1970,7 +2025,7 @@ static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const
19702025
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)),
19712026
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
19722027
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q4_K>(
1973-
vx, y, dst, ncols, nrows, lmem, it);
2028+
esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, lmem, it);
19742029
});
19752030
});
19762031
}
@@ -1987,7 +2042,7 @@ static void dequantize_mul_mat_vec_q5_K_sycl_reorder_esimd(const void *vx, const
19872042
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)),
19882043
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
19892044
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q5_K>(
1990-
vx, y, dst, ncols, nrows, lmem, it);
2045+
esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, lmem, it);
19912046
});
19922047
});
19932048
}
@@ -2004,7 +2059,7 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const
20042059
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)),
20052060
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
20062061
dequantize_mul_mat_vec_reorder_esimd<GGML_TYPE_Q6_K>(
2007-
vx, y, dst, ncols, nrows, lmem, it);
2062+
esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, lmem, it);
20082063
});
20092064
});
20102065
}
@@ -2045,9 +2100,9 @@ ESIMD_INLINE void q8_0_mac_stripe(
20452100
}
20462101
}
20472102

2048-
template <int WG>
2103+
template <int WG, typename Policy>
20492104
ESIMD_INLINE void dequantize_mul_mat_vec_q8_0_reorder_esimd(
2050-
const void * vx, const float * y, float * dst,
2105+
const Policy & policy, const float * y,
20512106
const int ncols, const int nrows,
20522107
sycl::local_accessor<float, 1> lmem,
20532108
const sycl::nd_item<1> & it) {
@@ -2057,84 +2112,137 @@ ESIMD_INLINE void dequantize_mul_mat_vec_q8_0_reorder_esimd(
20572112

20582113
const int nblk_row = ncols / QK8_0;
20592114
const size_t nb = (size_t) nrows * nblk_row;
2060-
const int8_t * qs = (const int8_t *) vx;
2061-
const sycl::half * d = (const sycl::half *) (qs + nb * QK8_0);
2115+
const int8_t * qs_a = (const int8_t *) policy.w_a();
2116+
const int8_t * qs_b = (const int8_t *) policy.w_b();
2117+
const sycl::half * d_a = (const sycl::half *) (qs_a + nb * QK8_0);
2118+
const sycl::half * d_b = (const sycl::half *) (qs_b + nb * QK8_0);
20622119

2063-
const int tid = it.get_local_id(0);
2064-
const int row_pair = it.get_group(0);
2065-
const int row0 = row_pair * 2;
2066-
const bool has_row1 = row0 + 1 < nrows;
2120+
const int tid = it.get_local_id(0);
2121+
const int g = it.get_group(0);
2122+
const bool has_b = policy.has_b(g, nrows);
20672123

2068-
const size_t base0 = (size_t) row0 * nblk_row;
2069-
const size_t base1 = has_row1 ? (size_t) (row0 + 1) * nblk_row : base0;
2124+
const size_t base_a = (size_t) policy.row_a(g) * nblk_row;
2125+
const size_t base_b = (size_t) policy.row_b(g) * nblk_row;
20702126

2071-
simd<float, 32> acc0 = 0.0f;
2072-
simd<float, 32> acc1 = 0.0f;
2127+
simd<float, 32> acc_a = 0.0f;
2128+
simd<float, 32> acc_b = 0.0f;
20732129

20742130
// Each thread processes one contiguous stripe.
20752131
int ib = 0;
20762132
for (; ib + WG * STRIPE <= nblk_row; ib += WG * STRIPE) {
20772133
const int b = ib + tid * STRIPE;
20782134
simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) b * QK8_0);
2079-
q8_0_mac_stripe<STRIPE>(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0,
2080-
d + base0 + b, d + base1 + b, has_row1, y_vec, acc0, acc1);
2135+
q8_0_mac_stripe<STRIPE>(qs_a + (base_a + b) * QK8_0, qs_b + (base_b + b) * QK8_0,
2136+
d_a + base_a + b, d_b + base_b + b, has_b, y_vec, acc_a, acc_b);
20812137
}
20822138

20832139
// Distribute remaining blocks across the work-group.
20842140
for (int b = ib + tid; b < nblk_row; b += WG) {
20852141
simd<float, 32> y_vec = block_load<float, 32>(y + (size_t) b * QK8_0);
2086-
q8_0_mac_stripe<1>(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0,
2087-
d + base0 + b, d + base1 + b, has_row1, y_vec, acc0, acc1);
2142+
q8_0_mac_stripe<1>(qs_a + (base_a + b) * QK8_0, qs_b + (base_b + b) * QK8_0,
2143+
d_a + base_a + b, d_b + base_b + b, has_b, y_vec, acc_a, acc_b);
20882144
}
20892145

2090-
lmem[tid * 2 + 0] = reduce<float>(acc0, std::plus<>{});
2091-
lmem[tid * 2 + 1] = reduce<float>(acc1, std::plus<>{});
2146+
lmem[tid * 2 + 0] = reduce<float>(acc_a, std::plus<>{});
2147+
lmem[tid * 2 + 1] = reduce<float>(acc_b, std::plus<>{});
20922148
it.barrier(sycl::access::fence_space::local_space);
20932149

20942150
if (tid == 0) {
2095-
float sum0 = 0.0f;
2096-
float sum1 = 0.0f;
2151+
float sum_a = 0.0f;
2152+
float sum_b = 0.0f;
20972153
for (int p = 0; p < WG; ++p) {
2098-
sum0 += lmem[p * 2 + 0];
2099-
sum1 += lmem[p * 2 + 1];
2100-
}
2101-
dst[row0 + 0] = sum0;
2102-
if (has_row1) {
2103-
dst[row0 + 1] = sum1;
2154+
sum_a += lmem[p * 2 + 0];
2155+
sum_b += lmem[p * 2 + 1];
21042156
}
2157+
2158+
policy.epilogue(g, nrows, sum_a, sum_b);
21052159
}
21062160
}
21072161

2108-
template <int WG>
2109-
static void q8_0_esimd_launch(const void * vx, const float * y, float * dst, const int ncols,
2110-
const int nrows, dpct::queue_ptr stream) {
2111-
const int workgroups = (nrows + 1) / 2;
2162+
template <int WG, typename Policy>
2163+
static void q8_0_esimd_launch(const Policy & policy, const float * y, const int ncols, const int nrows,
2164+
const int workgroups, dpct::queue_ptr stream) {
21122165
stream->submit([&](sycl::handler & h) {
21132166
sycl::local_accessor<float, 1> lmem(sycl::range<1>(WG * 2), h);
21142167
h.parallel_for(
21152168
sycl::nd_range<1>(sycl::range<1>((size_t) workgroups * WG), sycl::range<1>(WG)),
21162169
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
2117-
dequantize_mul_mat_vec_q8_0_reorder_esimd<WG>(vx, y, dst, ncols, nrows, lmem, it);
2170+
dequantize_mul_mat_vec_q8_0_reorder_esimd<WG>(policy, y, ncols, nrows, lmem, it);
21182171
});
21192172
});
21202173
}
21212174

2122-
static void dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(const void *vx, const float *y,
2123-
float *dst, const int ncols,
2124-
const int nrows,
2125-
dpct::queue_ptr stream) {
2175+
// Scale the work-group with the number of blocks per row.
2176+
template <typename Policy>
2177+
static void q8_0_esimd_launch_scaled(const Policy & policy, const float * y, const int ncols, const int nrows,
2178+
const int workgroups, dpct::queue_ptr stream) {
21262179
GGML_ASSERT(ncols % QK8_0 == 0);
21272180

2128-
// Scale the work-group with the number of blocks per row.
21292181
const int nblk_row = ncols / QK8_0;
21302182
if (nblk_row >= 64) {
2131-
q8_0_esimd_launch<8>(vx, y, dst, ncols, nrows, stream);
2183+
q8_0_esimd_launch<8>(policy, y, ncols, nrows, workgroups, stream);
21322184
} else if (nblk_row >= 32) {
2133-
q8_0_esimd_launch<4>(vx, y, dst, ncols, nrows, stream);
2185+
q8_0_esimd_launch<4>(policy, y, ncols, nrows, workgroups, stream);
21342186
} else if (nblk_row >= 16) {
2135-
q8_0_esimd_launch<2>(vx, y, dst, ncols, nrows, stream);
2187+
q8_0_esimd_launch<2>(policy, y, ncols, nrows, workgroups, stream);
21362188
} else {
2137-
q8_0_esimd_launch<1>(vx, y, dst, ncols, nrows, stream);
2189+
q8_0_esimd_launch<1>(policy, y, ncols, nrows, workgroups, stream);
2190+
}
2191+
}
2192+
2193+
static void dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(const void *vx, const float *y,
2194+
float *dst, const int ncols,
2195+
const int nrows,
2196+
dpct::queue_ptr stream) {
2197+
q8_0_esimd_launch_scaled(esimd_dmmv_row_pair_policy{ vx, dst }, y, ncols, nrows, (nrows + 1) / 2, stream);
2198+
}
2199+
2200+
// ESIMD fused gate+up+GLU on reordered K-quant SOA weights: one output row (gate
2201+
// and up) per work-group, two independent 32-wide accumulators per super-block.
2202+
// Reuses the DMMV kernels' dequant+MAC primitive (esimd_reorder_q_traits<T>::mac_pair).
2203+
template <ggml_type T>
2204+
static void dequantize_mul_mat_vec_glu_reorder_esimd(const void * vx, const void * vgate,
2205+
const float * y, float * dst,
2206+
const int ncols, const int nrows,
2207+
const ggml_glu_op glu_op,
2208+
dpct::queue_ptr stream) {
2209+
GGML_ASSERT(ncols % QK_K == 0);
2210+
const int workgroups = nrows; // one output row (gate+up) per work-group
2211+
stream->submit([&](sycl::handler & h) {
2212+
sycl::local_accessor<float, 1> lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h);
2213+
h.parallel_for(
2214+
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)),
2215+
[=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] {
2216+
dequantize_mul_mat_vec_reorder_esimd<T>(
2217+
esimd_dmmv_glu_policy{ vgate, vx, dst, glu_op }, y, ncols, nrows, lmem, it);
2218+
});
2219+
});
2220+
}
2221+
2222+
bool ggml_sycl_dequantize_mul_mat_vec_glu_reorder_esimd(enum ggml_type src0_type, enum ggml_glu_op glu_op,
2223+
const void * vx, const void * vgate, const float * y,
2224+
float * dst, int ncols, int nrows, dpct::queue_ptr stream) {
2225+
switch (src0_type) {
2226+
case GGML_TYPE_Q2_K:
2227+
dequantize_mul_mat_vec_glu_reorder_esimd<GGML_TYPE_Q2_K>(vx, vgate, y, dst, ncols, nrows, glu_op, stream);
2228+
return true;
2229+
case GGML_TYPE_Q3_K:
2230+
dequantize_mul_mat_vec_glu_reorder_esimd<GGML_TYPE_Q3_K>(vx, vgate, y, dst, ncols, nrows, glu_op, stream);
2231+
return true;
2232+
case GGML_TYPE_Q4_K:
2233+
dequantize_mul_mat_vec_glu_reorder_esimd<GGML_TYPE_Q4_K>(vx, vgate, y, dst, ncols, nrows, glu_op, stream);
2234+
return true;
2235+
case GGML_TYPE_Q5_K:
2236+
dequantize_mul_mat_vec_glu_reorder_esimd<GGML_TYPE_Q5_K>(vx, vgate, y, dst, ncols, nrows, glu_op, stream);
2237+
return true;
2238+
case GGML_TYPE_Q6_K:
2239+
dequantize_mul_mat_vec_glu_reorder_esimd<GGML_TYPE_Q6_K>(vx, vgate, y, dst, ncols, nrows, glu_op, stream);
2240+
return true;
2241+
case GGML_TYPE_Q8_0:
2242+
q8_0_esimd_launch_scaled(esimd_dmmv_glu_policy{ vgate, vx, dst, glu_op }, y, ncols, nrows, nrows, stream);
2243+
return true;
2244+
default:
2245+
return false;
21382246
}
21392247
}
21402248

0 commit comments

Comments
 (0)