From f91123d2d02907c665cee42d6428eee5ee35d356 Mon Sep 17 00:00:00 2001 From: Thiago Padilha Date: Fri, 28 Aug 2026 05:33:37 -0300 Subject: [PATCH] qwen4exp: fix sparse-attention block selection Build QSA blocks per sequence in token order and select complete blocks before expanding them to cache cells. Keep only the incomplete tail unconditionally visible and rotate pooled keys with the first token's full M-RoPE position. This prevents unified-cache sequences from sharing pooled indexer keys and avoids replacing padded tail entries with extra history tokens. Synthetic Qwen4 architecture, exact mask, F16 and Q8_0 state, sequence-copy, Metal, and AddressSanitizer checks pass. Assisted-by: Codex qwen4exp: support independent PLE embedding widths Size the PLE key and value projections from the concatenated n-gram embedding instead of assuming it matches the model hidden width. Validate the head count before narrowing it to the stored type. Add a synthetic PLE model with a 64-wide embedding and a 256-wide hidden state, then verify inference and model roundtrip. Assisted-by: Codex qwen4exp: validate model metadata Reject invalid GDN, hyper-connection, QSA, and PLE dimensions during model loading instead of aborting later while building the graph. Validate PLE array lengths before copying them into fixed storage. The released configuration and synthetic Qwen4 architecture tests pass. Assisted-by: Codex qwen4exp: update indexer cache after sequence copies Treat cached indexer keys as unrotated data and apply pending cache updates alongside the attention and recurrent state. This copies indexer data during non-unified cross-stream sequence copies without applying RoPE shifts to raw keys. Assisted-by: Codex qwen4exp: enable recurrent state rollback Assisted-by: Codex qwen4exp: disable tensor split Assisted-by: Codex metal: align dynamic threadgroup memory Assisted-by: Codex metal: widen expert matmul thread index Assisted-by: Codex --- ggml/src/ggml-metal/ggml-metal-device.m | 2 +- ggml/src/ggml-metal/kernels/mul_mm.metal | 2 +- src/llama-arch.cpp | 2 + src/llama-memory-hybrid-idx.cpp | 445 ++++++++++++++++------- src/llama-memory-hybrid-idx.h | 39 +- src/models/models.h | 14 +- src/models/qwen4exp.cpp | 414 ++++++++++----------- tests/CMakeLists.txt | 10 + tests/test-backend-ops.cpp | 1 + tests/test-llama-archs.cpp | 100 ++++- 10 files changed, 657 insertions(+), 372 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 41ce90dc8a..7ad83ae4ce 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -800,7 +800,7 @@ void ggml_metal_encoder_set_buffer(ggml_metal_encoder_t encoder, struct ggml_met } void ggml_metal_encoder_set_threadgroup_memory_size(ggml_metal_encoder_t encoder, size_t size, int idx) { - [encoder->obj setThreadgroupMemoryLength:size atIndex:idx]; + [encoder->obj setThreadgroupMemoryLength:GGML_PAD(size, 16) atIndex:idx]; } void ggml_metal_encoder_dispatch_threadgroups(ggml_metal_encoder_t encoder, int tg0, int tg1, int tg2, int tptg0, int tptg1, int tptg2) { diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal index ee848eed6d..1339525b45 100644 --- a/ggml/src/ggml-metal/kernels/mul_mm.metal +++ b/ggml/src/ggml-metal/kernels/mul_mm.metal @@ -435,7 +435,7 @@ kernel void kernel_mul_mm_id( device char * dst, threadgroup char * shmem [[threadgroup(0)]], uint3 tgpig[[threadgroup_position_in_grid]], - ushort tiitg[[thread_index_in_threadgroup]], + uint tiitg[[thread_index_in_threadgroup]], ushort tiisg[[thread_index_in_simdgroup]], ushort sgitg[[simdgroup_index_in_threadgroup]]) { threadgroup S0 * sa = (threadgroup S0 *)(shmem); diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 5e61f61f7f..6f781949d2 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -1100,6 +1100,7 @@ bool llm_arch_supports_rs_rollback(const llm_arch & arch) { switch (arch) { case LLM_ARCH_QWEN35: case LLM_ARCH_QWEN35MOE: + case LLM_ARCH_QWEN4EXP: case LLM_ARCH_DEEPSEEK4: case LLM_ARCH_NEMOTRON_H: case LLM_ARCH_NEMOTRON_H_MOE: @@ -1141,6 +1142,7 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) { case LLM_ARCH_KIMI_LINEAR: case LLM_ARCH_BAILINGMOE3: case LLM_ARCH_KIMI_K3: + case LLM_ARCH_QWEN4EXP: case LLM_ARCH_QWEN3TTS: return false; default: diff --git a/src/llama-memory-hybrid-idx.cpp b/src/llama-memory-hybrid-idx.cpp index d4e59d77e5..1575f01893 100644 --- a/src/llama-memory-hybrid-idx.cpp +++ b/src/llama-memory-hybrid-idx.cpp @@ -9,7 +9,9 @@ #include #include #include +#include #include +#include // // llama_memory_hybrid_idx @@ -47,6 +49,7 @@ llama_memory_hybrid_idx::llama_memory_hybrid_idx( hparams_idx(model.hparams), mem_idx(filter_idx == nullptr ? nullptr : [&] { // MQA with a single key head of indexer_head_size, as llama_kv_cache_dsa shapes its own + hparams_idx.rope_type = LLAMA_ROPE_TYPE_NONE; std::fill(hparams_idx.n_head_kv_arr.begin(), hparams_idx.n_head_kv_arr.end(), 1); hparams_idx.n_embd_head_k_full = model.hparams.indexer_head_size; @@ -137,6 +140,8 @@ void llama_memory_hybrid_idx::clear(bool data) { if (mem_idx) { mem_idx->clear(data); } + + qsa_histories.clear(); } bool llama_memory_hybrid_idx::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) { @@ -149,15 +154,96 @@ bool llama_memory_hybrid_idx::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_po mem_idx->seq_rm(seq_id, p0, p1); } - return get_mem_attn()->seq_rm(seq_id, p0, p1); + const bool res = get_mem_attn()->seq_rm(seq_id, p0, p1); + if (!res) { + return false; + } + + auto remove = [&](qsa_history & history) { + history.erase(std::remove_if(history.begin(), history.end(), [&](const qsa_token & token) { + return (p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1); + }), history.end()); + }; + + if (seq_id < 0) { + for (auto & item : qsa_histories) { + remove(item.second); + } + } else { + auto it = qsa_histories.find(seq_id); + if (it != qsa_histories.end()) { + remove(it->second); + if (it->second.empty()) { + qsa_histories.erase(it); + } + } + } + + return true; } void llama_memory_hybrid_idx::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) { + if (seq_id_src == seq_id_dst) { + return; + } + + qsa_history copied; + const auto & cells_src = get_mem_attn()->get_cells(seq_id_src); + const auto & cells_dst = get_mem_attn()->get_cells(seq_id_dst); + const bool replace = &cells_src != &cells_dst; + const auto src = qsa_histories.find(seq_id_src); + if (src != qsa_histories.end()) { + using pos_key = std::tuple; + std::map> cells_by_pos; + for (uint32_t cell = 0; cell < cells_src.size(); ++cell) { + if (cells_src.is_empty(cell) || !cells_src.seq_has(cell, seq_id_src)) { + continue; + } + + const llama_pos pos = cells_src.pos_get(cell); + if ((p0 >= 0 && pos < p0) || (p1 >= 0 && pos >= p1)) { + continue; + } + + const auto & ext = cells_src.ext_get(cell); + cells_by_pos[{ pos, ext.y, ext.x }].push_back(!replace && cells_src.seq_has(cell, seq_id_dst)); + } + + std::map next_cell; + for (const auto & token : src->second) { + if ((p0 >= 0 && token.pos[0] < p0) || (p1 >= 0 && token.pos[0] >= p1)) { + continue; + } + + const pos_key key = { token.pos[0], token.pos[1], token.pos[2] }; + auto cells = cells_by_pos.find(key); + if (cells == cells_by_pos.end()) { + continue; + } + + size_t & index = next_cell[key]; + if (index < cells->second.size() && !cells->second[index++]) { + copied.push_back(token); + } + } + } + llama_memory_hybrid::seq_cp(seq_id_src, seq_id_dst, p0, p1); if (mem_idx) { mem_idx->seq_cp(seq_id_src, seq_id_dst, p0, p1); } + + if (replace) { + if (copied.empty()) { + qsa_histories.erase(seq_id_dst); + } else { + qsa_histories[seq_id_dst] = std::move(copied); + } + } else if (!copied.empty()) { + auto & dst = qsa_histories[seq_id_dst]; + dst.insert(dst.end(), copied.begin(), copied.end()); + } } void llama_memory_hybrid_idx::seq_keep(llama_seq_id seq_id) { @@ -166,6 +252,13 @@ void llama_memory_hybrid_idx::seq_keep(llama_seq_id seq_id) { if (mem_idx) { mem_idx->seq_keep(seq_id); } + + auto it = qsa_histories.find(seq_id); + qsa_history keep = it == qsa_histories.end() ? qsa_history{} : std::move(it->second); + qsa_histories.clear(); + if (!keep.empty()) { + qsa_histories.emplace(seq_id, std::move(keep)); + } } void llama_memory_hybrid_idx::seq_add(llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) { @@ -174,6 +267,15 @@ void llama_memory_hybrid_idx::seq_add(llama_seq_id seq_id, llama_pos p0, llama_p if (mem_idx) { mem_idx->seq_add(seq_id, p0, p1, shift); } + + auto it = qsa_histories.find(seq_id); + if (it != qsa_histories.end()) { + for (auto & token : it->second) { + if ((p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1)) { + token.pos[0] += shift; + } + } + } } void llama_memory_hybrid_idx::seq_div(llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) { @@ -182,6 +284,15 @@ void llama_memory_hybrid_idx::seq_div(llama_seq_id seq_id, llama_pos p0, llama_p if (mem_idx) { mem_idx->seq_div(seq_id, p0, p1, d); } + + auto it = qsa_histories.find(seq_id); + if (it != qsa_histories.end()) { + for (auto & token : it->second) { + if ((p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1)) { + token.pos[0] /= d; + } + } + } } std::map llama_memory_hybrid_idx::memory_breakdown() const { @@ -205,8 +316,28 @@ void llama_memory_hybrid_idx::state_write(llama_io_write_i & io, llama_seq_id se if (mem_idx) { mem_idx->state_write(io, seq_id, flags); } - } + uint32_t n_histories = 0; + if (seq_id < 0) { + n_histories = (uint32_t) qsa_histories.size(); + } else if (qsa_histories.count(seq_id) != 0) { + n_histories = 1; + } + io.write(&n_histories, sizeof(n_histories)); + + for (const auto & item : qsa_histories) { + if (seq_id >= 0 && item.first != seq_id) { + continue; + } + + io.write(&item.first, sizeof(item.first)); + const uint64_t n_tokens = item.second.size(); + io.write(&n_tokens, sizeof(n_tokens)); + for (const auto & token : item.second) { + io.write(token.pos.data(), sizeof(token.pos)); + } + } + } } void llama_memory_hybrid_idx::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) { @@ -230,6 +361,34 @@ void llama_memory_hybrid_idx::state_read(llama_io_read_i & io, llama_seq_id seq_ if (mem_idx) { mem_idx->state_read_sinfo(io, seq_id, flags, nullptr, &sinfos_attn); } + + uint32_t n_histories; + io.read(&n_histories, sizeof(n_histories)); + if (n_histories > LLAMA_MAX_SEQ) { + throw std::runtime_error("invalid QSA history count"); + } + + if (seq_id < 0) { + qsa_histories.clear(); + } else { + qsa_histories.erase(seq_id); + } + + for (uint32_t ih = 0; ih < n_histories; ++ih) { + llama_seq_id stored_seq; + uint64_t n_tokens; + io.read(&stored_seq, sizeof(stored_seq)); + io.read(&n_tokens, sizeof(n_tokens)); + if (stored_seq < 0 || stored_seq >= LLAMA_MAX_SEQ || n_tokens > get_mem_attn()->get_size()) { + throw std::runtime_error("invalid QSA history"); + } + + auto & history = qsa_histories[seq_id < 0 ? stored_seq : seq_id]; + history.resize(n_tokens); + for (auto & token : history) { + io.read(token.pos.data(), sizeof(token.pos)); + } + } } } catch (...) { @@ -255,12 +414,156 @@ void llama_memory_hybrid_idx::state_drop(llama_seq_id seq_id) { if (mem_idx) { mem_idx->seq_rm(seq_id, -1, -1); } + + qsa_histories.erase(seq_id); } llama_kv_cache * llama_memory_hybrid_idx::get_mem_idx() const { return mem_idx.get(); } +void llama_memory_hybrid_idx::commit_qsa_tokens(const llama_ubatch & ubatch) { + if (!mem_idx) { + return; + } + + for (uint32_t i = 0; i < ubatch.n_tokens; ++i) { + qsa_token token = {}; + if (ubatch.token) { + token.pos = { ubatch.pos[i], ubatch.pos[i], ubatch.pos[i], 0 }; + } else { + for (uint32_t ip = 0; ip < token.pos.size(); ++ip) { + token.pos[ip] = ip < ubatch.n_pos ? ubatch.pos[i + ip*ubatch.n_tokens] : ubatch.pos[i]; + } + } + + for (int32_t is = 0; is < ubatch.n_seq_id[i]; ++is) { + qsa_histories[ubatch.seq_id[i][is]].push_back(token); + } + } +} + +void llama_memory_hybrid_idx::set_input_qsa( + ggml_tensor * block_cells, + ggml_tensor * block_pos, + ggml_tensor * block_mask, + ggml_tensor * selected, + const ggml_tensor * kq_mask, + const llama_ubatch * ubatch, + uint32_t ratio, + uint32_t block_topk) const { + GGML_ASSERT(ggml_backend_buffer_is_host(block_cells->buffer)); + GGML_ASSERT(ggml_backend_buffer_is_host(block_pos->buffer)); + GGML_ASSERT(ggml_backend_buffer_is_host(block_mask->buffer)); + GGML_ASSERT(ggml_backend_buffer_is_host(selected->buffer)); + GGML_ASSERT(ggml_backend_buffer_is_host(kq_mask->buffer)); + + const int64_t n_blocks = block_cells->ne[1]; + const int64_t n_tokens = ubatch->n_tokens; + const int64_t n_pos = block_pos->ne[1]; + const int64_t n_kv = selected->ne[0]; + + GGML_ASSERT(block_cells->type == GGML_TYPE_I32); + GGML_ASSERT(block_pos->type == GGML_TYPE_I32); + GGML_ASSERT(block_mask->type == GGML_TYPE_F32); + GGML_ASSERT(selected->type == GGML_TYPE_F32); + GGML_ASSERT(block_cells->ne[0] == ratio && block_cells->ne[2] == n_tokens); + GGML_ASSERT(block_pos->ne[0] == n_blocks && block_pos->ne[2] == n_tokens); + GGML_ASSERT(block_mask->ne[0] == n_blocks && block_mask->ne[1] == n_tokens); + GGML_ASSERT(selected->ne[1] == n_tokens); + GGML_ASSERT(kq_mask->ne[0] == n_kv); + + int32_t * cell_data = (int32_t *) block_cells->data; + int32_t * pos_data = (int32_t *) block_pos->data; + float * mask_data = (float *) block_mask->data; + float * selected_data = (float *) selected->data; + + std::fill(cell_data, cell_data + ggml_nelements(block_cells), 0); + std::fill(pos_data, pos_data + ggml_nelements(block_pos), 0); + std::fill(mask_data, mask_data + ggml_nelements(block_mask), -INFINITY); + std::fill(selected_data, selected_data + ggml_nelements(selected), 0.0f); + + auto mask_visible = [&](int64_t query, uint32_t cell) { + const int64_t index = query*n_kv + cell; + if (kq_mask->type == GGML_TYPE_F16) { + return std::isfinite(ggml_fp16_to_fp32(((const ggml_fp16_t *) kq_mask->data)[index])); + } + return std::isfinite(((const float *) kq_mask->data)[index]); + }; + + for (int64_t iq = 0; iq < n_tokens; ++iq) { + const llama_seq_id seq_id = ubatch->seq_id[iq][0]; + const auto found = qsa_histories.find(seq_id); + if (found == qsa_histories.end()) { + continue; + } + + const auto & cells = get_mem_attn()->get_cells(seq_id); + using pos_key = std::tuple; + std::map> cells_by_pos; + for (uint32_t cell = 0; cell < cells.size() && cell < (uint32_t) n_kv; ++cell) { + if (cells.is_empty(cell) || !cells.seq_has(cell, seq_id)) { + continue; + } + + const auto & ext = cells.ext_get(cell); + cells_by_pos[{ cells.pos_get(cell), ext.y, ext.x }].push_back(cell); + } + + std::map next_cell; + std::vector> visible; + for (const auto & token : found->second) { + const pos_key key = { token.pos[0], token.pos[1], token.pos[2] }; + auto cells = cells_by_pos.find(key); + if (cells == cells_by_pos.end()) { + continue; + } + + size_t & index = next_cell[key]; + if (index >= cells->second.size()) { + continue; + } + + const uint32_t cell = cells->second[index++]; + if (mask_visible(iq, cell)) { + visible.emplace_back(&token, cell); + } + } + + const size_t n_complete = visible.size()/ratio; + const size_t n_write = std::min(n_complete, n_blocks); + std::vector used_cells(n_kv, 0); + for (size_t ib = 0; ib < n_write; ++ib) { + mask_data[iq*n_blocks + ib] = 0.0f; + for (uint32_t ir = 0; ir < ratio; ++ir) { + const uint32_t cell = visible[ib*ratio + ir].second; + cell_data[(iq*n_blocks + ib)*ratio + ir] = cell; + used_cells[cell] = 1; + } + for (int64_t ip = 0; ip < n_pos; ++ip) { + pos_data[(iq*n_pos + ip)*n_blocks + ib] = visible[ib*ratio].first->pos[ip]; + } + } + + uint32_t fallback_cell = 0; + for (size_t ib = n_write; ib < (size_t) n_blocks; ++ib) { + for (uint32_t ir = 0; ir < ratio; ++ir) { + while (fallback_cell < used_cells.size() && used_cells[fallback_cell]) { + ++fallback_cell; + } + GGML_ASSERT(fallback_cell < used_cells.size()); + cell_data[(iq*n_blocks + ib)*ratio + ir] = fallback_cell; + used_cells[fallback_cell++] = 1; + } + } + + const size_t selected_start = n_complete <= block_topk ? 0 : n_complete*ratio; + for (size_t iv = selected_start; iv < visible.size(); ++iv) { + selected_data[iq*n_kv + visible[iv].second] = 1.0f; + } + } +} + // // llama_memory_hybrid_idx_context // @@ -295,7 +598,9 @@ llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context( llama_context * lctx, bool optimize) : llama_memory_hybrid_context(mem, lctx, optimize), - mem(mem) {} + mem(mem), + ctx_idx(mem->get_mem_idx() == nullptr ? nullptr : + mem->get_mem_idx()->init_update(lctx, optimize)) {} llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context( llama_memory_hybrid_idx * mem, @@ -307,7 +612,8 @@ llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context( mem(mem), ns_ubatch(llama_memory_hybrid_idx_ns(sinfos_idx)), ctx_idx(mem->get_mem_idx() == nullptr ? nullptr : - new llama_kv_cache_context(mem->get_mem_idx(), std::move(sinfos_idx), ubatches)) {} + new llama_kv_cache_context(mem->get_mem_idx(), std::move(sinfos_idx), ubatches)), + has_ubatches(true) {} bool llama_memory_hybrid_idx_context::next() { if (ctx_idx) { @@ -326,6 +632,10 @@ bool llama_memory_hybrid_idx_context::apply() { res = res & ctx_idx->apply(); } + if (res && ctx_idx && has_ubatches) { + mem->commit_qsa_tokens(ctx_idx->get_ubatch()); + } + return res; } @@ -340,126 +650,17 @@ uint32_t llama_memory_hybrid_idx_context::get_n_stream() const { } void llama_memory_hybrid_idx_context::set_input_qsa( - ggml_tensor * cell_blk, - ggml_tensor * blk_cells, - ggml_tensor * blk_pos, - ggml_tensor * bias, + ggml_tensor * block_cells, + ggml_tensor * block_pos, + ggml_tensor * block_mask, + ggml_tensor * selected, + const ggml_tensor * kq_mask, const llama_ubatch * ubatch, - uint32_t ratio, - bool blk_bias) const { + uint32_t ratio, + uint32_t block_topk) const { GGML_ASSERT(ratio > 0); GGML_ASSERT(mem != nullptr && mem->get_mem_idx() != nullptr); - GGML_ASSERT(ggml_backend_buffer_is_host(cell_blk->buffer)); - - const int64_t n_kv = cell_blk->ne[0]; - const int64_t n_ns = cell_blk->ne[1]; // streams in this ubatch - const int64_t n_blocks = blk_pos->ne[0]/(4*n_ns); - const int64_t n_tokens = ubatch->n_tokens; - const int64_t r = ratio; - - GGML_ASSERT(n_tokens % n_ns == 0); - const int64_t n_tps = n_tokens/n_ns; // tokens per stream - - int32_t * dst_cell_blk = (int32_t *) cell_blk->data; - int32_t * dst_blk_cells = (int32_t *) blk_cells->data; - int32_t * dst_blk_pos = (int32_t *) blk_pos->data; - float * dst_bias = (float *) bias->data; - - // block b covers [b*ratio, (b+1)*ratio), so its first token is at b*ratio - // all mrope sections carry it: exact for text, approximate for images - for (int64_t sec = 0; sec < 4; ++sec) { - for (int64_t s = 0; s < n_ns; ++s) { - for (int64_t b = 0; b < n_blocks; ++b) { - dst_blk_pos[sec*(n_blocks*n_ns) + s*n_blocks + b] = (int32_t) (b*r); - } - } - } - - // one pass per stream: cell j is a different token in each, so no mapping is shared - std::vector blk_of(n_kv); - std::vector filled(n_blocks); - - for (int64_t s = 0; s < n_ns; ++s) { - // ubatch index s*n_tps belongs to this stream; ask which cells array it uses - const llama_seq_id seq_of_stream = ubatch->seq_id[s*n_tps][0]; - const auto & cells = mem->get_mem_idx()->get_cells(seq_of_stream); - - int32_t * cur_cell_blk = dst_cell_blk + s*n_kv; - int32_t * cur_blk_cells = dst_blk_cells + s*(r*n_blocks); - - // an incomplete block cannot be pooled; the bias below forces those tail cells in - // -1 means no usable block, and block 0 only keeps the gather in range - std::fill(blk_of.begin(), blk_of.end(), -1); - std::fill(filled.begin(), filled.end(), 0); - std::fill(cur_blk_cells, cur_blk_cells + r*n_blocks, 0); - - // a cell no block covers needs its own -inf, which a per-block bias cannot carry - // every cache path keeps the position below the cell window, so this stays false - bool oor = false; - - for (int64_t j = 0; j < n_kv; ++j) { - if (cells.is_empty(j)) { - continue; - } - - const llama_pos p = cells.pos_get(j); - const int64_t b = p/r; - - if (b >= n_blocks) { - oor = true; - continue; - } - - blk_of[j] = (int32_t) b; - cur_blk_cells[b*r + (p%r)] = (int32_t) j; - filled[b]++; - } - - GGML_ASSERT((!blk_bias || !oor) && "qsa: cell position runs past the cell window"); - - // per-block mode keeps an unpooled cell's real block, so the block's own -inf reaches it - // per-cell mode carries that -inf itself and only needs the gather in range - for (int64_t j = 0; j < n_kv; ++j) { - if (blk_of[j] >= 0 && filled[blk_of[j]] < r && !blk_bias) { - blk_of[j] = -1; - } - cur_cell_blk[j] = blk_of[j] < 0 ? 0 : blk_of[j]; - } - - for (int64_t ii = 0; ii < n_tps; ++ii) { - const int64_t i = s*n_tps + ii; - const llama_seq_id seq_id = ubatch->seq_id[i][0]; - const llama_pos q = ubatch->pos[i]; - - // the tail is an incomplete block and is always visible, as in the reference - const llama_pos tail_start = (q + 1)/r*r; - - if (blk_bias) { - // a block sits wholly inside or outside the tail, so one value covers it - // the caller adds the attention mask, which drops empty, foreign and future cells - float * cur_blk_bias = dst_bias + i*n_blocks; - - for (int64_t b = 0; b < n_blocks; ++b) { - // finite, so it can never meet a -inf and produce a nan - cur_blk_bias[b] = b*r >= tail_start ? 1e9f : (filled[b] < r ? -INFINITY : 0.0f); - } - - continue; - } - - float * cur_bias = dst_bias + i*n_kv; - - for (int64_t j = 0; j < n_kv; ++j) { - float v = -INFINITY; - - if (!cells.is_empty(j) && cells.seq_has(j, seq_id) && cells.pos_get(j) <= q) { - // finite, so it can never meet a -inf and produce a nan - v = cells.pos_get(j) >= tail_start ? 1e9f : (blk_of[j] < 0 ? -INFINITY : 0.0f); - } - - cur_bias[j] = v; - } - } - } + mem->set_input_qsa(block_cells, block_pos, block_mask, selected, + kq_mask, ubatch, ratio, block_topk); } diff --git a/src/llama-memory-hybrid-idx.h b/src/llama-memory-hybrid-idx.h index e3472646d0..746ce76ab1 100644 --- a/src/llama-memory-hybrid-idx.h +++ b/src/llama-memory-hybrid-idx.h @@ -2,6 +2,8 @@ #include "llama-memory-hybrid.h" +#include +#include #include #include @@ -75,7 +77,20 @@ public: llama_kv_cache * get_mem_idx() const; // nullptr when the model carries no indexer + void set_input_qsa(ggml_tensor * block_cells, ggml_tensor * block_pos, + ggml_tensor * block_mask, ggml_tensor * selected, + const ggml_tensor * kq_mask, const llama_ubatch * ubatch, + uint32_t ratio, uint32_t block_topk) const; + + void commit_qsa_tokens(const llama_ubatch & ubatch); + private: + struct qsa_token { + std::array pos; + }; + + using qsa_history = std::vector; + // forget seq_id (all of it if seq_id < 0) in every cache at once, so a failed restore cannot leave the caches out of step // seq_id < 0 drops the whole context, as the caches themselves do on a failed restore void state_drop(llama_seq_id seq_id); @@ -85,6 +100,8 @@ private: llama_hparams hparams_idx; const std::unique_ptr mem_idx; + + std::map qsa_histories; }; class llama_memory_hybrid_idx_context : public llama_memory_hybrid_context { @@ -123,26 +140,20 @@ public: // llama_memory_hybrid_idx_context specific API // - // nullptr with no indexer, and for the update context, which builds no sparse graph + // nullptr with no indexer const llama_kv_cache_context * get_idx() const; // streams in the current slot info, the `ns` of get_k/get_v; 1 if unified uint32_t get_n_stream() const; - // block-compressed sparse attention (qwen4exp QSA) over the cells of the indexer cache. - // Blocks cut the position line, not the cell array, so no caller assumes a contiguous layout: - // cell_blk I32 [n_kv, ns] block each cell belongs to - // blk_cells I32 [ratio*n_blocks, ns] cells making up each block - // blk_pos I32 [4*n_blocks*ns] mrope position rows of each block's first token - // bias F32 [n_kv, n_tokens/ns, ns] -inf where invisible, large where always visible - // blk_bias asks for the bias per block instead: [n_blocks, n_tokens/ns, ns] - // the caller then adds the attention mask, the only part of the bias that varies within a block - void set_input_qsa(ggml_tensor * cell_blk, ggml_tensor * blk_cells, ggml_tensor * blk_pos, - ggml_tensor * bias, const llama_ubatch * ubatch, uint32_t ratio, - bool blk_bias) const; + // QSA blocks follow each sequence's token order, not physical cells or scalar positions. + void set_input_qsa(ggml_tensor * block_cells, ggml_tensor * block_pos, + ggml_tensor * block_mask, ggml_tensor * selected, + const ggml_tensor * kq_mask, const llama_ubatch * ubatch, + uint32_t ratio, uint32_t block_topk) const; private: - const llama_memory_hybrid_idx * mem = nullptr; + llama_memory_hybrid_idx * mem = nullptr; // streams per ubatch, read from the slot infos before ctx_idx takes them // declared first, so it is initialised while sinfos_idx is still intact @@ -151,6 +162,8 @@ private: // null unless the model has an indexer and this is a batch or full context const llama_memory_context_ptr ctx_idx; + const bool has_ubatches = false; + // mirrors the base class's ubatch cursor, which is private there size_t i_cur = 0; }; diff --git a/src/models/models.h b/src/models/models.h index af60764c2f..0283500a53 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -2310,22 +2310,12 @@ struct llama_model_qwen4exp : public llama_model_base { int * sections, int il); - // dense self-attention restricted to the cells that top_k names - ggml_tensor * build_attn_qsa( - llm_graph_input_attn_kv * inp, - ggml_tensor * q_cur, - ggml_tensor * k_cur, - ggml_tensor * v_cur, - ggml_tensor * top_k, - float kq_scale, - int il); - // the QSA cache layout inputs do not depend on the layer, only on its compress ratio, // so the layers sharing a ratio share one input set std::map qsa_inps; - // QSA: token indices this layer's queries may attend to, or nullptr for dense - ggml_tensor * build_qsa_top_k( + // QSA mask for this layer, or nullptr for dense attention + ggml_tensor * build_qsa_mask( const llama_memory_hybrid_idx_context * mctx_hyb, ggml_tensor * cur, ggml_tensor * inp_pos, diff --git a/src/models/qwen4exp.cpp b/src/models/qwen4exp.cpp index acfdd5b500..cb18ca0366 100644 --- a/src/models/qwen4exp.cpp +++ b/src/models/qwen4exp.cpp @@ -18,21 +18,29 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) { ml.get_key(LLM_KV_SSM_STATE_SIZE, hparams.ssm_d_state); ml.get_key(LLM_KV_SSM_TIME_STEP_RANK, hparams.ssm_dt_rank); ml.get_key(LLM_KV_SSM_GROUP_COUNT, hparams.ssm_n_group); - GGML_ASSERT(hparams.ssm_d_conv > 0 && hparams.ssm_d_inner > 0 && hparams.ssm_d_state > 0 && - hparams.ssm_dt_rank > 0 && hparams.ssm_n_group > 0); + if (hparams.ssm_d_conv == 0 || hparams.ssm_d_inner == 0 || hparams.ssm_d_state == 0 || + hparams.ssm_dt_rank == 0 || hparams.ssm_n_group == 0 || + hparams.ssm_dt_rank % hparams.ssm_n_group != 0 || + (uint64_t) hparams.ssm_d_state * hparams.ssm_dt_rank != hparams.ssm_d_inner) { + throw std::runtime_error("invalid Qwen4-Exp gated delta net dimensions"); + } // HC; low_rank is qwen4exp-specific, DeepSeek-V4 leaves it absent (full rank) ml.get_key(LLM_KV_HYPER_CONNECTION_COUNT, hparams.dsv4_hc_mult); ml.get_key(LLM_KV_HYPER_CONNECTION_LOW_RANK, hparams.hc_low_rank); - GGML_ASSERT(hparams.dsv4_hc_mult > 0 && hparams.hc_low_rank > 0); + if (hparams.n_embd == 0 || hparams.dsv4_hc_mult <= 1 || hparams.hc_low_rank == 0 || + hparams.dsv4_hc_mult > UINT32_MAX/hparams.n_embd) { + throw std::runtime_error("invalid Qwen4-Exp hyper-connection dimensions"); + } hparams.n_embd_out_impl = hparams.dsv4_hc_mult * hparams.n_embd; ml.get_key(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, hparams.indexer_n_head); ml.get_key(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, hparams.indexer_head_size); ml.get_key(LLM_KV_ATTENTION_INDEXER_TOP_K, hparams.indexer_top_k); - GGML_ASSERT(hparams.indexer_n_head > 0 - && hparams.indexer_head_size > 0 - && hparams.indexer_top_k > 0); + if (hparams.indexer_n_head == 0 || hparams.indexer_head_size == 0 || hparams.indexer_top_k == 0 || + hparams.n_rot_full > hparams.indexer_head_size) { + throw std::runtime_error("invalid Qwen4-Exp sparse-attention dimensions"); + } ml.get_key_or_arr(LLM_KV_ATTENTION_COMPRESS_RATIOS, hparams.dsv4_compress_ratios, hparams.n_layer_all, false); // PLE n-gram hash embeddings; if the key group is absent every field stays zero @@ -44,9 +52,11 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) { if (n_ple > 0) { std::vector ple_layers; ml.get_arr(LLM_KV_PLE_LAYERS, ple_layers); - GGML_ASSERT(n_ple == 1 && "qwen4exp supports only one PLE layer"); + if (n_ple != 1 || ple_layers.size() != n_ple) { + throw std::runtime_error("Qwen4-Exp supports exactly one PLE layer"); + } for (uint32_t il : ple_layers) { - if (il >= hparams.n_layer_all) { + if (il >= hparams.n_layer()) { throw std::runtime_error(format("PLE layer %u is out of range", il)); } hparams.is_ple_impl.set(il); @@ -59,15 +69,30 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) { // optional: files written before this key fall back to the EOS token ml.get_key(LLM_KV_PLE_IMAGE_TOKEN_ID, hparams.ple_image_token_id, false); ml.get_key(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, hparams.n_embd_per_layer); - GGML_ASSERT(hparams.ple_conv_kernel > 0 && hparams.n_embd_per_layer > 0); + if (hparams.ple_conv_kernel == 0 || hparams.n_embd_per_layer == 0) { + throw std::runtime_error("invalid Qwen4-Exp PLE dimensions"); + } - hparams.ple_n_heads = (hparams.ple_ngram_size - 1) * hparams.ple_heads_per_ngram; - hparams.ple_head_dim = hparams.n_embd_per_layer; if (hparams.ple_ngram_size < 2 || hparams.ple_ngram_size > LLAMA_MAX_PLE_NGRAM) { throw std::runtime_error(format("PLE n-gram size %u is out of range", hparams.ple_ngram_size)); } - if (hparams.ple_n_heads == 0 || hparams.ple_n_heads > LLAMA_MAX_PLE_HEADS) { - throw std::runtime_error(format("PLE head count %u is out of range", hparams.ple_n_heads)); + + const uint64_t ple_n_heads = (uint64_t) (hparams.ple_ngram_size - 1) * hparams.ple_heads_per_ngram; + hparams.ple_head_dim = hparams.n_embd_per_layer; + if (ple_n_heads == 0 || ple_n_heads > LLAMA_MAX_PLE_HEADS) { + throw std::runtime_error(format("PLE head count %" PRIu64 " is out of range", ple_n_heads)); + } + hparams.ple_n_heads = (uint32_t) ple_n_heads; + + uint32_t n_multipliers = 0; + uint32_t n_offsets = 0; + uint32_t n_vocab_sizes = 0; + ml.get_arr_n(LLM_KV_PLE_LAYER_MULTIPLIERS, n_multipliers); + ml.get_arr_n(LLM_KV_PLE_HEAD_OFFSETS, n_offsets); + ml.get_arr_n(LLM_KV_PLE_HEAD_VOCAB_SIZES, n_vocab_sizes); + if (n_multipliers != hparams.ple_ngram_size || + n_offsets != hparams.ple_n_heads || n_vocab_sizes != hparams.ple_n_heads) { + throw std::runtime_error("invalid Qwen4-Exp PLE metadata lengths"); } ml.get_arr(LLM_KV_PLE_LAYER_MULTIPLIERS, hparams.ple_layer_multipliers); @@ -93,12 +118,28 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) { if (!ml.get_key_or_arr(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, hparams.n_layer_all, false)) { uint32_t full_attn_interval = 4; ml.get_key(LLM_KV_FULL_ATTENTION_INTERVAL, full_attn_interval, false); - GGML_ASSERT(full_attn_interval > 0); + if (full_attn_interval == 0) { + throw std::runtime_error("invalid Qwen4-Exp full-attention interval"); + } for (uint32_t i = 0; i < hparams.n_layer_all; ++i) { hparams.is_recr_impl[i] = (i < hparams.n_layer()) && ((i + 1) % full_attn_interval != 0); } } + for (uint32_t il = 0; il < hparams.n_layer(); ++il) { + const uint32_t ratio = hparams.dsv4_compress_ratios[il]; + if (hparams.is_recr(il)) { + if (ratio != 0) { + throw std::runtime_error(format("Qwen4-Exp recurrent layer %u has a QSA compression ratio", il)); + } + } else if (ratio == 0 || hparams.indexer_top_k % ratio != 0) { + throw std::runtime_error(format("invalid Qwen4-Exp QSA compression ratio %u at layer %u", ratio, il)); + } + if (hparams.is_ple(il) && !hparams.is_recr(il)) { + throw std::runtime_error(format("Qwen4-Exp PLE layer %u is not recurrent", il)); + } + } + switch (hparams.n_layer()) { case 48: type = LLM_TYPE_A3B; break; default: type = LLM_TYPE_UNKNOWN; @@ -127,8 +168,15 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) { // flat [ple_head_dim, n_rows] gather target; n_rows is padded, so read it back if (hparams.ple_n_heads > 0) { const std::string ple_name = tn(LLM_TENSOR_PER_LAYER_TOKEN_EMBD, "weight").str(); - const auto & ple_w = ml.require_weight(ple_name.c_str()); - const int64_t ple_rows = ple_w.tensor->ne[1]; + const auto * ple_w = ml.get_weight(ple_name.c_str()); + int64_t ple_rows = 0; + for (uint32_t h = 0; h < hparams.ple_n_heads; ++h) { + ple_rows = std::max(ple_rows, + (int64_t) hparams.ple_head_offsets[h] + hparams.ple_head_vocab_sizes[h]); + } + if (ple_w) { + ple_rows = ple_w->tensor->ne[1]; + } // sanity check for (uint32_t h = 0; h < hparams.ple_n_heads; ++h) { @@ -190,8 +238,9 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) { } if (hparams.is_ple(il)) { - layer.ple_key = create_tensor(tn(LLM_TENSOR_PLE_KEY, "weight", il), { n_embd, hc_dim }, 0); - layer.ple_value = create_tensor(tn(LLM_TENSOR_PLE_VALUE, "weight", il), { n_embd, n_embd }, 0); + const int64_t ple_dim = (int64_t) hparams.ple_head_dim * hparams.ple_n_heads; + layer.ple_key = create_tensor(tn(LLM_TENSOR_PLE_KEY, "weight", il), { ple_dim, hc_dim }, 0); + layer.ple_value = create_tensor(tn(LLM_TENSOR_PLE_VALUE, "weight", il), { ple_dim, n_embd }, 0); layer.ple_norm_key = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY, "weight", il), { hc_dim }, 0); layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { hc_dim }, 0); layer.ple_norm_conv = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV, "weight", il), { hc_dim }, 0); @@ -416,13 +465,17 @@ ggml_tensor * llama_model_qwen4exp::graph::build_norm_gated( // one mean-pooled indexer key scores each block; set_input resolves the cache layout class llama_model_qwen4exp::llm_graph_input_qsa : public llm_graph_input_i { public: - llm_graph_input_qsa(const llama_memory_hybrid_idx_context * mctx, uint32_t ratio, bool blk_bias) : - mctx(mctx), ratio(ratio), blk_bias(blk_bias) {} + llm_graph_input_qsa( + const llama_memory_hybrid_idx_context * mctx, + ggml_tensor * kq_mask, + uint32_t ratio, + uint32_t block_topk) : + mctx(mctx), kq_mask(kq_mask), ratio(ratio), block_topk(block_topk) {} virtual ~llm_graph_input_qsa() = default; void set_input(const llama_ubatch * ubatch) override { - mctx->get_idx()->set_input_k_idxs(k_idxs, ubatch); - mctx->set_input_qsa(cell_blk, blk_cells, blk_pos, bias, ubatch, ratio, blk_bias); + mctx->set_input_qsa(block_cells, block_pos, block_mask, selected, + kq_mask, ubatch, ratio, block_topk); } bool can_reuse(const llm_graph_params & params) override { @@ -434,39 +487,33 @@ public: } const int64_t n_kv = idx->get_n_kv(); - const int64_t n_stream = mctx->get_n_stream(); - const int64_t n_blocks = (n_kv + ratio - 1)/ratio; + const int64_t n_blocks = n_kv/ratio; bool res = true; - res &= params.ubatch.n_tokens % n_stream == 0; - - res &= k_idxs->ne[0] == params.ubatch.n_tokens; - res &= cell_blk->ne[0] == n_kv; - res &= cell_blk->ne[1] == n_stream; - res &= blk_cells->ne[0] == (int64_t) ratio*n_blocks; - res &= blk_pos->ne[0] == 4*n_blocks*n_stream; - res &= bias->ne[0] == (blk_bias ? n_blocks : n_kv); - res &= bias->ne[1] == params.ubatch.n_tokens/n_stream; + res &= n_kv > (int64_t) block_topk*ratio + ratio - 1; + res &= block_cells->ne[1] == n_blocks; + res &= block_cells->ne[2] == params.ubatch.n_tokens; + res &= block_pos->ne[2] == params.ubatch.n_tokens; + res &= block_mask->ne[1] == params.ubatch.n_tokens; + res &= selected->ne[0] == n_kv; + res &= selected->ne[1] == params.ubatch.n_tokens; return res; } - // per stream: a cell index names a different token in each stream - ggml_tensor * k_idxs = nullptr; // I32 [n_tokens] - ggml_tensor * cell_blk = nullptr; // I32 [n_kv, n_stream] - ggml_tensor * blk_cells = nullptr; // I32 [ratio*n_blocks, n_stream] - ggml_tensor * blk_pos = nullptr; // I32 [4*n_blocks*n_stream] - ggml_tensor * bias = nullptr; // F32 [n_blocks or n_kv, n_tokens/n_stream, n_stream] + ggml_tensor * block_cells = nullptr; + ggml_tensor * block_pos = nullptr; + ggml_tensor * block_mask = nullptr; + ggml_tensor * selected = nullptr; const llama_memory_hybrid_idx_context * mctx; + ggml_tensor * kq_mask; const uint32_t ratio; - - // the per-cell half of the bias is the attention mask, so only the per-block half is uploaded - const bool blk_bias; + const uint32_t block_topk; }; -ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k( +ggml_tensor * llama_model_qwen4exp::graph::build_qsa_mask( const llama_memory_hybrid_idx_context * mctx_hyb, ggml_tensor * cur, ggml_tensor * inp_pos, @@ -482,21 +529,13 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k( GGML_ASSERT(r > 0); - const int64_t n_blocks = (n_kv + r - 1)/r; + const int64_t n_blocks = n_kv/r; + const int64_t block_topk = hparams.indexer_top_k/r; - // build_attn_qsa and the KQ mask need the tokens to divide evenly across the streams const int64_t n_stream = mctx_hyb->get_n_stream(); GGML_ASSERT(n_tokens % n_stream == 0); const int64_t n_tps = n_tokens/n_stream; - // only the "which block is visible" half of the bias varies per block - // the rest is the visible/not test the attention mask already carries, so upload the per-block half only: 1/ratio of the cells - // alibi writes distances instead of a mask and non-causal keeps future cells, so both opt out - // the mask also holds an mrope rule for the query's own position, but only 2d image positions can differ there - const bool blk_bias = kq_mask != nullptr && - kq_mask->ne[0] == n_kv && kq_mask->ne[1] == n_tps && kq_mask->ne[3] == n_stream && - cparams.causal_attn && !hparams.use_alibi; - // nothing above depends on the layer, so the layers sharing a ratio share one input set llm_graph_input_qsa * inp = nullptr; @@ -504,59 +543,29 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k( if (it != qsa_inps.end()) { inp = it->second; } else { - auto qsa = std::make_unique(mctx_hyb, (uint32_t) r, blk_bias); + auto qsa = std::make_unique( + mctx_hyb, kq_mask, (uint32_t) r, (uint32_t) block_topk); - qsa->k_idxs = mctx_idx->build_input_k_idxs(ctx0, ubatch); - qsa->cell_blk = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_kv, n_stream); - qsa->blk_cells = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, r*n_blocks, n_stream); - qsa->blk_pos = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, 4*n_blocks*n_stream); - qsa->bias = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, blk_bias ? n_blocks : n_kv, n_tps, n_stream); + qsa->block_cells = ggml_new_tensor_3d(ctx0, GGML_TYPE_I32, r, n_blocks, n_tokens); + qsa->block_pos = ggml_new_tensor_3d(ctx0, GGML_TYPE_I32, n_blocks, 4, n_tokens); + qsa->block_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_blocks, n_tokens); + qsa->selected = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_kv, n_tokens); - ggml_set_input(qsa->cell_blk); - ggml_set_input(qsa->blk_cells); - ggml_set_input(qsa->blk_pos); - ggml_set_input(qsa->bias); + ggml_set_input(qsa->block_cells); + ggml_set_input(qsa->block_pos); + ggml_set_input(qsa->block_mask); + ggml_set_input(qsa->selected); inp = qsa.get(); res->add_input(std::move(qsa)); qsa_inps.emplace((uint32_t) r, inp); } - // cached indexer keys are raw: pooling precedes norm and rotation, so apply neither - ggml_tensor * k_raw = build_lora_mm(model.layers[il].index_k_proj, cur); - k_raw = ggml_reshape_3d(ctx0, k_raw, idx_dim, 1, n_tokens); - cb(k_raw, "indexer_k_raw", il); - - ggml_build_forward_expand(gf, mctx_idx->cpy_k(ctx0, k_raw, inp->k_idxs, il)); + kq_mask = inp->kq_mask; // one key head, so rows are contiguous. get_k gives [idx_dim, n_head_kv, n_kv, n_stream]. ggml_tensor * k_all = mctx_idx->get_k(ctx0, il); k_all = ggml_view_3d(ctx0, k_all, idx_dim, n_kv, n_stream, k_all->nb[2], k_all->nb[3], 0); - - // gathers per stream: blk_cells row s indexes stream s's own cells - ggml_tensor * members = ggml_get_rows(ctx0, k_all, inp->blk_cells); - members = ggml_reshape_4d(ctx0, members, idx_dim, r, n_blocks, n_stream); - - // mean over the block members; r is small, so summing slices beats a transpose plus sum_rows - ggml_tensor * pooled = nullptr; - for (int64_t i = 0; i < r; ++i) { - ggml_tensor * slice = ggml_cont(ctx0, - ggml_view_3d(ctx0, members, idx_dim, n_blocks, n_stream, - members->nb[2], members->nb[3], i*members->nb[1])); - pooled = pooled ? ggml_add(ctx0, pooled, slice) : slice; - } - pooled = ggml_scale(ctx0, pooled, 1.0f/(float) r); - cb(pooled, "indexer_k_pooled", il); - - // rope wants [n_dims, n_head, n_tokens]: lay every stream's blocks flat, split after. - pooled = ggml_reshape_3d(ctx0, pooled, idx_dim, 1, n_blocks*n_stream); - pooled = build_norm(pooled, model.layers[il].index_k_norm, nullptr, LLM_NORM_RMS, il); - pooled = ggml_rope_multi(ctx0, pooled, inp->blk_pos, nullptr, - n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, - ext_factor, attn_factor, beta_fast, beta_slow); - pooled = ggml_reshape_3d(ctx0, pooled, idx_dim, n_blocks, n_stream); - cb(pooled, "indexer_k", il); - ggml_tensor * q = build_lora_mm(model.layers[il].index_q_proj, cur); q = ggml_reshape_3d(ctx0, q, idx_dim, n_idx_h, n_tokens); q = build_norm(q, model.layers[il].index_q_norm, nullptr, LLM_NORM_RMS, il); @@ -565,128 +574,86 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k( ext_factor, attn_factor, beta_fast, beta_slow); cb(q, "indexer_q", il); - // rectify each head dot product before the sum, as in the DeepSeek lightning indexer - // mul_mat matches ne[2], so the queries of stream s only meet the blocks of stream s - ggml_tensor * score = ggml_mul_mat(ctx0, pooled, - ggml_reshape_3d(ctx0, ggml_cont(ctx0, q), idx_dim, n_idx_h*n_tps, n_stream)); - score = ggml_reshape_4d(ctx0, score, n_blocks, n_idx_h, n_tps, n_stream); - score = ggml_relu(ctx0, score); - score = ggml_cont(ctx0, ggml_permute(ctx0, score, 1, 0, 2, 3)); - score = ggml_sum_rows(ctx0, score); - score = ggml_reshape_3d(ctx0, score, n_blocks, n_tps, n_stream); - cb(score, "indexer_score", il); + const ggml_type activation_type = mctx_idx->type_k(); + std::vector selected_streams; + selected_streams.reserve(n_stream); - // one value per block, so it is cheaper to bias here than after the cells are expanded - if (blk_bias) { - score = ggml_add(ctx0, score, inp->bias); + for (int64_t is = 0; is < n_stream; ++is) { + ggml_tensor * cache = ggml_view_2d(ctx0, k_all, idx_dim, n_kv, + k_all->nb[1], is*k_all->nb[2]); + ggml_tensor * block_cells = ggml_view_3d(ctx0, inp->block_cells, r, n_blocks, n_tps, + inp->block_cells->nb[1], inp->block_cells->nb[2], is*n_tps*inp->block_cells->nb[2]); + ggml_tensor * block_keys = ggml_get_rows(ctx0, cache, + ggml_reshape_1d(ctx0, block_cells, r*n_blocks*n_tps)); + block_keys = ggml_reshape_4d(ctx0, block_keys, idx_dim, r, n_blocks, n_tps); + block_keys = ggml_cont(ctx0, ggml_transpose(ctx0, block_keys)); + block_keys = ggml_mean(ctx0, block_keys); + block_keys = ggml_cont(ctx0, ggml_transpose(ctx0, block_keys)); + block_keys = ggml_reshape_3d(ctx0, block_keys, idx_dim, 1, n_blocks*n_tps); + + if (block_keys->type != activation_type) { + block_keys = ggml_cast(ctx0, block_keys, activation_type); + block_keys = ggml_cast(ctx0, block_keys, GGML_TYPE_F32); + } + block_keys = build_norm(block_keys, model.layers[il].index_k_norm, nullptr, LLM_NORM_RMS, il); + + ggml_tensor * block_pos = ggml_view_3d(ctx0, inp->block_pos, n_blocks, 4, n_tps, + inp->block_pos->nb[1], inp->block_pos->nb[2], is*n_tps*inp->block_pos->nb[2]); + block_keys = ggml_rope_multi(ctx0, block_keys, + ggml_reshape_1d(ctx0, block_pos, n_blocks*4*n_tps), nullptr, + n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(block_keys, "indexer_k", il); + + block_keys = ggml_reshape_4d(ctx0, block_keys, idx_dim, n_blocks, 1, n_tps); + ggml_tensor * query = ggml_view_3d(ctx0, q, idx_dim, n_idx_h, n_tps, + q->nb[1], q->nb[2], is*n_tps*q->nb[2]); + query = ggml_reshape_4d(ctx0, query, idx_dim, n_idx_h, 1, n_tps); + + ggml_tensor * scores = ggml_mul_mat(ctx0, block_keys, query); + ggml_mul_mat_set_prec(scores, GGML_PREC_F32); + scores = ggml_relu(ctx0, scores); + scores = ggml_sum_rows(ctx0, ggml_cont(ctx0, ggml_permute(ctx0, scores, 1, 0, 2, 3))); + scores = ggml_scale(ctx0, ggml_reshape_2d(ctx0, scores, n_blocks, n_tps), + 1.0f/sqrtf((float) idx_dim)); + + ggml_tensor * block_mask = ggml_view_2d(ctx0, inp->block_mask, n_blocks, n_tps, + inp->block_mask->nb[1], is*n_tps*inp->block_mask->nb[1]); + scores = ggml_add(ctx0, scores, block_mask); + cb(scores, "indexer_score", il); + + ggml_tensor * top_blocks = ggml_top_k(ctx0, scores, block_topk); + ggml_tensor * top_cells = ggml_get_rows(ctx0, block_cells, top_blocks); + top_cells = ggml_reshape_2d(ctx0, top_cells, r*block_topk, n_tps); + + ggml_tensor * base_selected = ggml_view_2d(ctx0, inp->selected, n_kv, n_tps, + inp->selected->nb[1], is*n_tps*inp->selected->nb[1]); + base_selected = ggml_reshape_3d(ctx0, base_selected, 1, n_kv, n_tps); + ggml_tensor * selected_top = ggml_fill(ctx0, base_selected, 0.0f); + ggml_tensor * ones = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, 1, r*block_topk, n_tps); + ones = ggml_fill(ctx0, ones, 1.0f); + selected_top = ggml_set_rows(ctx0, selected_top, ones, top_cells); + ggml_tensor * selected_stream = ggml_clamp( + ctx0, ggml_add(ctx0, base_selected, selected_top), 0.0f, 1.0f); + selected_streams.push_back(ggml_reshape_2d(ctx0, selected_stream, n_kv, n_tps)); } - // every token of a block gets the block score; the budget is whole blocks, so top-k cuts on a block boundary - ggml_tensor * expanded = ggml_get_rows(ctx0, - ggml_cont(ctx0, ggml_permute(ctx0, score, 1, 0, 2, 3)), inp->cell_blk); - expanded = ggml_cont(ctx0, ggml_permute(ctx0, expanded, 1, 0, 2, 3)); - - if (blk_bias) { - // flash attention keeps the mask in f16; the scores are f32 - ggml_tensor * mask = kq_mask->type == GGML_TYPE_F32 ? kq_mask : ggml_cast(ctx0, kq_mask, GGML_TYPE_F32); - expanded = ggml_add(ctx0, expanded, ggml_reshape_3d(ctx0, mask, n_kv, n_tps, n_stream)); - } else { - expanded = ggml_add(ctx0, expanded, inp->bias); + ggml_tensor * selected = selected_streams[0]; + for (int64_t is = 1; is < n_stream; ++is) { + selected = ggml_concat(ctx0, selected, selected_streams[is], 1); } - cb(expanded, "indexer_score_tokens", il); + selected = ggml_scale_bias(ctx0, selected, 1e30f, -1e30f); - // the reference returns indexer_top_k + compress_ratio - 1: whole blocks plus the tail - const int64_t width = std::min(n_kv, (int64_t) hparams.indexer_top_k + r - 1); - - ggml_tensor * top_k = ggml_cont(ctx0, ggml_top_k(ctx0, expanded, width)); - - // build_attn_qsa reads [n_top_k, n_batch, 1, n_stream], matching the KQ mask. - top_k = ggml_reshape_4d(ctx0, top_k, width, n_tps, 1, n_stream); - cb(top_k, "indexer_top_k", il); - - return top_k; -} - -// Dense GQA self-attention restricted to the cells that top_k names. -// The mask build below copies the MLA sparse path in llm_graph_context::build_attn. -ggml_tensor * llama_model_qwen4exp::graph::build_attn_qsa( - llm_graph_input_attn_kv * inp, - ggml_tensor * q_cur, - ggml_tensor * k_cur, - ggml_tensor * v_cur, - ggml_tensor * top_k, - float kq_scale, - int il) { - // rotate q/k/v before they reach a quantized cache, as the dense path does. the indexer - // has already scored with its own query in build_qsa_top_k, so top_k is unaffected. - if (inp->self_k_rot) { - q_cur = llama_mul_mat_hadamard(ctx0, q_cur, inp->self_k_rot); - k_cur = llama_mul_mat_hadamard(ctx0, k_cur, inp->self_k_rot); + ggml_tensor * base_mask = ggml_reshape_2d(ctx0, kq_mask, n_kv, n_tokens); + if (base_mask->type != GGML_TYPE_F32) { + base_mask = ggml_cast(ctx0, base_mask, GGML_TYPE_F32); } - - if (inp->self_v_rot) { - v_cur = llama_mul_mat_hadamard(ctx0, v_cur, inp->self_v_rot); + ggml_tensor * mask = ggml_add(ctx0, base_mask, selected); + if (cparams.flash_attn) { + mask = ggml_cast(ctx0, mask, GGML_TYPE_F16); } - - // these nodes are added to the graph together so that they are not reordered - // by doing so, the number of splits in the graph is reduced - // expand k later to enable rope fusion which directly writes into k-v cache - ggml_build_forward_expand(gf, q_cur); - ggml_build_forward_expand(gf, v_cur); - ggml_build_forward_expand(gf, k_cur); - - const auto * mctx_cur = inp->mctx; - - // store to KV cache - { - const auto & k_idxs = inp->get_k_idxs(); - const auto & v_idxs = inp->get_v_idxs(); - - ggml_build_forward_expand(gf, mctx_cur->cpy_k(ctx0, k_cur, k_idxs, il)); - ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, v_idxs, il)); - } - - ggml_tensor * kq_mask = inp->get_kq_mask(); - - // prepare new kq mask - starts filled with -INFINITY - ggml_tensor * kq_mask_all = ggml_fill(ctx0, kq_mask, -INFINITY); - - // reshape KQ mask into tensor with rows of size 1: - // [n_kv, n_batch, 1, n_stream] -> [1, n_kv, n_batch, n_stream] - kq_mask_all = ggml_view_4d(ctx0, kq_mask_all, 1, kq_mask_all->ne[0], kq_mask_all->ne[1], kq_mask_all->ne[3], kq_mask_all->nb[0], kq_mask_all->nb[1], kq_mask_all->nb[2], 0); - - // reshape top_k indices: [n_top_k, n_batch, 1, n_stream] -> [n_top_k, n_batch, n_stream, 1] - ggml_tensor * top_k_3d = ggml_view_4d(ctx0, top_k, top_k->ne[0], top_k->ne[1], top_k->ne[3], 1, top_k->nb[1], top_k->nb[2], top_k->ne[3]*top_k->nb[3], 0); - - // prepare zero-filled tensor with rows of size 1: [1, n_top_k, n_batch, n_stream] - // this will be our source of zero values for unmasking top k mask elements - ggml_tensor * zeros = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, 1, top_k_3d->ne[0], top_k_3d->ne[1], top_k_3d->ne[2]); - zeros = ggml_fill(ctx0, zeros, 0.0f); - - // modify KQ mask by unmasking elements that are in top_k indices - // ggml_set_rows([1, n_kv, n_batch, n_stream], [1, n_top_k, n_batch, n_stream], [n_top_k, n_batch, n_stream, 1]) - ggml_tensor * kq_mask_top_k = ggml_set_rows(ctx0, kq_mask_all, zeros, top_k_3d); - - // reshape to restore the original shape of KQ mask: - // [1, n_kv, n_batch, n_stream] -> [n_kv, n_batch, 1, n_stream] - kq_mask_top_k = ggml_view_4d(ctx0, kq_mask_top_k, kq_mask_top_k->ne[1], kq_mask_top_k->ne[2], 1, kq_mask_top_k->ne[3], kq_mask_top_k->nb[2], kq_mask_top_k->nb[3], kq_mask_top_k->nb[3], 0); - - // combine with the original kq mask - kq_mask_top_k = ggml_add(ctx0, kq_mask_top_k, kq_mask); - - ggml_tensor * q = q_cur; - ggml_tensor * k = mctx_cur->get_k(ctx0, il); - ggml_tensor * v = mctx_cur->get_v(ctx0, il); - - ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, kq_scale, il); - cb(cur, "kqv_out", il); - - // the rotation is its own inverse, so undo it on the value side of the output - if (inp->self_v_rot) { - cur = llama_mul_mat_hadamard(ctx0, cur, inp->self_v_rot); - } - - return cur; + cb(mask, "qsa_mask", il); + return mask; } ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn( @@ -699,10 +666,23 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn( const int64_t n_embd_head = hparams.n_embd_head_v(); GGML_ASSERT(n_embd_head == hparams.n_embd_head_k()); - // indexer reads the same block input as q/k/v; no cache or no ratio means dense - const bool qsa = mctx_hyb->get_idx() != nullptr && hparams.dsv4_compress_ratios[il] > 0; + const llama_kv_cache_context * mctx_idx = mctx_hyb->get_idx(); + if (mctx_idx) { + ggml_tensor * index_k = build_lora_mm(model.layers[il].index_k_proj, cur); + index_k = ggml_reshape_3d(ctx0, index_k, hparams.indexer_head_size, 1, n_tokens); + cb(index_k, "indexer_k_raw", il); + ggml_build_forward_expand(gf, mctx_idx->cpy_k(ctx0, index_k, inp->get_k_idxs(), il)); + } - ggml_tensor * top_k = qsa ? build_qsa_top_k(mctx_hyb, cur, inp_pos, inp->get_kq_mask(), sections, il) : nullptr; + const int64_t ratio = hparams.dsv4_compress_ratios[il]; + const bool qsa = mctx_idx && ratio > 0 && + mctx_idx->get_n_kv() > (int64_t) hparams.indexer_top_k + ratio - 1; + if (qsa) { + inp->self_kq_mask_cnv = build_qsa_mask( + mctx_hyb, cur, inp_pos, inp->self_kq_mask, sections, il); + } else { + inp->self_kq_mask_cnv = inp->self_kq_mask; + } // Qwen3Next uses a single Q projection that outputs query + gate ggml_tensor * Qcur_full = build_lora_mm(model.layers[il].wq, cur, model.layers[il].wq_s); // [ (n_embd_head * 2) * n_head, n_tokens ] @@ -754,13 +734,9 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn( const float kq_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale; - if (top_k) { - cur = build_attn_qsa(inp, Qcur, Kcur, Vcur, top_k, kq_scale, il); - } else { - cur = build_attn(inp, - nullptr, nullptr, nullptr, - Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); - } + cur = build_attn(inp, + nullptr, nullptr, nullptr, + Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il); cb(cur, "attn_pregate", il); ggml_tensor * gate_sigmoid = ggml_sigmoid(ctx0, gate); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index fe3d14ffc5..7067f5f546 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -220,6 +220,16 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS) FIXTURES_REQUIRED generate-models ) + llama_test( + test-recurrent-state-rollback + NAME test-recurrent-state-rollback-qwen4exp + LABEL main + ARGS -m "${MODEL_DIR}/qwen4exp-moe.gguf" + ) + set_tests_properties(test-recurrent-state-rollback-qwen4exp PROPERTIES + FIXTURES_REQUIRED generate-models + ) + llama_test( test-recurrent-state-rollback NAME test-recurrent-state-rollback-nemotron-h diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index df0b9aa666..9bb6becaed 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -9358,6 +9358,7 @@ static std::vector> make_test_cases_eval() { // test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 512, 262144, 9216, {1, 1}, {1, 1})); // test large experts*tokens + test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 512, 10, false, 64, 512, 64)); for (bool b : {false, true}) { test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 16, 16, b, 32, 1024, 16)); test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 2, 2, b, 32, 8192, 64)); diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp index 35a3286e4a..f84bca8a59 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -79,7 +79,7 @@ static std::vector get_tokens(const uint32_t n_tokens, const uint32 return ret; } -static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { +static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe, const bool qwen_ple = false) { gguf_context_ptr ret(gguf_init_empty()); llama_model_saver ms(arch, ret.get()); const uint32_t n_ctx = 256; @@ -252,8 +252,23 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) { if (arch == LLM_ARCH_QWEN4EXP) { ms.add_kv(LLM_KV_HYPER_CONNECTION_COUNT, uint32_t(4)); ms.add_kv(LLM_KV_HYPER_CONNECTION_LOW_RANK, uint32_t(8)); - // without this the QSA layers fall back to dense and go uncovered - ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, std::vector(n_layer, 4)); + std::vector ratios(n_layer, 0); + for (uint32_t il = 1; il < n_layer; il += 2) { + ratios[il] = 4; + } + ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, ratios); + + if (qwen_ple) { + ms.add_kv(LLM_KV_PLE_LAYERS, std::vector({0})); + ms.add_kv(LLM_KV_PLE_NGRAM_SIZE, uint32_t(2)); + ms.add_kv(LLM_KV_PLE_HEADS_PER_NGRAM, uint32_t(1)); + ms.add_kv(LLM_KV_PLE_CONV_KERNEL, uint32_t(2)); + ms.add_kv(LLM_KV_PLE_EOS_TOKEN_ID, uint32_t(1)); + ms.add_kv(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, uint32_t(64)); + ms.add_kv(LLM_KV_PLE_LAYER_MULTIPLIERS, std::vector({1, 3})); + ms.add_kv(LLM_KV_PLE_HEAD_OFFSETS, std::vector({0})); + ms.add_kv(LLM_KV_PLE_HEAD_VOCAB_SIZES, std::vector({16})); + } } // minimax-m3 keeps one indexer head per GQA head; the rest use a fixed 64 to match the fused @@ -347,7 +362,8 @@ static bool silent_model_load_progress(float /*progress*/, void * /*user_data*/) static std::pair get_model_and_ctx( struct gguf_context * gguf_ctx, FILE * file, const size_t seed, const std::vector & devs, - const llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER, bool encode = false) { + const llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER, bool encode = false, + ggml_backend_sched_eval_callback cb_eval = nullptr, void * cb_eval_user_data = nullptr) { GGML_ASSERT((gguf_ctx == nullptr) != (file == nullptr)); llama_model_params model_params = llama_model_default_params(); model_params.progress_callback = silent_model_load_progress; @@ -360,6 +376,8 @@ static std::pair get_model_and_ctx( ctx_params.n_ctx = 0; ctx_params.n_threads = 4; ctx_params.n_threads_batch = 4; + ctx_params.cb_eval = cb_eval; + ctx_params.cb_eval_user_data = cb_eval_user_data; if (!encode) { ctx_params.n_ubatch = 64; } @@ -412,6 +430,54 @@ static std::vector get_logits( return ret; } +struct qwen4_qsa_mask_check { + int64_t n_seen = 0; + int64_t n_tokens_seen = 0; + bool ok = true; +}; + +static bool check_qwen4_qsa_mask(ggml_tensor * tensor, bool ask, void * user_data) { + if (strncmp(tensor->name, "qsa_mask", 8) != 0) { + return false; + } + if (ask) { + return true; + } + + auto * check = (qwen4_qsa_mask_check *) user_data; + const int64_t n_kv = tensor->ne[0]; + const int64_t n_tokens = tensor->ne[1]; + std::vector mask(ggml_nelements(tensor)); + if (tensor->type == GGML_TYPE_F32) { + ggml_backend_tensor_get(tensor, mask.data(), 0, ggml_nbytes(tensor)); + } else { + GGML_ASSERT(tensor->type == GGML_TYPE_F16); + std::vector mask_f16(ggml_nelements(tensor)); + ggml_backend_tensor_get(tensor, mask_f16.data(), 0, ggml_nbytes(tensor)); + for (size_t i = 0; i < mask.size(); ++i) { + mask[i] = ggml_fp16_to_fp32(mask_f16[i]); + } + } + + for (int64_t it = 0; it < n_tokens; ++it) { + const int64_t n_visible = check->n_tokens_seen + it + 1; + const int64_t n_complete = n_visible/4; + const int64_t expected = n_complete <= 2 ? n_visible : 8 + n_visible%4; + int64_t actual = 0; + for (int64_t ikv = 0; ikv < n_kv; ++ikv) { + actual += mask[it*n_kv + ikv] > -1e20f; + } + if (actual != expected) { + fprintf(stderr, "Qwen4 QSA mask row %lld selects %lld tokens, expected %lld\n", + (long long) (check->n_tokens_seen + it), (long long) actual, (long long) expected); + } + check->ok = check->ok && actual == expected; + } + check->n_tokens_seen += n_tokens; + check->n_seen++; + return true; +} + static bool moe_mandatory(const llm_arch arch) { switch (arch) { case LLM_ARCH_LLAMA4: @@ -686,6 +752,32 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg if (arch == LLM_ARCH_BAILINGMOE3) { GGML_ASSERT(gguf_remove_key(gguf_ctx.get(), "bailingmoe3.kda.safe_gate") >= 0); } + if (arch == LLM_ARCH_QWEN4EXP) { + qwen4_qsa_mask_check check; + auto model_and_ctx = get_model_and_ctx( + gguf_ctx.get(), nullptr, seed, {}, LLAMA_SPLIT_MODE_LAYER, false, + check_qwen4_qsa_mask, &check); + get_logits(model_and_ctx.first.get(), model_and_ctx.second.get(), tokens); + GGML_ASSERT(check.ok && check.n_seen > 0); + + gguf_context_ptr gguf_ctx_ple = get_gguf_ctx(arch, moe, true); + auto model_and_ctx_ple = get_model_and_ctx(gguf_ctx_ple.get(), nullptr, seed, {}); + const std::vector logits_ple = get_logits( + model_and_ctx_ple.first.get(), model_and_ctx_ple.second.get(), tokens); + + FILE * file_ple = tmpfile(); + GGML_ASSERT(file_ple); + llama_model_saver saver_ple(model_and_ctx_ple.first.get()); + saver_ple.add_kv_from_model(); + saver_ple.add_tensors_from_model(); + saver_ple.save(file_ple); + rewind(file_ple); + + auto model_and_ctx_ple_saved = get_model_and_ctx(nullptr, file_ple, seed, {}); + const std::vector logits_ple_saved = get_logits( + model_and_ctx_ple_saved.first.get(), model_and_ctx_ple_saved.second.get(), tokens); + GGML_ASSERT(logits_ple == logits_ple_saved); + } std::pair model_and_ctx_cpu; std::vector logits_cpu; for (device_config & dc : dev_configs) {