Commit 3bb63cbcb for llama.cpp
commit 3bb63cbcbfb62056c84fe9e521b0f3d756d29c32
Author: Sam Murdoch <hello@burl.au>
Date: Tue Oct 20 10:47:21 2026 +1100
sycl : fix get_rows on reordered quantized weights (#30430)
Co-authored-by: Sam <sam@garagellm.tail7dd858.ts.net>
diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp
index 8ba26d3a9..45c60e8e5 100644
--- a/ggml/src/ggml-sycl/convert.cpp
+++ b/ggml/src/ggml-sycl/convert.cpp
@@ -85,7 +85,7 @@ static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
[=](sycl::nd_item<3> item_ct1) {
- dequantize_block_q2_K_reorder(vx, y, item_ct1, nb);
+ dequantize_block_q2_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
});
}
@@ -129,7 +129,7 @@ static void dequantize_row_q3_K_sycl_reorder(const void *vx, dst_t *y, const int
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
[=](sycl::nd_item<3> item_ct1) {
- dequantize_block_q3_K_reorder(vx, y, item_ct1, nb);
+ dequantize_block_q3_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb);
});
}
@@ -231,7 +231,7 @@ static void dequantize_row_q4_K_sycl_reorder(const void * vx, dst_t * y, const i
cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(global_size), sycl::range<1>(local_size)),
[=](sycl::nd_item<1> item_ct1) {
- dequantize_block_q4_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
+ dequantize_block_q4_K_reorder(vx, y + item_ct1.get_group(0) * QK_K, item_ct1.get_group(0), get_pointer(scale_local_acc), item_ct1, nb);
});
});
}
@@ -258,7 +258,7 @@ static void dequantize_row_q5_K_sycl_reorder(const void * vx, dst_t * y, const i
cgh.parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
[=](sycl::nd_item<3> item_ct1) {
- dequantize_block_q5_K_reorder(vx, y, get_pointer(scale_local_acc), item_ct1, nb);
+ dequantize_block_q5_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), get_pointer(scale_local_acc), item_ct1, nb);
});
});
}
@@ -303,7 +303,7 @@ static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const i
stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
- [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
+ [=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y + item_ct1.get_group(2) * QK_K, item_ct1.get_group(2), item_ct1, nb); });
}
template <typename dst_t>
diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp
index 48ee651ad..f00f97ad6 100644
--- a/ggml/src/ggml-sycl/dequantize.hpp
+++ b/ggml/src/ggml-sycl/dequantize.hpp
@@ -926,10 +926,9 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri
}
template<typename dst_t>
-static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
#if QK_K == 256
- const int64_t i = item_ct1.get_group(2);
if (i >= n_blocks) {
return;
}
@@ -949,7 +948,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
const int64_t is = 8 * n + l / 16;
const uint8_t q = qs[32 * n + l];
- dst_t * y = yy + i * QK_K + 128 * n;
+ dst_t * y = yy + 128 * n;
const float dall = (*dm)[0];
const float dmin = (*dm)[1];
@@ -960,6 +959,7 @@ static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t *
#else
GGML_UNUSED(vx);
GGML_UNUSED(yy);
+ GGML_UNUSED(i);
GGML_UNUSED(item_ct1);
GGML_UNUSED(n_blocks);
GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256");
@@ -1022,10 +1022,9 @@ static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restri
}
template<typename dst_t>
-static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
#if QK_K == 256
- const int64_t i = item_ct1.get_group(2);
if (i >= n_blocks) {
return;
}
@@ -1062,7 +1061,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
const float dl = d_all * (us - 32);
- dst_t * y = yy + i * QK_K + 128 * n + 32 * j;
+ dst_t * y = yy + 128 * n + 32 * j;
const uint8_t * q = qs + 32 * n;
const uint8_t * hm = hmask;
@@ -1072,6 +1071,7 @@ static void dequantize_block_q3_K_reorder(const void * __restrict__ vx, dst_t *
#else
GGML_UNUSED(vx);
GGML_UNUSED(yy);
+ GGML_UNUSED(i);
GGML_UNUSED(item_ct1);
GGML_UNUSED(n_blocks);
GGML_ABORT("Q3_K reorder dequantize not supported for QK_K != 256");
@@ -1160,14 +1160,13 @@ inline void dequantize_q4_K_common(dst_t * __restrict__ y, const uint8_t * __res
}
template <typename dst_t>
-static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, uint8_t * scales_local,
- const sycl::nd_item<1> & item_ct1, int64_t nb) {
- const int64_t i = item_ct1.get_group(0); // block index
+static void dequantize_block_q4_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t i,
+ uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) {
const int64_t tid = item_ct1.get_local_id(0); // thread index within block
const int64_t il = tid / 8;
const int64_t ir = tid % 8;
- dst_t * y = yy + i * QK_K + 64 * il + 4 * ir;
+ dst_t * y = yy + 64 * il + 4 * ir;
const uint8_t * base = static_cast<const uint8_t *>(vx);
const size_t qs_offset = i * (QK_K / 2);
@@ -1237,9 +1236,8 @@ static void dequantize_block_q5_K_wide(const void * __restrict__ vx, dst_t * __r
}
template <typename dst_t>
-static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
- const int64_t ib = item_ct1.get_group(2);
#if QK_K == 256
// assume 64 threads
@@ -1248,7 +1246,7 @@ static void dequantize_block_q5_K_reorder(const void * __restrict__ vx, dst_t *
const int64_t ir = tid % 16; // 0...15
const int64_t is = 2 * il;
- dst_t * y = yy + ib * QK_K + 64 * il + 2 * ir;
+ dst_t * y = yy + 64 * il + 2 * ir;
const uint8_t * base = static_cast<const uint8_t *>(vx);
@@ -1340,9 +1338,8 @@ static void dequantize_block_q6_K(const void * __restrict__ vx, dst_t * __restri
}
template <typename dst_t>
-static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy,
+static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, const int64_t ib,
const sycl::nd_item<3> & item_ct1, int64_t n_blocks) {
- const int64_t ib = item_ct1.get_group(2);
const int64_t tid = item_ct1.get_local_id(2);
const int64_t ip = tid / 32; // ip is 0 or 1
@@ -1359,7 +1356,7 @@ static void dequantize_block_q6_K_reorder(const void * __restrict__ vx, dst_t *
const uint8_t * scales_ptr = base_ptr + base_scales_offset;
const ggml_half * d = (const ggml_half *) (base_ptr + base_d_offset) + ib;
- dst_t * y = yy + ib * QK_K + 128 * ip + il;
+ dst_t * y = yy + 128 * ip + il;
const uint8_t * ql = ql_ptr + 64 * ip + il;
const uint8_t qh = *(qh_ptr + 32 * ip + il);
diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp
index 4e84bd222..9c17daf86 100644
--- a/ggml/src/ggml-sycl/getrows.cpp
+++ b/ggml/src/ggml-sycl/getrows.cpp
@@ -137,6 +137,47 @@ static void k_get_rows_float(
dst_row[i00] = src0_row[i00];
}
+// src0 is in reorder layout: all qs of the tensor, then all d
+template<int qk, int qr, dequantize_kernel_t_reorder dequantize_kernel, typename dst_t>
+static void k_get_rows_reorder(
+ const void * src0_qs, const void * src0_d, const int32_t * src1, dst_t * dst,
+ int64_t ne00, int64_t ne12,
+ size_t s1, size_t s2, size_t s3,
+ size_t s10, size_t s11, size_t s12,
+ const sycl::nd_item<3> &item_ct1) {
+
+ const int i00 = (item_ct1.get_group(2) * item_ct1.get_local_range(2) +
+ item_ct1.get_local_id(2)) *
+ 2;
+ const int i10 = item_ct1.get_local_range(1) * item_ct1.get_group(1) +
+ item_ct1.get_local_id(1);
+ const int i11 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
+ item_ct1.get_local_id(0)) /
+ ne12;
+ const int i12 = (item_ct1.get_group(0) * item_ct1.get_local_range(0) +
+ item_ct1.get_local_id(0)) %
+ ne12;
+
+ if (i00 >= ne00) {
+ return;
+ }
+
+ const int i01 = src1[i10*s10 + i11*s11 + i12*s12];
+
+ dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3;
+
+ const int64_t ib = (int64_t) i01*(ne00/qk) + i00/qk; // block index in the whole tensor
+ const int iqs = (i00%qk)/qr;
+ const int iybs = i00 - i00%qk;
+ const int y_offset = qr == 1 ? 1 : qk/2;
+
+ dfloat2 v;
+ dequantize_kernel(src0_d, ib, src0_qs, ib*(qk/qr) + iqs, v);
+
+ dst_row[iybs + iqs + 0] = v.x();
+ dst_row[iybs + iqs + y_offset] = v.y();
+}
+
template <int qk, int qr, dequantize_kernel_t dq>
static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
ggml_tensor *dst, const void *src0_dd,
@@ -173,6 +214,93 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr
GGML_UNUSED(ctx);
}
+template <int qk, int qr, dequantize_kernel_t_reorder dq>
+static void get_rows_sycl_reorder(const ggml_tensor *src0, const ggml_tensor *src1,
+ ggml_tensor *dst, const void *src0_dd,
+ const int32_t *src1_dd, float *dst_dd,
+ queue_ptr stream) {
+
+ GGML_TENSOR_BINARY_OP_LOCALS
+
+ // reorder is done only on 2D weights
+ GGML_ASSERT(ne02 == 1 && ne03 == 1);
+ GGML_ASSERT(ne00 % 2 == 0);
+
+ const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE);
+ const int block_num_x = (ne00 + 2*SYCL_GET_ROWS_BLOCK_SIZE - 1) / (2*SYCL_GET_ROWS_BLOCK_SIZE);
+ const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x);
+
+ const size_t s1 = nb1 / ggml_element_size(dst);
+ const size_t s2 = nb2 / ggml_element_size(dst);
+ const size_t s3 = nb3 / ggml_element_size(dst);
+
+ const size_t s10 = nb10 / ggml_element_size(src1);
+ const size_t s11 = nb11 / ggml_element_size(src1);
+ const size_t s12 = nb12 / ggml_element_size(src1);
+
+ const void * src0_d = (const char *) src0_dd + ne00*ne01/qr;
+
+ stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
+ [=](sycl::nd_item<3> item_ct1) {
+ k_get_rows_reorder<qk, qr, dq>(
+ src0_dd, src0_d, src1_dd, dst_dd, ne00, ne12, s1, s2,
+ s3, s10, s11, s12, item_ct1);
+ });
+
+ GGML_UNUSED(src0);
+}
+
+// K-quants in reorder layout: one work-group dequantizes one QK_K block of one row
+template <int dims, int nthreads, int nlocal, typename F>
+static void get_rows_sycl_reorder_k(const ggml_tensor *src0, const ggml_tensor *src1,
+ ggml_tensor *dst, const int32_t *src1_dd, float *dst_dd,
+ queue_ptr stream, F dequantize_block) {
+
+ GGML_TENSOR_BINARY_OP_LOCALS
+
+ // reorder is done only on 2D weights
+ GGML_ASSERT(ne02 == 1 && ne03 == 1);
+ GGML_ASSERT(ne00 % QK_K == 0);
+
+ const int64_t nblk = ne00 / QK_K;
+ const int64_t nb = nblk * ne01; // blocks in the whole tensor
+ const int64_t ngroups = ne10 * ne11 * ne12 * nblk;
+
+ const size_t s1 = nb1 / ggml_element_size(dst);
+ const size_t s2 = nb2 / ggml_element_size(dst);
+ const size_t s3 = nb3 / ggml_element_size(dst);
+
+ const size_t s10 = nb10 / ggml_element_size(src1);
+ const size_t s11 = nb11 / ggml_element_size(src1);
+ const size_t s12 = nb12 / ggml_element_size(src1);
+
+ stream->submit([&](sycl::handler & cgh) {
+ sycl::local_accessor<uint8_t, 1> scales_local(sycl::range<1>(nlocal), cgh);
+
+ auto kernel = [=](sycl::nd_item<dims> item_ct1) {
+ const int64_t g = item_ct1.get_group(dims - 1);
+ const int64_t b = g % nblk;
+ const int64_t r = g / nblk;
+ const int64_t i10 = r % ne10;
+ const int64_t i11 = (r / ne10) % ne11;
+ const int64_t i12 = r / (ne10 * ne11);
+
+ const int64_t i01 = src1_dd[i10*s10 + i11*s11 + i12*s12];
+
+ float * y = dst_dd + i10*s1 + i11*s2 + i12*s3 + b*QK_K;
+ dequantize_block(y, i01*nblk + b, get_pointer(scales_local), item_ct1, nb);
+ };
+
+ if constexpr (dims == 1) {
+ cgh.parallel_for(sycl::nd_range<1>(sycl::range<1>(ngroups * nthreads), sycl::range<1>(nthreads)), kernel);
+ } else {
+ cgh.parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, ngroups * nthreads), sycl::range<3>(1, 1, nthreads)), kernel);
+ }
+ });
+
+ GGML_UNUSED(src0);
+}
+
template <int qk, int qr, dequantize_kernel_f32_t dq>
static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1,
ggml_tensor *dst, const void *src0_dd,
@@ -332,6 +460,9 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
GGML_ASSERT(dst->nb[0] == ggml_type_size(dst->type));
const int32_t * src1_i32 = (const int32_t *) dst->src[1]->data;
+ // a weight shared with MUL_MAT (e.g. tied embeddings) can be in reorder layout
+ const bool src0_reordered = dst->src[0]->extra &&
+ ((const ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder;
/* TODO: Refactor and remove duplicates */
switch (dst->src[0]->type) {
case GGML_TYPE_F16:
@@ -399,14 +530,37 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q2_K:
+ if (src0_reordered) {
+ const void * vx = dst->src[0]->data;
+ get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+ [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+ dequantize_block_q2_K_reorder(vx, y, ib, item_ct1, nb);
+ GGML_UNUSED(scales_local);
+ });
+ break;
+ }
get_rows_sycl_f32<QK_K, 1, dequantize_q2_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q3_K:
+ if (src0_reordered) {
+ const void * vx = dst->src[0]->data;
+ get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+ [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+ dequantize_block_q3_K_reorder(vx, y, ib, item_ct1, nb);
+ GGML_UNUSED(scales_local);
+ });
+ break;
+ }
get_rows_sycl<QK_K, 1, dequantize_q3_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q4_0:
+ if (src0_reordered) {
+ get_rows_sycl_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(dst->src[0], dst->src[1], dst, dst->src[0]->data,
+ src1_i32, (float *)dst->data, ctx.stream());
+ break;
+ }
get_rows_sycl<QK4_0, QR4_0, dequantize_q4_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
@@ -415,6 +569,14 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q4_K:
+ if (src0_reordered) {
+ const void * vx = dst->src[0]->data;
+ get_rows_sycl_reorder_k<1, 32, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+ [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<1> & item_ct1, int64_t nb) {
+ dequantize_block_q4_K_reorder(vx, y, ib, scales_local, item_ct1, nb);
+ });
+ break;
+ }
get_rows_sycl_f32<QK_K, 1, dequantize_q4_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
@@ -427,14 +589,36 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q5_K:
+ if (src0_reordered) {
+ const void * vx = dst->src[0]->data;
+ get_rows_sycl_reorder_k<3, 64, K_SCALE_SIZE>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+ [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+ dequantize_block_q5_K_reorder(vx, y, ib, scales_local, item_ct1, nb);
+ });
+ break;
+ }
get_rows_sycl_f32<QK_K, 1, dequantize_q5_K_f32>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q6_K:
+ if (src0_reordered) {
+ const void * vx = dst->src[0]->data;
+ get_rows_sycl_reorder_k<3, 64, 1>(dst->src[0], dst->src[1], dst, src1_i32, (float *)dst->data, ctx.stream(),
+ [=](float * y, int64_t ib, uint8_t * scales_local, const sycl::nd_item<3> & item_ct1, int64_t nb) {
+ dequantize_block_q6_K_reorder(vx, y, ib, item_ct1, nb);
+ GGML_UNUSED(scales_local);
+ });
+ break;
+ }
get_rows_sycl<QK_K, 1, dequantize_q6_K>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
case GGML_TYPE_Q8_0:
+ if (src0_reordered) {
+ get_rows_sycl_reorder<QK8_0, QR8_0, dequantize_q8_0_reorder>(dst->src[0], dst->src[1], dst, dst->src[0]->data,
+ src1_i32, (float *)dst->data, ctx.stream());
+ break;
+ }
get_rows_sycl<QK8_0, QR8_0, dequantize_q8_0>(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data,
src1_i32, (float *)dst->data, ctx.stream());
break;
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 8fc90e037..1df1129cd 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -2453,6 +2453,63 @@ struct test_get_rows : public test_case {
}
};
+// GGML_OP_GET_ROWS on a weight also used by GGML_OP_MUL_MAT (tied embeddings)
+// the mul_mat runs first, so a backend that changes the weight layout must still return the right rows
+struct test_get_rows_after_mul_mat : public test_case {
+ const ggml_type type;
+ const int n; // cols
+ const int m; // rows
+ const int r; // rows to get
+
+ std::string vars() override {
+ return VARS_TO_STR4(type, n, m, r);
+ }
+
+ test_get_rows_after_mul_mat(ggml_type type = GGML_TYPE_Q8_0, int n = 256, int m = 16, int r = 4)
+ : type(type), n(n), m(m), r(r) {}
+
+ double max_nmse_err() override {
+ return 5e-4; // same as test_mul_mat
+ }
+
+ ggml_tensor * build_graph(ggml_context * ctx) override {
+ ggml_tensor * w = ggml_new_tensor_2d(ctx, type, n, m);
+ ggml_set_name(w, "w");
+
+ ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, n, 1);
+ ggml_set_name(x, "x");
+
+ ggml_tensor * rows = ggml_new_tensor_1d(ctx, GGML_TYPE_I32, r);
+ ggml_set_name(rows, "rows");
+
+ ggml_tensor * mm = ggml_mul_mat(ctx, w, x);
+ ggml_set_name(mm, "mm");
+
+ ggml_tensor * gr = ggml_get_rows(ctx, w, rows);
+ ggml_set_name(gr, "gr");
+
+ // repeat(mm) is the first source, so mul_mat is computed before get_rows
+ ggml_tensor * out = ggml_add(ctx, ggml_repeat(ctx, mm, gr), gr);
+ ggml_set_name(out, "out");
+
+ return out;
+ }
+
+ void initialize_tensors(ggml_context * ctx) override {
+ for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
+ if (t->type == GGML_TYPE_I32) {
+ std::vector<int> data(ggml_nelements(t));
+ for (size_t i = 0; i < data.size(); i++) {
+ data[i] = rand() % m;
+ }
+ ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(int));
+ } else {
+ init_tensor_uniform(t);
+ }
+ }
+ }
+};
+
// GGML_OP_GET_ROWS_BACK
struct test_get_rows_back : public test_case {
const ggml_type type;
@@ -9828,6 +9885,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
}
test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false));
+ for (ggml_type type : {GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K}) {
+ test_cases.emplace_back(new test_get_rows_after_mul_mat(type));
+ }
for (ggml_type type : all_types) {
for (int b : {1, 7}) {
for (bool v : {false, true}) {