sycl : fix get_rows on reordered quantized weights (#30430)

Co-authored-by: Sam <sam@garagellm.tail7dd858.ts.net>
2026-10-20 07:47:21 +08:00
1 parent f711c8f303
commit 3bb63cbcbf
@@ -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>
@@ -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);
@@ -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;
@@ -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}) {