diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 3b807c7cb..b97c9e185 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3111,8 +3111,25 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten if (src1_on_device && src1_is_contiguous) { scope_op_debug_print scope_dbg_print(__func__, "/quantize_row_q8_1_sycl", dst, /*num_src=*/2, " : converting src1 to Q8_1"); + fprintf(stderr, "[DBG quantize] line 3103: quantize_f called, ne10=%lld nrows1=%lld padded=%lld\n", + (long long)ne10, (long long)nrows1, (long long)src1_padded_col_size); try { quantize_row_q8_1_sycl(dev[i].src1_ddf, dev[i].src1_ddq, ne10, nrows1, src1_padded_col_size, stream); + // Dump first 16 bytes of quantized output + { + char dbg_buf[16]; + SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dbg_buf, dev[i].src1_ddq, 16).wait())); + fprintf(stderr, "[DBG quantize] first 16 bytes: "); + for (int dbg = 0; dbg < 16; dbg++) fprintf(stderr, "%02x ", (unsigned char)dbg_buf[dbg]); + fprintf(stderr, "\n"); + // Also dump bytes at offset ncols (ds section) + int dbg_ds_off = ne10; + char dbg_ds[8]; + SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dbg_ds, (char*)dev[i].src1_ddq + dbg_ds_off, 8).wait())); + fprintf(stderr, "[DBG quantize] ds at offset %d: ", dbg_ds_off); + for (int dbg = 0; dbg < 8; dbg++) fprintf(stderr, "%02x ", (unsigned char)dbg_ds[dbg]); + fprintf(stderr, "\n"); + } } catch (sycl::exception const &exc) { std::cerr << "Quantize_row_q8_1_sycl error" << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl; @@ -3206,11 +3223,17 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten } if constexpr (quantize_enabled) { - scope_op_debug_print scope_dbg_print(__func__, "/quantize_row_q8_1_sycl", dst, - /*num_src=*/2, " : converting src1 to Q8_1"); + fprintf(stderr, "[DBG quantize] quantize_f called, ne10=%lld src1_ncols=%lld padded=%lld\n", + (long long)ne10, (long long)src1_ncols, (long long)src1_padded_col_size); try { - quantize_row_q8_1_sycl(src1_ddf_i, src1_ddq_i, ne10, src1_ncols, + quantize_row_q8_1_sycl(src1_ddf_i, src1_ddq_i, ne10, src1_ncols, src1_padded_col_size, stream); + // dump first 16 bytes of quantized output + std::vector dbg_buf(16); + SYCL_CHECK(CHECK_TRY_ERROR(stream->memcpy(dbg_buf.data(), src1_ddq_i, 16).wait())); + fprintf(stderr, "[DBG quantize] first 16 bytes: "); + for (int dbg = 0; dbg < 16; dbg++) fprintf(stderr, "%02x ", (unsigned char)dbg_buf[dbg]); + fprintf(stderr, "\n"); } catch (const sycl::exception & exc) { std::cerr << "Quantize_row_q8_1_sycl error" << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl; @@ -4333,7 +4356,11 @@ static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor * } ggml_tensor_extra_gpu * extra = static_cast(src0->extra); - if (!extra || extra->optimized_feature.reorder) { + if (!extra) { + fprintf(stderr, "[DBG opt_for_reorder] src0->extra is NULL! type=%d\n", (int)src0->type); + return; + } + if (extra->optimized_feature.reorder) { return; // Skip permutations and already reordered tensors } @@ -4355,7 +4382,11 @@ static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor * break; } - if (reorder_qw(src0, ctx->stream())) { + bool ok = reorder_qw(src0, ctx->stream()); + fprintf(stderr, "[DBG opt_for_reorder] type=%d algo=%d reorder_qw=%d ncols=%lld nrows=%lld\n", + (int)src0->type, (int)mm_algorithm, (int)ok, + (long long)src0->ne[0], (long long)src0->ne[1]); + if (ok) { extra->optimized_feature.reorder = true; // Used to decode/dequan in next steps and avoid re-reordering } } @@ -4469,13 +4500,31 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor } else if (use_mul_mat_vec_q) { opt_for_reorder(&ctx, src0, src1, dst, mul_mat_algo::MMVQ); ggml_tensor_extra_gpu * extra = static_cast(src0->extra); + fprintf(stderr, "[DBG dispatch] MMVQ: type=%d extra=%p reorder=%d src1_ncols=%lld\n", + (int)src0->type, (void*)extra, extra ? (int)extra->optimized_feature.reorder : -1, + (long long)src1->ne[1]); if (extra && extra->optimized_feature.reorder) { + fprintf(stderr, "[DBG dispatch] → reorder path (quantize_and_reorder_q8_1_soa)\n"); ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_vec_q); } else { + fprintf(stderr, "[DBG dispatch] → non-reorder path (quantize_q8_1)\n"); ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_vec_q); } } else if (use_mul_mat_q) { - ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_q); + // If the weight tensor was reordered by a prior MMVQ/DMMV opt_for_reorder + // call (in-place split qs+d layout), quantize_q8_1 would read the wrong + // layout (it expects original block_q8_0 interleaved d+qs). Fall back to + // the dequantize path (no_quantize_q8_1), which dequantizes first then + // does a normal matmul — layout-agnostic. This happens when the same + // weight is used by both the encoder (mul_mat_q, many tokens) and the + // decoder (MMVQ, single token): the decoder's opt_for_reorder corrupts + // the weight for the next encoder call. + ggml_tensor_extra_gpu * mmq_extra = static_cast(src0->extra); + if (mmq_extra && mmq_extra->optimized_feature.reorder) { + ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_sycl); + } else { + ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_q); + } } else { ggml_sycl_op_mul_mat(ctx, src0, src1, dst, ggml_sycl_op_mul_mat_sycl); }