From 64c9fb7c1437bbd05566c6a78ff56dc8b10f254b Mon Sep 17 00:00:00 2001 From: Daniel Wymark Date: Tue, 22 Sep 2026 22:54:37 -0700 Subject: [PATCH 1/3] Rebase the Bonsai decode work onto the PR tip Carry the explicit SIMD PQ2_0 and PTQ1_0 decode kernels, vectorized single-row state copies, the FP16-input GEMM switch, the attention decode knob, the logit probe, the comparison and perf scripts, and the model-shaped test cases onto de7a01b's 16-lane PQ2_0 kernel and gating fixes. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01WskiHpEsZkidixZe6TfM5R --- .gitignore | 2 + examples/CMakeLists.txt | 1 + examples/bonsai-logits/CMakeLists.txt | 4 + examples/bonsai-logits/bonsai-logits.cpp | 88 +++++++++ ggml/src/ggml-sycl/fattn.cpp | 9 + ggml/src/ggml-sycl/getrows.cpp | 31 +++ ggml/src/ggml-sycl/ggml-sycl.cpp | 7 +- ggml/src/ggml-sycl/mmvq.cpp | 240 ++++++++++++++++++++++- scripts/bonsai-compare-logits.py | 69 +++++++ scripts/bonsai-compare-perplexity.py | 51 +++++ scripts/bonsai-perf.py | 112 +++++++++++ tests/test-backend-ops.cpp | 46 +++++ 12 files changed, 655 insertions(+), 5 deletions(-) create mode 100644 examples/bonsai-logits/CMakeLists.txt create mode 100644 examples/bonsai-logits/bonsai-logits.cpp create mode 100644 scripts/bonsai-compare-logits.py create mode 100644 scripts/bonsai-compare-perplexity.py create mode 100644 scripts/bonsai-perf.py diff --git a/.gitignore b/.gitignore index 9b589615a402..aeda72bb0fd1 100644 --- a/.gitignore +++ b/.gitignore @@ -153,3 +153,5 @@ a.out.* AGENTS.local.md .pi/SYSTEM.md + +/scratch/ diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index f24052c8b9e8..c785709df73b 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -15,6 +15,7 @@ llama_add_compile_flags() if (EMSCRIPTEN) else() add_subdirectory(batched) + add_subdirectory(bonsai-logits) add_subdirectory(debug) add_subdirectory(embedding) add_subdirectory(eval-callback) diff --git a/examples/bonsai-logits/CMakeLists.txt b/examples/bonsai-logits/CMakeLists.txt new file mode 100644 index 000000000000..e85349235687 --- /dev/null +++ b/examples/bonsai-logits/CMakeLists.txt @@ -0,0 +1,4 @@ +set(TARGET llama-bonsai-logits) +add_executable(${TARGET} bonsai-logits.cpp) +target_link_libraries(${TARGET} PRIVATE llama ${CMAKE_THREAD_LIBS_INIT}) +target_compile_features(${TARGET} PRIVATE cxx_std_17) diff --git a/examples/bonsai-logits/bonsai-logits.cpp b/examples/bonsai-logits/bonsai-logits.cpp new file mode 100644 index 000000000000..a51ca0e15b45 --- /dev/null +++ b/examples/bonsai-logits/bonsai-logits.cpp @@ -0,0 +1,88 @@ +#include "llama.h" +#include "ggml-backend.h" +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +static void log_message(ggml_log_level level, const char * text, void *) { + if (level <= GGML_LOG_LEVEL_WARN) std::fputs(text, stderr); +} + +int main(int argc, char ** argv) { + if (argc != 8 && argc != 9) { + std::fprintf(stderr, "usage: llama-bonsai-logits model backend-dir corpus output prompt-tokens decode-tokens interval [batch]\n"); + return 2; + } + int prompt_count, decode_count, interval, batch; + try { + prompt_count = std::stoi(argv[5]); + decode_count = std::stoi(argv[6]); + interval = std::stoi(argv[7]); + batch = argc == 9 ? std::stoi(argv[8]) : 512; + } catch (const std::exception &) { + std::fprintf(stderr, "token counts, interval, and batch must be integers\n"); + return 2; + } + if (batch < 1 || batch > 4096) return 2; + if (prompt_count < 1 || decode_count < 1 || interval < 1 || decode_count % interval || + prompt_count > INT_MAX - decode_count) return 2; + std::ifstream corpus_file(argv[3], std::ios::binary); + if (!corpus_file) return 3; + const std::string corpus((std::istreambuf_iterator(corpus_file)), std::istreambuf_iterator()); + if (corpus.size() > INT_MAX) return 3; + llama_log_set(log_message, nullptr); + ggml_backend_load_all_from_path(argv[2]); + llama_backend_init(); + auto mp = llama_model_default_params(); mp.n_gpu_layers = 99; + auto * model = llama_model_load_from_file(argv[1], mp); + if (!model) return 4; + auto cp = llama_context_default_params(); + cp.n_ctx = prompt_count + decode_count; cp.n_batch = batch; cp.n_ubatch = batch; + cp.n_threads = 8; cp.n_threads_batch = 8; + cp.flash_attn_type = LLAMA_FLASH_ATTN_TYPE_ENABLED; + auto * ctx = llama_init_from_model(model, cp); + if (!ctx) return 5; + const auto * vocab = llama_model_get_vocab(model); + const int needed = -llama_tokenize(vocab, corpus.data(), int(corpus.size()), nullptr, 0, true, true); + if (needed < prompt_count + decode_count) return 6; + std::vector tokens(needed); + if (llama_tokenize(vocab, corpus.data(), int(corpus.size()), tokens.data(), needed, true, true) != needed) return 6; + auto * out = std::fopen(argv[4], "wb"); + if (!out) return 7; + const int32_t nv = llama_vocab_n_tokens(vocab), steps = 1 + decode_count / interval; + // The output contains two int32 dimensions followed by full FP32 vocabulary rows. + if (std::fwrite(&nv, sizeof(nv), 1, out) != 1 || std::fwrite(&steps, sizeof(steps), 1, out) != 1) return 9; + auto dump = [&]() { + const float * logits = llama_get_logits(ctx); + for (int i = 0; i < nv; ++i) if (!std::isfinite(logits[i])) return false; + return std::fwrite(logits, sizeof(float), nv, out) == size_t(nv); + }; + std::fprintf(stderr, "prefill batch=%d\n", batch); + const auto begin = std::chrono::steady_clock::now(); + for (int offset = 0; offset < prompt_count; offset += batch) { + const int count = std::min(batch, prompt_count - offset); + if (llama_decode(ctx, llama_batch_get_one(tokens.data() + offset, count))) return 8; + } + if (!dump()) return 9; + for (int i = 0; i < decode_count; ++i) { + if (llama_decode(ctx, llama_batch_get_one(tokens.data() + prompt_count + i, 1))) return 8; + if ((i + 1) % interval == 0) { + if (!dump()) return 9; + std::fprintf(stderr, "checked position=%d elapsed=%.3f seconds\n", prompt_count + i + 1, + std::chrono::duration(std::chrono::steady_clock::now() - begin).count()); + std::fflush(stderr); + } + } + if (std::fclose(out)) return 10; + std::fprintf(stderr, "complete prompt=%d decode=%d positions=%d vocabulary=%d\n", prompt_count, decode_count, steps, nv); + llama_free(ctx); llama_model_free(model); llama_backend_free(); + return 0; +} diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index a85eb721f6af..c41b99fb0cd0 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -247,6 +247,15 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0 && !has_bf16; + // Single-query decode normally bypasses the vector kernel whenever the grouped-query optimization applies, + // because that rule assumes a matrix-engine kernel takes those shapes. This backend has none, so the + // vector kernel can serve decode directly on request. + static const int decode_vec = ggml_sycl_get_env("GGML_SYCL_FA_DECODE_VEC", 0); + if (decode_vec && Q->ne[1] == 1 && can_use_vector_kernel && + !ggml_is_quantized(K->type) && !ggml_is_quantized(V->type)) { + return BEST_FATTN_KERNEL_VEC; + } + // Fused-XMX path: oneDNN Graph SDPA (flash attention). Strictly // additive -- taken only when statically supported, otherwise falls through to VEC/TILE below. if (ggml_sycl_flash_attn_ext_onednn_supported(dst)) { diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 36f840e6f5d1..1e96d8b7880c 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -214,6 +214,37 @@ static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tens GGML_TENSOR_BINARY_OP_LOCALS + static const int single_row_copy = ggml_sycl_get_env("GGML_SYCL_SINGLE_ROW_COPY", 0); + if constexpr (std::is_same_v && std::is_same_v) { + if (single_row_copy && ne01 == 1 && ne02 == 1 && ne03 == 1 && + ggml_nelements(src1) == 1 && ggml_is_contiguous(src0) && ggml_is_contiguous(dst)) { + // A valid index into a single-row tensor is zero. + if (dst_dd == src0_dd) { + return; + } + if (single_row_copy == 2) { + const int64_t vectors = (ne00 + 3) / 4; + const sycl::range<1> local(SYCL_GET_ROWS_BLOCK_SIZE); + const sycl::range<1> global(((vectors + local[0] - 1) / local[0]) * local[0]); + stream->parallel_for(sycl::nd_range<1>(global, local), [=](sycl::nd_item<1> item) { + const int64_t i = item.get_global_id(0); + if (4 * i + 3 < ne00) { + sycl::vec values; + values.load(i, src0_dd); + values.store(i, dst_dd); + } else { + for (int64_t j = 4 * i; j < ne00; ++j) { + dst_dd[j] = src0_dd[j]; + } + } + }); + } else { + stream->memcpy(dst_dd, src0_dd, ne00 * sizeof(src0_t)); + } + return; + } + } + const sycl::range<3> block_dims(1, 1, SYCL_GET_ROWS_BLOCK_SIZE); const int block_num_x = (ne00 + SYCL_GET_ROWS_BLOCK_SIZE - 1) / SYCL_GET_ROWS_BLOCK_SIZE; const sycl::range<3> block_nums(ne11 * ne12, ne10, block_num_x); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 761b7f9c664d..39f1520ac6fe 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -2682,8 +2682,13 @@ inline void ggml_sycl_op_mul_mat_sycl( #ifdef GGML_SYCL_F16 bool use_fp16 = true; // TODO(Yu) SYCL capability check #else - bool use_fp16 = false; + bool use_fp16 = src0->type == GGML_TYPE_PQ2_0 || src0->type == GGML_TYPE_PTQ1_0; #endif + const bool bonsai = src0->type == GGML_TYPE_PQ2_0 || src0->type == GGML_TYPE_PTQ1_0; + static const int bonsai_fp16 = ggml_sycl_get_env("GGML_SYCL_BONSAI_F16", 1); + if (bonsai && !bonsai_fp16) { + use_fp16 = false; + } #if GGML_SYCL_DNNL && defined(GGML_SYCL_HAS_BF16) // Fast path for bf16 src0 diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 36c4c1778769..46795dca2b10 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -6,6 +6,10 @@ #include "quants.hpp" #include "vecdotq.hpp" +#if defined(__INTEL_LLVM_COMPILER) +#include +#endif + template static void mul_mat_vec_q_reorder(const void * __restrict__ vx, const void * __restrict__ vy, float * __restrict__ dst, const int ncols, const int nrows, const sycl::nd_item<3> & nd_item) { @@ -1278,11 +1282,113 @@ static void mul_mat_vec_q1_0_q8_1_sycl_switch_ncols( } } +#if defined(__INTEL_LLVM_COMPILER) +static void mul_mat_vec_ptq1_0_q8_1_esimd(const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, dpct::queue_ptr stream) { + static const int rows = ggml_sycl_get_env("GGML_SYCL_PTQ1_ROWS", 4); + GGML_ASSERT(rows > 0 && rows <= 16 && (rows & (rows - 1)) == 0); + const sycl::range<1> local(rows); + const sycl::range<1> global(((nrows + rows - 1) / rows) * rows); + stream->parallel_for(sycl::nd_range<1>(global, local), + [=](sycl::nd_item<1> item) [[intel::sycl_explicit_simd]] { + using namespace sycl::ext::intel::esimd; + constexpr int lanes = 16; + const int row = item.get_global_id(0); + if (row >= nrows) return; + const uint8_t * weights = static_cast(vx) + + static_cast(row) * (ncols / QK_PTQ1_0) * sizeof(block_ptq1_0); + const uint8_t * input = static_cast(vy); + const simd lane(0, 1); + simd acc = 0.0f; + for (int i = 0; i < ncols / QK_PTQ1_0; i += lanes) { + const simd block = lane + i; + const simd_mask valid = block < static_cast(ncols / QK_PTQ1_0); + const simd wb = block * sizeof(block_ptq1_0); + const simd scale_offsets = wb + offsetof(block_ptq1_0, d); + const simd wd = gather( + reinterpret_cast(weights), scale_offsets, valid, simd(sycl::half(0))); + simd input_words[4]; + simd input_scales[4]; + simd sums[4]; +#pragma unroll + for (int chunk = 0; chunk < 4; ++chunk) { + const simd ab = (block * 4 + chunk) * sizeof(block_q8_1); + const simd aq = ab + 4; + input_words[chunk] = gather(reinterpret_cast(input), + aq, valid, simd(0)); + input_scales[chunk] = gather( + reinterpret_cast(input), ab, valid, simd(sycl::half(0))); + sums[chunk] = 0; + } + simd first_words = gather( + reinterpret_cast(weights), wb, valid, simd(0)); + const simd last_offsets = wb + 16; + simd last_words = gather( + reinterpret_cast(weights), last_offsets, valid, simd(0)); +#pragma unroll + for (int group = 0; group < 6; ++group) { + simd packed; + if (group < 4) { + packed = first_words.select(group * lanes); + } else { + packed = last_words.select((group - 4) * lanes); + } + simd v_lo = (packed & 0xff) | ((packed & 0xff00) << 8); + simd v_hi = ((packed >> 16) & 0xff) | ((packed & 0xff000000) >> 8); +#pragma unroll + for (int trit = 0; trit < 5; ++trit) { + const simd w_lo = v_lo * 3; + const simd w_hi = v_hi * 3; + v_lo = w_lo & 0x00ff00ff; v_hi = w_hi & 0x00ff00ff; + simd values = ((w_lo >> 8) & 0xff) | ((w_lo >> 16) & 0xff00) | + ((w_hi << 8) & 0xff0000) | (w_hi & 0xff000000); + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.bit_cast_view(); + const int element = group < 4 ? trit * 16 + 4 * group : 80 + trit * 8 + 4 * (group - 4); + const int chunk = element / 32; + const simd a = input_words[chunk].select((element % 32) / 4 * lanes); + sums[chunk] = dp4a(sums[chunk], w, a); + } + } + const simd tail_offset = wb + offsetof(block_ptq1_0, qh); + const simd tail = gather( + reinterpret_cast(weights), tail_offset, valid, simd(0)); + simd v = (tail & 0xff) | ((tail & 0xff00) << 8); +#pragma unroll + for (int trit = 0; trit < 4; trit += 2) { + const simd w0 = v * 3; v = w0 & 0x00ff00ff; + const simd w1 = v * 3; v = w1 & 0x00ff00ff; + simd values = ((w0 >> 8) & 0xff) | ((w0 >> 16) & 0xff00) | + ((w1 << 8) & 0xff0000) | (w1 & 0xff000000); + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.bit_cast_view(); + const simd a = input_words[3].select((6 + trit / 2) * lanes); + sums[3] = dp4a(sums[3], w, a); + } + simd weighted_sum = 0.0f; +#pragma unroll + for (int chunk = 0; chunk < 4; ++chunk) { + weighted_sum += input_scales[chunk] * simd(sums[chunk]); + } + acc += simd(wd) * weighted_sum; + } + dst[row] = reduce(acc, std::plus<>{}); + }); +} +#endif + static void mul_mat_vec_ptq1_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_PTQ1_0 == 0); +#if defined(__INTEL_LLVM_COMPILER) + static const int use_esimd = ggml_sycl_get_env("GGML_SYCL_PTQ1_ESIMD", 0); + if (use_esimd && g_ggml_sycl_enable_esimd) { + mul_mat_vec_ptq1_0_q8_1_esimd(vx, vy, dst, ncols, nrows, stream); + return; + } +#endif const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE); @@ -1338,14 +1444,139 @@ static void mul_mat_vec_ptq1_0_q8_1_sycl_switch_ncols( } } +static int pq2_rows_per_group() { + static const int rows = ggml_sycl_get_env("GGML_SYCL_PQ2_ROWS", GGML_SYCL_MMV_Y); + GGML_ASSERT(rows > 0 && rows <= 16 && (rows & (rows - 1)) == 0); + return rows; +} + +#if defined(__INTEL_LLVM_COMPILER) +template +static void mul_mat_vec_pq2_0_q8_1_esimd_impl(const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, dpct::queue_ptr stream) { + const int rows = pq2_rows_per_group(); + const sycl::range<1> local(rows); + const sycl::range<1> global(((nrows + rows - 1) / rows) * rows); + stream->parallel_for(sycl::nd_range<1>(global, local), + [=](sycl::nd_item<1> item) [[intel::sycl_explicit_simd]] { + using namespace sycl::ext::intel::esimd; + const int row = item.get_global_id(0); + if (row >= nrows) { + return; + } + const uint8_t * weights = static_cast(vx) + + static_cast(row) * (ncols / QK_PQ2_0) * sizeof(block_pq2_0); + const uint8_t * input = static_cast(vy); + simd acc = 0.0f; + const simd lane(0, 1); + const simd zero_half = sycl::half(0.0f); + // Each vector lane owns one 32-value activation block. + for (int i = 0; i < ncols / QK8_1; i += lanes) { + const simd q = lane + i; + const simd_mask valid = q < static_cast(ncols / QK8_1); + const simd wb = (q / 4) * sizeof(block_pq2_0); + const simd wo = wb + 2 + (q % 4) * 8; + const simd ab = q * sizeof(block_q8_1); + const simd wd = gather( + reinterpret_cast(weights), wb, valid, zero_half); + const simd ad = gather( + reinterpret_cast(input), ab, valid, zero_half); + simd input_words; + if constexpr (bulk_input) { + const simd input_offsets = ab + 4; + input_words = gather(reinterpret_cast(input), + input_offsets, valid, simd(0)); + } + simd packed_weights; + if constexpr (aligned_weights) { + const simd alignment = (wo + uint32_t(reinterpret_cast(weights) & 3)) & 3; + const simd aligned_offsets = wo - alignment; + packed_weights = gather(reinterpret_cast(weights), + aligned_offsets, valid, simd(0)); + const simd_mask shifted = valid & (alignment == 2); + const simd tail_offsets = wo + 6; + const simd tail = gather( + reinterpret_cast(weights), tail_offsets, shifted, simd(0)); + simd first = packed_weights.template select(0); + simd second = packed_weights.template select(lanes); + // The final halfword completes an unaligned payload without crossing its block boundary. + first.merge((first >> 16) | (second << 16), shifted); + second.merge((second >> 16) | (tail << 16), shifted); + packed_weights.template select(0) = first; + packed_weights.template select(lanes) = second; + } + simd sum = 0; +#pragma unroll + for (int j = 0; j < 4; ++j) { + simd pair; + if constexpr (aligned_weights) { + pair = (packed_weights.template select((j / 2) * lanes) >> ((j % 2) * 16)) & 0xffff; + } else { + const simd offsets = wo + 2 * j; + pair = gather(reinterpret_cast(weights), + offsets, valid, simd(0)); + } +#pragma unroll + for (int k = 0; k < 2; ++k) { + simd values = (pair >> (8 * k)) & 0xff; + values = (values | (values << 12)) & 0x000f000f; + values = (values | (values << 6)) & 0x03030303; + values = ((values | 0x80808080u) - 0x01010101u) ^ 0x80808080u; + const simd w = values.template bit_cast_view(); + const simd offsets_a = ab + 4 + 4 * (2 * j + k); + simd a; + if constexpr (bulk_input) { + a = input_words.template select((2 * j + k) * lanes); + } else { + a = gather(reinterpret_cast(input), + offsets_a, valid, simd(0)); + } + sum = dp4a(sum, w, a); + } + } + acc += simd(wd) * simd(ad) * simd(sum); + } + dst[row] = reduce(acc, std::plus<>{}); + }); +} +template +static void mul_mat_vec_pq2_0_q8_1_esimd(const void * vx, const void * vy, float * dst, + const int ncols, const int nrows, dpct::queue_ptr stream) { + static const int bulk_input = ggml_sycl_get_env("GGML_SYCL_PQ2_BULK_INPUT", 0); + static const int aligned_weights = ggml_sycl_get_env("GGML_SYCL_PQ2_ALIGNED_LOADS", 0); + if (bulk_input && aligned_weights) { + mul_mat_vec_pq2_0_q8_1_esimd_impl(vx, vy, dst, ncols, nrows, stream); + } else if (bulk_input) { + mul_mat_vec_pq2_0_q8_1_esimd_impl(vx, vy, dst, ncols, nrows, stream); + } else { + mul_mat_vec_pq2_0_q8_1_esimd_impl(vx, vy, dst, ncols, nrows, stream); + } +} +#endif + static void mul_mat_vec_pq2_0_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_PQ2_0 == 0); - const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; +#if defined(__INTEL_LLVM_COMPILER) + static const int use_esimd = ggml_sycl_get_env("GGML_SYCL_PQ2_ESIMD", 0); + if (use_esimd && g_ggml_sycl_enable_esimd) { + static const int lanes = ggml_sycl_get_env("GGML_SYCL_PQ2_LANES", 16); + switch (lanes) { + case 8: mul_mat_vec_pq2_0_q8_1_esimd<8> (vx, vy, dst, ncols, nrows, stream); break; + case 16: mul_mat_vec_pq2_0_q8_1_esimd<16>(vx, vy, dst, ncols, nrows, stream); break; + case 32: mul_mat_vec_pq2_0_q8_1_esimd<32>(vx, vy, dst, ncols, nrows, stream); break; + case 64: mul_mat_vec_pq2_0_q8_1_esimd<64>(vx, vy, dst, ncols, nrows, stream); break; + default: GGML_ABORT("unsupported PQ2_0 vector width: %d", lanes); + } + return; + } +#endif + const int rows = pq2_rows_per_group(); + const int block_num_y = (nrows + rows - 1) / rows; const sycl::range<3> block_nums(1, 1, block_num_y); - const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE); + const sycl::range<3> block_dims(1, rows, WARP_SIZE); stream->submit([&](sycl::handler & cgh) { cgh.parallel_for( @@ -1365,9 +1596,10 @@ static void mul_mat_vec_pq2_0_q8_1_sycl_ncols( const int stride_col_y, const int stride_col_dst, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_PQ2_0 == 0); - const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y; + const int rows = pq2_rows_per_group(); + const int block_num_y = (nrows + rows - 1) / rows; const sycl::range<3> block_nums(1, 1, block_num_y); - const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE); + const sycl::range<3> block_dims(1, rows, WARP_SIZE); stream->submit([&](sycl::handler & cgh) { cgh.parallel_for( diff --git a/scripts/bonsai-compare-logits.py b/scripts/bonsai-compare-logits.py new file mode 100644 index 000000000000..27d45a4404a3 --- /dev/null +++ b/scripts/bonsai-compare-logits.py @@ -0,0 +1,69 @@ +#!/usr/bin/env python3 +"""Compare float32 logit dumps with int32 vocabulary-size and row-count headers.""" + +import argparse +from array import array +import json +import math +from pathlib import Path +import struct + + +def read(path): + with path.open("rb") as source: + vocabulary, rows = struct.unpack("= 0 for value in values) or values[-2] <= 0: + raise ValueError(f"Invalid perplexity values: {path}") + return {"path": str(path), "context": int(setup[2]), "chunks": chunks, + "predictions": chunks * (int(setup[2]) // 2 - 1), + "perplexity": float(final[1]), "standard_error": float(final[2])} + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("reference", type=Path) + parser.add_argument("candidate", type=Path) + parser.add_argument("output", type=Path) + parser.add_argument("--chunks", type=int, required=True) + parser.add_argument("--max-relative-increase", type=float, required=True) + args = parser.parse_args() + if args.chunks <= 0 or not math.isfinite(args.max_relative_increase) or args.max_relative_increase < 0: + parser.error("chunks must be positive and the limit must be finite and nonnegative") + reference = read_result(args.reference, args.chunks) + candidate = read_result(args.candidate, args.chunks) + if reference["context"] != candidate["context"]: + raise ValueError("Context sizes differ") + change = candidate["perplexity"] / reference["perplexity"] - 1 + result = {"reference": reference, "candidate": candidate, "relative_increase": change, + "limit": args.max_relative_increase, "passed": change <= args.max_relative_increase} + args.output.write_text(json.dumps(result, indent=2) + "\n", encoding="utf-8") + print(json.dumps(result)) + return 0 if result["passed"] else 1 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/bonsai-perf.py b/scripts/bonsai-perf.py new file mode 100644 index 000000000000..4172c3d9f2f2 --- /dev/null +++ b/scripts/bonsai-perf.py @@ -0,0 +1,112 @@ +#!/usr/bin/env python3 +"""Run sequential benchmark jobs from a JSON specification with bounded lifetimes.""" + +import argparse +import datetime as dt +import hashlib +import json +import os +from pathlib import Path +import signal +import subprocess +import time + + +def now(): + return dt.datetime.now(dt.timezone.utc) + + +def write_json(path, value): + temporary = path.with_suffix(path.suffix + ".tmp") + temporary.write_text(json.dumps(value, indent=2) + "\n", encoding="utf-8") + temporary.replace(path) + + +def terminate_tree(process): + if process.poll() is not None: + return + if os.name == "nt": + subprocess.run(["taskkill", "/PID", str(process.pid), "/T", "/F"], check=False, stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL) + else: + os.killpg(process.pid, signal.SIGKILL) + process.wait(timeout=30) + + +def power_state(): + if os.name != "nt": + return None + result = subprocess.run([ + "powershell.exe", "-NoProfile", "-Command", + "Get-CimInstance -Namespace root/wmi -ClassName BatteryStatus | Select-Object PowerOnline,Discharging | ConvertTo-Json -Compress", + ], capture_output=True, text=True, timeout=20, check=False) + return {"returncode": result.returncode, "output": result.stdout.strip(), "error": result.stderr.strip()} + + +def run(spec_path): + spec = json.loads(spec_path.read_text(encoding="utf-8")) + output = Path(spec["output"]) + if (output / "summary.json").exists(): + raise FileExistsError(output / "summary.json") + output.mkdir(parents=True, exist_ok=True) + deadline = dt.datetime.fromisoformat(spec["deadline_utc"]) + if deadline.utcoffset() is None: + raise ValueError("deadline_utc must include a timezone") + jobs = spec["jobs"] + names = [job["name"] for job in jobs] + if len(set(names)) != len(names) or any(Path(name).name != name for name in names): + raise ValueError("job names must be unique file names") + environment = os.environ.copy() + environment.update(spec.get("env", {})) + summary = {"started_utc": now().isoformat(), "spec": spec, "jobs": []} + write_json(output / "summary.json", summary) + for job in jobs: + record_path = output / (job["name"] + ".json") + if record_path.exists(): + raise FileExistsError(record_path) + record = {"name": job["name"], "argv": job["argv"], "started_utc": now().isoformat(), "source_commit": job.get("source_commit")} + job_environment = environment.copy() + job_environment.update(job.get("env", {})) + record["env"] = dict(spec.get("env", {}), **job.get("env", {})) + record["power_before"] = power_state() + timeout = min(float(job["timeout_seconds"]), (deadline - now()).total_seconds()) + if timeout <= 0: + record["status"] = "deadline" + record["returncode"] = None + else: + executable = Path(job["argv"][0]) + if executable.is_file(): + record["executable_sha256"] = hashlib.sha256(executable.read_bytes()).hexdigest() + options = {"creationflags": subprocess.CREATE_NEW_PROCESS_GROUP} if os.name == "nt" else {"start_new_session": True} + started = time.monotonic() + with (output / (job["name"] + ".log")).open("wb") as log: + process = subprocess.Popen(job["argv"], cwd=spec["cwd"], env=job_environment, stdout=log, stderr=subprocess.STDOUT, **options) + record["pid"] = process.pid + record["status"] = "running" + write_json(record_path, record) + print(json.dumps(record), flush=True) + try: + record["returncode"] = process.wait(timeout=timeout) + record["status"] = "passed" if process.returncode == 0 else "failed" + except subprocess.TimeoutExpired: + terminate_tree(process) + record["returncode"] = process.returncode + record["status"] = "timeout" + finally: + terminate_tree(process) + record["elapsed_seconds"] = time.monotonic() - started + record["finished_utc"] = now().isoformat() + record["power_after"] = power_state() + write_json(record_path, record) + summary["jobs"].append(record) + summary["finished_utc"] = now().isoformat() + write_json(output / "summary.json", summary) + print(json.dumps(record), flush=True) + if record["status"] != "passed" and not spec.get("continue_on_failure"): + return 1 + return 0 + + +if __name__ == "__main__": + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("spec", type=Path) + raise SystemExit(run(parser.parse_args().spec)) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 1363837294df..aee5fe2fafbb 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4633,6 +4633,28 @@ struct test_mul_mat : public test_case { } }; +struct test_mul_mat_small_f16 : public test_mul_mat { + explicit test_mul_mat_small_f16(int64_t tokens) + : test_mul_mat(GGML_TYPE_PQ2_0, GGML_TYPE_F32, 67, tokens, 256, {1, 1}, {1, 1}) {} + + std::string vars() override { return test_mul_mat::vars() + ",small_f16=1"; } + + void initialize_tensors(ggml_context * ctx) override { + for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) { + if (strcmp(t->name, "b") == 0) { + std::vector data(ggml_nelements(t)); + for (size_t i = 0; i < data.size(); ++i) { + const int value = i % 32 == 0 ? 127 : int((i * 17) % 255) - 127; + data[i] = std::ldexp(float(value), -24); + } + ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float)); + } else { + init_tensor_uniform(t); + } + } + } +}; + // GGML_HINT_SRC0_IS_HADAMARD struct test_mul_mat_hadamard : public test_mul_mat { test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32, @@ -8505,6 +8527,10 @@ static std::vector> make_test_cases_eval() { } test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false)); + for (int n : {127, 30720, 786432}) { + test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, n, 1, 1, 1, 1, false)); + } + for (ggml_type type : all_types) { for (int b : {1, 7}) { for (bool v : {false, true}) { @@ -9445,6 +9471,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 6, 4096, 5120, {1, 1}, {1, 1})); + for (int64_t tokens : {9, 17, 65, 256}) { + test_cases.emplace_back(new test_mul_mat_small_f16(tokens)); + } + // K not a multiple of 32 test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 64, 32, 65, {1, 1}, {1, 1})); test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 64, 32, 80, {1, 1}, {1, 1})); @@ -10290,6 +10320,22 @@ static std::vector> make_test_cases_eval() { static std::vector> make_test_cases_perf() { std::vector> test_cases; + for (int n : {127, 30720, 786432}) { + test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, n, 1, 1, 1, 1, false)); + } + + // Ternary Bonsai 2 27B single-stream decode shapes: grouped-query attention at depth and the PQ2_0 + // matrix-vector products behind one token. + for (int kv : { 2048, 4096, 8192 }) { + test_cases.emplace_back(new test_flash_attn_ext(256, 256, 4, {6, 1}, kv, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + } + for (int64_t m : { 17408, 12288, 10240, 6144, 1024, 248320 }) { + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_PQ2_0, GGML_TYPE_F32, m, 1, 5120, {1, 1}, {1, 1})); + } + for (int64_t k : { 17408, 6144 }) { + test_cases.emplace_back(new test_mul_mat(GGML_TYPE_PQ2_0, GGML_TYPE_F32, 5120, 1, k, {1, 1}, {1, 1})); + } + // SWIGLU at a 27B-class FFN width, fused [gate|up] vs split operands // note: same bytes either way, so a backend that indexes them differently shows it here for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) { From 01781f2829d0a9065ca514dd155f3de15f57f758 Mon Sep 17 00:00:00 2001 From: Daniel Wymark Date: Wed, 23 Sep 2026 00:00:05 -0700 Subject: [PATCH 2/3] examples: let the Bonsai logit probe run on the CPU backend alone Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01WskiHpEsZkidixZe6TfM5R --- examples/bonsai-logits/bonsai-logits.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/examples/bonsai-logits/bonsai-logits.cpp b/examples/bonsai-logits/bonsai-logits.cpp index a51ca0e15b45..25e166bb42f2 100644 --- a/examples/bonsai-logits/bonsai-logits.cpp +++ b/examples/bonsai-logits/bonsai-logits.cpp @@ -7,6 +7,7 @@ #include #include #include +#include #include #include #include @@ -41,7 +42,13 @@ int main(int argc, char ** argv) { llama_log_set(log_message, nullptr); ggml_backend_load_all_from_path(argv[2]); llama_backend_init(); - auto mp = llama_model_default_params(); mp.n_gpu_layers = 99; + auto mp = llama_model_default_params(); + // BONSAI_LOGITS_NGL selects the offload depth; zero also withholds every device so the + // scheduler cannot route prompt matrix products to an accelerator. + const char * ngl_env = std::getenv("BONSAI_LOGITS_NGL"); + mp.n_gpu_layers = ngl_env ? std::stoi(ngl_env) : 99; + static ggml_backend_dev_t no_devices[] = { nullptr }; + if (mp.n_gpu_layers == 0) { mp.devices = no_devices; } auto * model = llama_model_load_from_file(argv[1], mp); if (!model) return 4; auto cp = llama_context_default_params(); From 1bd1b5da3312e1bc8a78bd6b5bd98f17a76d59fb Mon Sep 17 00:00:00 2001 From: Daniel Wymark Date: Wed, 23 Sep 2026 08:30:53 -0700 Subject: [PATCH 3/3] Tighten the comments on the opt-in decode paths Each experimental kernel now carries a short signpost naming what it does and which environment variable turns it on, and the remaining comments state present behaviour rather than the case they guard against. Co-Authored-By: Claude Fable 5.1 Claude-Session: https://claude.ai/code/session_01WskiHpEsZkidixZe6TfM5R --- ggml/src/ggml-sycl/fattn.cpp | 5 ++--- ggml/src/ggml-sycl/getrows.cpp | 2 +- ggml/src/ggml-sycl/mmvq.cpp | 9 ++++++++- 3 files changed, 11 insertions(+), 5 deletions(-) diff --git a/ggml/src/ggml-sycl/fattn.cpp b/ggml/src/ggml-sycl/fattn.cpp index c41b99fb0cd0..d82ec5fe01ef 100644 --- a/ggml/src/ggml-sycl/fattn.cpp +++ b/ggml/src/ggml-sycl/fattn.cpp @@ -247,9 +247,8 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const const bool can_use_vector_kernel = Q->ne[0] <= 512 && Q->ne[0] % 64 == 0 && K->ne[1] % FATTN_KQ_STRIDE == 0 && !has_bf16; - // Single-query decode normally bypasses the vector kernel whenever the grouped-query optimization applies, - // because that rule assumes a matrix-engine kernel takes those shapes. This backend has none, so the - // vector kernel can serve decode directly on request. + // Serve single-query decode with the vector kernel on request. The grouped-query rule below assumes a + // matrix-engine kernel exists for those shapes, and this backend has none. static const int decode_vec = ggml_sycl_get_env("GGML_SYCL_FA_DECODE_VEC", 0); if (decode_vec && Q->ne[1] == 1 && can_use_vector_kernel && !ggml_is_quantized(K->type) && !ggml_is_quantized(V->type)) { diff --git a/ggml/src/ggml-sycl/getrows.cpp b/ggml/src/ggml-sycl/getrows.cpp index 1e96d8b7880c..256f5c0a4676 100644 --- a/ggml/src/ggml-sycl/getrows.cpp +++ b/ggml/src/ggml-sycl/getrows.cpp @@ -218,7 +218,7 @@ static void get_rows_sycl_float(ggml_backend_sycl_context & ctx, const ggml_tens if constexpr (std::is_same_v && std::is_same_v) { if (single_row_copy && ne01 == 1 && ne02 == 1 && ne03 == 1 && ggml_nelements(src1) == 1 && ggml_is_contiguous(src0) && ggml_is_contiguous(dst)) { - // A valid index into a single-row tensor is zero. + // The only valid row index is zero, so src1 is never read. if (dst_dd == src0_dd) { return; } diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 46795dca2b10..f33917d54466 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1283,6 +1283,8 @@ static void mul_mat_vec_q1_0_q8_1_sycl_switch_ncols( } #if defined(__INTEL_LLVM_COMPILER) +// Explicit-SIMD decode kernel: one work-item per weight row, with the vector lanes spread across that +// row's blocks. Opt-in through GGML_SYCL_PTQ1_ESIMD; GGML_SYCL_PTQ1_ROWS sets the work-group size. static void mul_mat_vec_ptq1_0_q8_1_esimd(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { static const int rows = ggml_sycl_get_env("GGML_SYCL_PTQ1_ROWS", 4); @@ -1444,6 +1446,7 @@ static void mul_mat_vec_ptq1_0_q8_1_sycl_switch_ncols( } } +// Rows per work-group for the PQ2_0 matrix-vector kernels, from GGML_SYCL_PQ2_ROWS. static int pq2_rows_per_group() { static const int rows = ggml_sycl_get_env("GGML_SYCL_PQ2_ROWS", GGML_SYCL_MMV_Y); GGML_ASSERT(rows > 0 && rows <= 16 && (rows & (rows - 1)) == 0); @@ -1451,6 +1454,10 @@ static int pq2_rows_per_group() { } #if defined(__INTEL_LLVM_COMPILER) +// Explicit-SIMD decode kernel: one work-item per weight row, with the vector lanes spread across that +// row's activation blocks. Opt-in through GGML_SYCL_PQ2_ESIMD. GGML_SYCL_PQ2_LANES picks the vector +// width, GGML_SYCL_PQ2_BULK_INPUT gathers each lane's activations in one load, and +// GGML_SYCL_PQ2_ALIGNED_LOADS reads weights as aligned 32-bit words. template static void mul_mat_vec_pq2_0_q8_1_esimd_impl(const void * vx, const void * vy, float * dst, const int ncols, const int nrows, dpct::queue_ptr stream) { @@ -1499,7 +1506,7 @@ static void mul_mat_vec_pq2_0_q8_1_esimd_impl(const void * vx, const void * vy, reinterpret_cast(weights), tail_offsets, shifted, simd(0)); simd first = packed_weights.template select(0); simd second = packed_weights.template select(lanes); - // The final halfword completes an unaligned payload without crossing its block boundary. + // A payload that starts mid-word is shifted down and completed from the halfword after it. first.merge((first >> 16) | (second << 16), shifted); second.merge((second >> 16) | (tail << 16), shifted); packed_weights.template select(0) = first;