From d311b882b3f8925db262d2bc8b315bc742476606 Mon Sep 17 00:00:00 2001 From: Xu Han Date: Thu, 30 Jul 2026 22:29:38 +0800 Subject: [PATCH 1/6] sycl : emit PDB for ggml-sycl.dll * add /Zi + /DEBUG so cdb resolves frames in ggml-sycl.dll --- ggml/src/ggml-sycl/CMakeLists.txt | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index a8d9c0d80..89eb350d2 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -105,6 +105,19 @@ endif() target_compile_options(ggml-sycl PRIVATE "-Wno-narrowing") +# Generate a PDB (debug symbols) for ggml-sycl.dll so cdb/WinDbg can resolve +# source lines + frames inside ggml_backend_sycl_init (otherwise the UAF/crash +# in this DLL shows only as a raw offset — no PDB, because icpx under a VS +# generator doesn't emit one by default). /Zi emits debug info during +# compile; /DEBUG at link writes the .pdb beside the .dll. Both are +# config-agnostic so Release builds also carry symbols (the PDB is read by +# cdb from the .dll's side; runtime perf is unaffected). icpx on Windows +# uses the MSVC codegen back-end, so these cdb-loadable PDBs carry line info. +if (WIN32) + target_compile_options(ggml-sycl PRIVATE "/Zi") + target_link_options(ggml-sycl PRIVATE "/DEBUG") +endif() + message(STATUS "GGML_SYCL_SUPPORT_LEVEL_ZERO_API ${GGML_SYCL_SUPPORT_LEVEL_ZERO_API}") if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API) # Link against Level Zero loader for direct device memory allocation. From db9bff463728af81616d03da77d44a70d35fb7fa Mon Sep 17 00:00:00 2001 From: Xu Han Date: Sun, 2 Aug 2026 16:34:51 +0800 Subject: [PATCH 2/6] sycl : fix quantize dispatch + reorder fallback for Q8_0 * use quantize_f instead of hardcoded quantize_q8_1 in mul_mat * restore Q8_0 in the reorder support lists * fallback mul_mat_q to dequantize when weight is reordered * add debug prints to isolate the reorder issue --- ggml/src/ggml-sycl/ggml-sycl.cpp | 61 ++++++++++++++++++++++++++++---- 1 file changed, 55 insertions(+), 6 deletions(-) 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); } From 0adcf51393e3a556607850e3cc6f58ea8158cb8f Mon Sep 17 00:00:00 2001 From: Xu Han Date: Mon, 3 Aug 2026 04:02:40 +0800 Subject: [PATCH 3/6] sycl : add reorder-aware get_rows for Q8_0/Q4_0 and k-quants * k_get_rows_reorder template for Q8_0/Q4_0 (per-element) * block kernels for Q3_K/Q4_K/Q5_K/Q6_K (32/64 threads) * dispatch reorder path when src0 is reordered --- ggml/src/ggml-sycl/getrows.cpp | 537 +++++++++++++++++++++++++++++---- 1 file changed, 473 insertions(+), 64 deletions(-) diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 2113f3563..37645f74f 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -13,9 +13,328 @@ #include "ggml-impl.h" #include "common.hpp" #include "dequantize.hpp" +#include "quants.hpp" #include "getrows.hpp" +// True if dst->src[0] has been in-place reordered by opt_for_reorder into a +// split layout. get_rows must use a reorder-aware dequantize in that case. +static bool ggml_sycl_get_rows_src0_reordered(const ggml_tensor * dst) { + ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *)dst->src[0]->extra; + return extra && extra->optimized_feature.reorder; +} + + +// Reorder-aware get_rows for src0 that has been in-place reordered by +// opt_for_reorder into a split layout (qs region, then scales/d region). +// Standard k_get_rows reads (block_q_t*)src0_row which assumes interleaved +// d+qs per block; that is wrong after reorder. This kernel indexes the split +// layout directly using the GLOBAL block index (i01*nblocks_per_row + ib) via +// block_q_t::get_block_offset / get_d_offset, then calls the per-element +// reorder dequantize. Covers Q1_0/Q4_0/Q8_0 (the types with a per-element +// dequantize_kernel_t_reorder). K-quants need block-cooperative dequantize and +// are handled separately. +template +static void k_get_rows_reorder( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne01, /*int64_t ne02, int64_t ne03,*/ + /*int64_t ne10, int64_t ne11,*/ int64_t ne12, /*int64_t ne13,*/ + /*size_t s0,*/ size_t s1, size_t s2, size_t s3, + /*size_t nb00,*/ size_t nb01, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + const sycl::nd_item<3> &item_ct1) { + + using block_type = ggml_sycl_reordered::block_q_t; + using block_traits = typename block_type::traits; + constexpr int qk = block_traits::qk; + constexpr int qr = block_traits::qr; + + 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; + + // For 3D/4D src0 (ne02/ne03 > 1) we offset by the slice; nb02/nb03 describe + // the per-slice byte offset of the START of that slice's split region. + const char * slice = (const char *)src0 + i11*nb02 + i12*nb03; + const int nblocks_per_row = ne00 / qk; + const int total_nblocks = ne01 * nblocks_per_row; + + const int ib = i00 / qk; // block index within the row + const int iqs = (i00 % qk) / qr; // quant index within block + const int gblock = i01 * nblocks_per_row + ib; // global block index + + // block_q_t gives the split-layout offsets for a given (global) block index. + // qs_ptr = this block's qs start (get_block_offset already includes gblock). + // d_ptr = the d-SECTION start (block_index 0); dq_reorder adds gblock + // internally to index the scale (see dequantize_q8_0_reorder: d = *(half*)(d_ptr+ib)). + const auto bx_off = block_type::get_block_offset(gblock, total_nblocks); + const auto d_sec = block_type::get_d_offset(ne01, ne00, 0); + const char * base = slice; + const void * qs_ptr = base + bx_off.first; + const void * d_ptr = base + d_sec.first; + + dfloat2 v; + dq_reorder(d_ptr, gblock, qs_ptr, iqs, v); + + const int iybs = i00 - i00 % qk; + const int y_offset = qr == 1 ? 1 : qk / 2; + dst_row[iybs + iqs + 0] = v.x(); + dst_row[iybs + iqs + y_offset] = v.y(); +} + + +// ── K-quant reorder get_rows kernels ── +// K-quants (Q3_K/Q4_K/Q5_K/Q6_K) use block-cooperative dequantize (scales +// shared across the 32/64/128-thread work-group). After opt_for_reorder the +// src0 buffer is split into [qs|qh/hmask|scales|d] regions indexed by GLOBAL +// block index. get_rows must read block (src1[batch]*nblocks_per_row + block_in_row) +// and write to dst_row + block_in_row*QK_K. These kernels launch one +// work-group per (batch, block_in_row); group(2)=block_in_row, group(1)=batch. + +template +static void k_get_rows_q4_K_reorder( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne01, int64_t ne12, + size_t s1, size_t s2, size_t s3, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + uint8_t * scales_local, const sycl::nd_item<3> &item_ct1) { +#if QK_K == 256 + const int64_t ib_row = item_ct1.get_group(2); // block index within row + const int64_t i10 = item_ct1.get_group(1); // batch + const int64_t tid_local = item_ct1.get_local_id(2); + const int64_t il = tid_local / 8; + const int64_t ir = tid_local % 8; + + const int64_t i11 = i10 / ne12; + const int64_t i12 = i10 % ne12; + const int64_t i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + dst_t * y = dst_row + ib_row * QK_K + 64 * il + 4 * ir; + + const char * slice = (const char *)src0 + i11*nb02 + i12*nb03; + const int nblocks_per_row = ne00 / QK_K; + const int total_nblocks = ne01 * nblocks_per_row; + const int gblock = i01 * nblocks_per_row + ib_row; // global block index + + const uint8_t * base = (const uint8_t *)slice; + const size_t qs_offset = (size_t)gblock * (QK_K / 2); + const size_t scales_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)gblock * K_SCALE_SIZE; + const size_t dm_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)total_nblocks * K_SCALE_SIZE + (size_t)gblock * sizeof(ggml_half2); + const uint8_t * qs_ptr = base + qs_offset; + const uint8_t * scales_ptr = base + scales_offset; + const ggml_half2 dm_values = *reinterpret_cast(base + dm_offset); + const float dall = dm_values.x(); + const float dmin = dm_values.y(); + + if (tid_local < 12) scales_local[tid_local] = scales_ptr[tid_local]; + item_ct1.barrier(sycl::access::fence_space::local_space); + dequantize_q4_K_common(y, qs_ptr, dall, dmin, scales_local, (int)il, (int)ir); +#else + GGML_UNUSED(src0);GGML_UNUSED(src1);GGML_UNUSED(dst);GGML_UNUSED(ne00);GGML_UNUSED(ne01); + GGML_UNUSED(ne12);GGML_UNUSED(s1);GGML_UNUSED(s2);GGML_UNUSED(s3);GGML_UNUSED(nb02);GGML_UNUSED(nb03); + GGML_UNUSED(s10);GGML_UNUSED(s11);GGML_UNUSED(s12);GGML_UNUSED(scales_local);GGML_UNUSED(item_ct1); + GGML_ABORT("Q4_K reorder get_rows not supported for QK_K != 256"); +#endif +} + +template +static void k_get_rows_q5_K_reorder( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne01, int64_t ne12, + size_t s1, size_t s2, size_t s3, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + uint8_t * scales_local, const sycl::nd_item<3> &item_ct1) { +#if QK_K == 256 + const int64_t ib_row = item_ct1.get_group(2); + const int64_t i10 = item_ct1.get_group(1); + const int64_t tid = item_ct1.get_local_id(2); + const int64_t il = tid / 16; + const int64_t ir = tid % 16; + const int64_t is = 2 * il; + + const int64_t i11 = i10 / ne12; + const int64_t i12 = i10 % ne12; + const int64_t i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + dst_t * y = dst_row + ib_row * QK_K + 64 * il + 2 * ir; + + const char * slice = (const char *)src0 + i11*nb02 + i12*nb03; + const int nblocks_per_row = ne00 / QK_K; + const int total_nblocks = ne01 * nblocks_per_row; + const int gblock = i01 * nblocks_per_row + ib_row; + + const uint8_t * base = (const uint8_t *)slice; + const size_t qs_offset = (size_t)gblock * (QK_K / 2); + const size_t qh_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)gblock * (QK_K / 8); + const size_t scales_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)total_nblocks * (QK_K / 8) + (size_t)gblock * K_SCALE_SIZE; + const size_t dm_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)total_nblocks * (QK_K / 8) + (size_t)total_nblocks * K_SCALE_SIZE + (size_t)gblock * sizeof(ggml_half2); + const uint8_t * qs_ptr = base + qs_offset; + const uint8_t * qh_ptr = base + qh_offset; + const uint8_t * scales_ptr = base + scales_offset; + const ggml_half2 dm_values = *reinterpret_cast(base + dm_offset); + const float dall = dm_values.x(); + const float dmin = dm_values.y(); + + const uint8_t * ql = qs_ptr + 32 * il + 2 * ir; + const uint8_t * qh = qh_ptr + 2 * ir; + if (tid < K_SCALE_SIZE) scales_local[tid] = scales_ptr[tid]; + item_ct1.barrier(sycl::access::fence_space::local_space); + + uint8_t sc, m; + get_scale_min_k4(is + 0, scales_local, sc, m); + const float d1 = dall * sc; const float m1 = dmin * m; + get_scale_min_k4(is + 1, scales_local, sc, m); + const float d2 = dall * sc; const float m2 = dmin * m; + + uint8_t hm = 1 << (2 * il); + y[ 0] = d1 * ((ql[ 0] & 0xF) + (qh[ 0] & hm ? 16 : 0)) - m1; + y[ 1] = d1 * ((ql[ 1] & 0xF) + (qh[ 1] & hm ? 16 : 0)) - m1; + hm <<= 1; + y[32] = d2 * ((ql[ 0] >> 4) + (qh[ 0] & hm ? 16 : 0)) - m2; + y[33] = d2 * ((ql[ 1] >> 4) + (qh[ 1] & hm ? 16 : 0)) - m2; +#else + GGML_UNUSED(src0);GGML_UNUSED(src1);GGML_UNUSED(dst);GGML_UNUSED(ne00);GGML_UNUSED(ne01); + GGML_UNUSED(ne12);GGML_UNUSED(s1);GGML_UNUSED(s2);GGML_UNUSED(s3);GGML_UNUSED(nb02);GGML_UNUSED(nb03); + GGML_UNUSED(s10);GGML_UNUSED(s11);GGML_UNUSED(s12);GGML_UNUSED(scales_local);GGML_UNUSED(item_ct1); + GGML_ABORT("Q5_K reorder get_rows not supported for QK_K != 256"); +#endif +} + +template +static void k_get_rows_q3_K_reorder( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne01, int64_t ne12, + size_t s1, size_t s2, size_t s3, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + const sycl::nd_item<3> &item_ct1) { +#if QK_K == 256 + const int64_t ib_row = item_ct1.get_group(2); + const int64_t i10 = item_ct1.get_group(1); + const int64_t tid = item_ct1.get_local_id(2); + + const int64_t i11 = i10 / ne12; + const int64_t i12 = i10 % ne12; + const int64_t i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + + const char * slice = (const char *)src0 + i11*nb02 + i12*nb03; + const int nblocks_per_row = ne00 / QK_K; + const int total_nblocks = ne01 * nblocks_per_row; + const int gblock = i01 * nblocks_per_row + ib_row; + + const uint8_t * base = (const uint8_t *)slice; + const size_t qs_offset = (size_t)gblock * (QK_K / 4); + const size_t hmask_offset = (size_t)total_nblocks * (QK_K / 4) + (size_t)gblock * (QK_K / 8); + const size_t scales_offset = (size_t)total_nblocks * (QK_K / 4) + (size_t)total_nblocks * (QK_K / 8) + (size_t)gblock * 12; + const size_t d_offset = (size_t)total_nblocks * (QK_K / 4) + (size_t)total_nblocks * (QK_K / 8) + (size_t)total_nblocks * 12 + (size_t)gblock * sizeof(ggml_half); + const uint8_t * qs = base + qs_offset; + const uint8_t * hmask = base + hmask_offset; + const uint8_t * scales = base + scales_offset; + const float d_all = static_cast(*reinterpret_cast(base + d_offset)); + + const int64_t r = tid / 4; + const int64_t tid4 = r / 2; + const int64_t is0 = r % 2; + const int64_t l0 = 16 * is0 + 4 * (tid % 4); + const int64_t n = tid4 / 4; + const int64_t j = tid4 - 4 * n; + const int64_t is = 8 * n + 2 * j + is0; + const int shift = 2 * j; + uint8_t m = 1 << (4 * n + j); + + uint8_t us = is < 4 + ? (scales[is - 0] & 0xF) | (((scales[is + 8] >> 0) & 3) << 4) + : is < 8 + ? (scales[is - 0] & 0xF) | (((scales[is + 4] >> 2) & 3) << 4) + : is < 12 + ? (scales[is - 8] >> 4) | (((scales[is + 0] >> 4) & 3) << 4) + : (scales[is - 8] >> 4) | (((scales[is - 4] >> 6) & 3) << 4); + const float dl = d_all * (us - 32); + + dst_t * y = dst_row + ib_row * QK_K + 128 * n + 32 * j; + const uint8_t * q = qs + 32 * n; + const uint8_t * hm = hmask; + for (int l = l0; l < l0 + 4; ++l) { + y[l] = dl * ((int8_t) ((q[l] >> shift) & 3) - ((hm[l] & m) ? 0 : 4)); + } +#else + GGML_UNUSED(src0);GGML_UNUSED(src1);GGML_UNUSED(dst);GGML_UNUSED(ne00);GGML_UNUSED(ne01); + GGML_UNUSED(ne12);GGML_UNUSED(s1);GGML_UNUSED(s2);GGML_UNUSED(s3);GGML_UNUSED(nb02);GGML_UNUSED(nb03); + GGML_UNUSED(s10);GGML_UNUSED(s11);GGML_UNUSED(s12);GGML_UNUSED(item_ct1); + GGML_ABORT("Q3_K reorder get_rows not supported for QK_K != 256"); +#endif +} + +template +static void k_get_rows_q6_K_reorder( + const void * src0, const int32_t * src1, dst_t * dst, + int64_t ne00, int64_t ne01, int64_t ne12, + size_t s1, size_t s2, size_t s3, size_t nb02, size_t nb03, + size_t s10, size_t s11, size_t s12, + const sycl::nd_item<3> &item_ct1) { +#if QK_K == 256 + const int64_t ib_row = item_ct1.get_group(2); + const int64_t i10 = item_ct1.get_group(1); + const int64_t tid = item_ct1.get_local_id(2); + const int64_t ip = tid / 32; + const int64_t il = tid - 32 * ip; + const int64_t is = 8 * ip + il / 16; + + const int64_t i11 = i10 / ne12; + const int64_t i12 = i10 % ne12; + const int64_t i01 = src1[i10*s10 + i11*s11 + i12*s12]; + + dst_t * dst_row = dst + i10*s1 + i11*s2 + i12*s3; + dst_t * y = dst_row + ib_row * QK_K + 128 * ip + il; + + const char * slice = (const char *)src0 + i11*nb02 + i12*nb03; + const int nblocks_per_row = ne00 / QK_K; + const int total_nblocks = ne01 * nblocks_per_row; + const int gblock = i01 * nblocks_per_row + ib_row; + + const uint8_t * base = (const uint8_t *)slice; + const size_t ql_offset = (size_t)gblock * (QK_K / 2); + const size_t qh_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)gblock * (QK_K / 4); + const size_t base_scales_offset = (size_t)total_nblocks * (QK_K / 2) + (size_t)total_nblocks * (QK_K / 4) + (size_t)gblock * (QK_K / 16); + const size_t base_d_offset = ((QK_K / 2) + (QK_K / 4) + (QK_K / 16)) * (size_t)total_nblocks; + const uint8_t * ql_ptr = base + ql_offset; + const uint8_t * qh_ptr = base + qh_offset; + const uint8_t * scales_ptr = base + base_scales_offset; + const ggml_half * d = (const ggml_half *)(base + base_d_offset) + gblock; + + const uint8_t * ql = ql_ptr + 64 * ip + il; + const uint8_t qh = *(qh_ptr + 32 * ip + il); + const int8_t * sc = reinterpret_cast(scales_ptr + is); + + y[0] = *d * sc[0] * ((int8_t) ((ql[0] & 0xF) | (((qh >> 0) & 3) << 4)) - 32); + y[32] = *d * sc[2] * ((int8_t) ((ql[32] & 0xF) | (((qh >> 2) & 3) << 4)) - 32); + y[64] = *d * sc[4] * ((int8_t) ((ql[0] >> 4) | (((qh >> 4) & 3) << 4)) - 32); + y[96] = *d * sc[6] * ((int8_t) ((ql[32] >> 4) | (((qh >> 6) & 3) << 4)) - 32); +#else + GGML_UNUSED(src0);GGML_UNUSED(src1);GGML_UNUSED(dst);GGML_UNUSED(ne00);GGML_UNUSED(ne01); + GGML_UNUSED(ne12);GGML_UNUSED(s1);GGML_UNUSED(s2);GGML_UNUSED(s3);GGML_UNUSED(nb02);GGML_UNUSED(nb03); + GGML_UNUSED(s10);GGML_UNUSED(s11);GGML_UNUSED(s12);GGML_UNUSED(item_ct1); + GGML_ABORT("Q6_K reorder get_rows not supported for QK_K != 256"); +#endif +} + + template static void k_get_rows( const void * src0, const int32_t * src1, dst_t * dst, @@ -60,50 +379,6 @@ static void k_get_rows( dst_row[iybs + iqs + y_offset] = v.y(); } -template -static void k_get_rows_f32( - const void * src0, const int32_t * src1, dst_t * dst, - int64_t ne00, - int64_t ne12, - size_t s1, size_t s2, size_t s3, - size_t nb01, size_t nb02, size_t nb03, - 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 void * src0_row = (const char *)src0 + i01*nb01 + i11*nb02 + i12*nb03; - - const int ib = i00/qk; - const int iqs = (i00%qk)/qr; - const int iybs = i00 - i00%qk; - const int y_offset = qr == 1 ? 1 : qk/2; - - float v0; - float v1; - dequantize_kernel(src0_row, ib, iqs, v0, v1); - - dst_row[iybs + iqs + 0] = (dst_t) v0; - dst_row[iybs + iqs + y_offset] = (dst_t) v1; -} - template static void k_get_rows_float( const src0_t * src0, const int32_t * src1, dst_t * dst, @@ -173,11 +448,13 @@ static void get_rows_sycl(ggml_backend_sycl_context & ctx, const ggml_tensor *sr GGML_UNUSED(ctx); } -template -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, - const int32_t *src1_dd, float *dst_dd, - queue_ptr stream) { +// Reorder-aware get_rows for src0 split into [qs|scales|d] by opt_for_reorder. +// Covers Q1_0/Q4_0/Q8_0 (types with a per-element dequantize_kernel_t_reorder). +template +static void get_rows_sycl_reorder(ggml_backend_sycl_context & ctx, 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 @@ -197,8 +474,8 @@ static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor stream->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> item_ct1) { - k_get_rows_f32( - src0_dd, src1_dd, dst_dd, ne00, ne12, s1, s2, + k_get_rows_reorder( + src0_dd, src1_dd, dst_dd, ne00, ne01, ne12, s1, s2, s3, nb01, nb02, nb03, s10, s11, s12, item_ct1); }); @@ -206,6 +483,105 @@ static void get_rows_sycl_f32(ggml_backend_sycl_context & ctx, const ggml_tensor GGML_UNUSED(ctx); } +// K-quant reorder get_rows launchers. grid = (1, n11*n12, nblocks_per_row); +// work-group size = 32 (Q4_K) / 64 (Q5_K, Q6_K) / 128 (Q3_K). Each group handles +// one (batch, block_in_row). +#define GGML_SYCL_GET_ROWS_K_LOCALS(ne00_, ne01_, ne12_, s1_, s2_, s3_, nb02_, nb03_, s10_, s11_, s12_) \ + const int64_t ne00 = ne00_; const int64_t ne01 = ne01_; const int64_t ne12 = ne12_; \ + const size_t s1 = s1_; const size_t s2 = s2_; const size_t s3 = s3_; \ + const size_t nb02 = nb02_; const size_t nb03 = nb03_; \ + const size_t s10 = s10_; const size_t s11 = s11_; const size_t s12 = s12_; + +static void get_rows_sycl_q4_K_reorder(ggml_backend_sycl_context & ctx, 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 + const int nblocks_per_row = ne00 / QK_K; + const sycl::range<3> block_dims(1, 1, 32); + const sycl::range<3> block_nums(1, ne11 * ne12, nblocks_per_row); + 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 scale_acc(sycl::range<1>(12), cgh); + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_q4_K_reorder(src0_dd, src1_dd, dst_dd, ne00, ne01, ne12, + s1, s2, s3, nb02, nb03, s10, s11, s12, get_pointer(scale_acc), item_ct1); + }); + }); + GGML_UNUSED(dst); GGML_UNUSED(ctx); +} + +static void get_rows_sycl_q5_K_reorder(ggml_backend_sycl_context & ctx, 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 + const int nblocks_per_row = ne00 / QK_K; + const sycl::range<3> block_dims(1, 1, 64); + const sycl::range<3> block_nums(1, ne11 * ne12, nblocks_per_row); + 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 scale_acc(sycl::range<1>(K_SCALE_SIZE), cgh); + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_q5_K_reorder(src0_dd, src1_dd, dst_dd, ne00, ne01, ne12, + s1, s2, s3, nb02, nb03, s10, s11, s12, get_pointer(scale_acc), item_ct1); + }); + }); + GGML_UNUSED(dst); GGML_UNUSED(ctx); +} + +static void get_rows_sycl_q3_K_reorder(ggml_backend_sycl_context & ctx, 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 + const int nblocks_per_row = ne00 / QK_K; + const sycl::range<3> block_dims(1, 1, 64); + const sycl::range<3> block_nums(1, ne11 * ne12, nblocks_per_row); + 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->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_q3_K_reorder(src0_dd, src1_dd, dst_dd, ne00, ne01, ne12, + s1, s2, s3, nb02, nb03, s10, s11, s12, item_ct1); + }); + GGML_UNUSED(dst); GGML_UNUSED(ctx); +} + +static void get_rows_sycl_q6_K_reorder(ggml_backend_sycl_context & ctx, 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 + const int nblocks_per_row = ne00 / QK_K; + const sycl::range<3> block_dims(1, 1, 64); + const sycl::range<3> block_nums(1, ne11 * ne12, nblocks_per_row); + 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->parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) { + k_get_rows_q6_K_reorder(src0_dd, src1_dd, dst_dd, ne00, ne01, ne12, + s1, s2, s3, nb02, nb03, s10, s11, s12, item_ct1); + }); + GGML_UNUSED(dst); GGML_UNUSED(ctx); +} + template static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst, @@ -273,6 +649,9 @@ void ggml_sycl_op_get_rows(ggml_backend_sycl_context & ctx, ggml_tensor * dst) { src1_i32, (int32_t *)dst->data, ctx.stream()); break; case GGML_TYPE_Q1_0: + if (ggml_sycl_get_rows_src0_reordered(dst)) { + GGML_ABORT("get_rows reorder for Q1_0 not yet implemented (block_q_t is unspecialized)"); + } get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, src1_i32, (float *)dst->data, ctx.stream()); break; @@ -321,24 +700,39 @@ 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: - get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + get_rows_sycl(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: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_q3_K_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(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: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(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_1: get_rows_sycl(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_K: - get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_q4_K_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(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_Q5_0: get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, @@ -349,16 +743,31 @@ 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: - get_rows_sycl_f32(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_q5_K_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(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: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_q6_K_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(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: - get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, - src1_i32, (float *)dst->data, ctx.stream()); + if (ggml_sycl_get_rows_src0_reordered(dst)) { + get_rows_sycl_reorder(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } else { + get_rows_sycl(ctx, dst->src[0], dst->src[1], dst, (const float *)dst->src[0]->data, + src1_i32, (float *)dst->data, ctx.stream()); + } break; default: // TODO: k-quants From 7a53572884c2bf642d4406c52aa4288bffccc07a Mon Sep 17 00:00:00 2001 From: Xu Han Date: Mon, 3 Aug 2026 04:05:36 +0800 Subject: [PATCH 4/6] sycl : drop debug prints from reorder mul_mat dispatch --- ggml/src/ggml-sycl/ggml-sycl.cpp | 34 -------------------------------- 1 file changed, 34 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index b97c9e185..ffad9f902 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3111,25 +3111,8 @@ 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; @@ -3223,17 +3206,9 @@ static void ggml_sycl_op_mul_mat(ggml_backend_sycl_context & ctx, const ggml_ten } if constexpr (quantize_enabled) { - 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, 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; @@ -4357,7 +4332,6 @@ static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor * ggml_tensor_extra_gpu * extra = static_cast(src0->extra); if (!extra) { - fprintf(stderr, "[DBG opt_for_reorder] src0->extra is NULL! type=%d\n", (int)src0->type); return; } if (extra->optimized_feature.reorder) { @@ -4383,9 +4357,6 @@ static void opt_for_reorder(ggml_backend_sycl_context * ctx, const ggml_tensor * } 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 } @@ -4500,14 +4471,9 @@ 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) { From 332a27ac8024d947b2949a894569e80ceb068b18 Mon Sep 17 00:00:00 2001 From: Xu Han Date: Mon, 3 Aug 2026 04:31:33 +0800 Subject: [PATCH 5/6] sycl : drop verbose comment block in k-quant get_rows kernels --- ggml/src/ggml-sycl/getrows.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 37645f74f..5b5a0d080 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -96,14 +96,6 @@ static void k_get_rows_reorder( } -// ── K-quant reorder get_rows kernels ── -// K-quants (Q3_K/Q4_K/Q5_K/Q6_K) use block-cooperative dequantize (scales -// shared across the 32/64/128-thread work-group). After opt_for_reorder the -// src0 buffer is split into [qs|qh/hmask|scales|d] regions indexed by GLOBAL -// block index. get_rows must read block (src1[batch]*nblocks_per_row + block_in_row) -// and write to dst_row + block_in_row*QK_K. These kernels launch one -// work-group per (batch, block_in_row); group(2)=block_in_row, group(1)=batch. - template static void k_get_rows_q4_K_reorder( const void * src0, const int32_t * src1, dst_t * dst, From cb74b304c861aca25e133fea7e7505c2e74da997 Mon Sep 17 00:00:00 2001 From: Xu Han Date: Mon, 3 Aug 2026 04:37:57 +0800 Subject: [PATCH 6/6] sycl : drop Windows PDB /Zi+/DEBUG cmake (debug-only, not for upstream) --- ggml/src/ggml-sycl/CMakeLists.txt | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index 89eb350d2..a8d9c0d80 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -105,19 +105,6 @@ endif() target_compile_options(ggml-sycl PRIVATE "-Wno-narrowing") -# Generate a PDB (debug symbols) for ggml-sycl.dll so cdb/WinDbg can resolve -# source lines + frames inside ggml_backend_sycl_init (otherwise the UAF/crash -# in this DLL shows only as a raw offset — no PDB, because icpx under a VS -# generator doesn't emit one by default). /Zi emits debug info during -# compile; /DEBUG at link writes the .pdb beside the .dll. Both are -# config-agnostic so Release builds also carry symbols (the PDB is read by -# cdb from the .dll's side; runtime perf is unaffected). icpx on Windows -# uses the MSVC codegen back-end, so these cdb-loadable PDBs carry line info. -if (WIN32) - target_compile_options(ggml-sycl PRIVATE "/Zi") - target_link_options(ggml-sycl PRIVATE "/DEBUG") -endif() - message(STATUS "GGML_SYCL_SUPPORT_LEVEL_ZERO_API ${GGML_SYCL_SUPPORT_LEVEL_ZERO_API}") if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API) # Link against Level Zero loader for direct device memory allocation.