Publish versioned FPX optimization and fix scripts 20261021-01

Add pinned source patch, verified application script, runtime environment preset and CK import support. Preserve existing runtime release artifacts; update only mutable README index. No model weights or full source snapshot.

AiTable-ROCmFPX-Optimizations-Fixes-20261021-01.zip

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:0b11b7f5b1c95f69524efcad44d4f1c3f1b9102d403380b1ca7dea0e5102c087
+size 44071

FPX-Fixes-SHA256SUMS-20261021-01.txt

@@ -0,0 +1,6 @@
+dcdd6a8ebca68467f3faa9fd07e2f69f11e25446f100b4c214cc2619fc681f6b  aitable-fpx-native-fixes-20261021-01.patch
+0b11b7f5b1c95f69524efcad44d4f1c3f1b9102d403380b1ca7dea0e5102c087  AiTable-ROCmFPX-Optimizations-Fixes-20261021-01.zip
+1972b0408bd74f1d0e5695bbd1dae20844d043528f464ae26ca6069564d89d4d  apply-fpx-fixes-20261021-01.py
+6ffc49b30005d3d0343710ee4f8ed4a28f4fc16d334db24fe251a9ef7d60dcd2  configure-fpx-env-20261021-01.cmd
+bee1ef0d41e4fe3b63bf5b95367ff7a760b63bc56a9296bb925e40e7f4a480f5  fpx-fixes-manifest-20261021-01.json
+1e52743e8668db18f93be8e2396ec0bddb5f9ca68d473a2fbf99be1b1092c761  FPX-Optimizations-20261021-01.md

FPX-Optimizations-20261021-01.md

@@ -0,0 +1,95 @@
+# AiTable ROCmFPX optimizations and fixes — 20261021-01
+
+## 下載與配對 / Downloads and pairing
+
+This is the **script/source-diff add-on** for the existing [Windows gfx1201 Runtime 20261021-01](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/resolve/main/AiTable-ROCmFPX-gfx1201-Windows-x64-20261021-01.zip). It is not a replacement model or a new inference binary.
+
+- [Complete scripts, patch and build support ZIP](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/resolve/main/AiTable-ROCmFPX-Optimizations-Fixes-20261021-01.zip)
+- [Application script](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/apply-fpx-fixes-20261021-01.py)
+- [Native patch](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/aitable-fpx-native-fixes-20261021-01.patch)
+- [Source/hash manifest](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/fpx-fixes-manifest-20261021-01.json)
+- [Runtime environment preset](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/configure-fpx-env-20261021-01.cmd)
+- [Add-on checksums](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/resolve/main/FPX-Fixes-SHA256SUMS-20261021-01.txt)
+- [Paired model weights](https://huggingface.co/Roland2531/Ornith-1.5-9B-uncensored-ROCmI4-Q6IO-A8)
+
+已經下載 Runtime 的使用者不用對 EXE 套 patch。Runtime 已包含該版原生修改；套用腳本是提供建置者將修改帶到指定上游來源，環境腳本則供現有 Runtime 設定目前命令列環境。所有檔案以固定版本名稱發布，之後新增版本，不覆蓋舊版。
+
+## Scope and provenance
+
+Upstream: [charlie12345/ROCmFPX](https://github.com/charlie12345/ROCmFPX), exact base **fb08d7cdb67082ddda15a25e64dc7ec413afc41c**. The native delta was taken from the source snapshot frozen alongside the already-published Runtime. All 23 differing source files also matched the current project source at preparation time. Existing upstream features, quantization formats and dependencies remain upstream work; the diff is an integration record, not a claim that AiTable invented the underlying methods.
+
+This add-on includes a source patch, application/environment scripts, license notices and the tiny Windows CK import library/DEF required for linking. It excludes the full C/C++ source snapshot, build caches/objects, SDK, private logs, prompts, conversations, model weights and credentials. It does not include any newly trained model or new quantization.
+
+## What is included
+
+| Area | Files / mechanism | Value and limits |
+| --- | --- | --- |
+| Short prefix checkpoints | `prefix-short-policy.h`, server context; 1024/2048 token boundaries, four checkpoint slots | Enables earlier compatible-prefix reuse on the supported hybrid model path. It cannot reuse changed token content or guarantee a fixed 3K-prompt latency. |
+| RAM checkpoint restore | `src/llama-context.cpp` plus server cache integration | Retains the installed restore correction. RAM prompt cache must be explicitly enabled; it is host memory, separate from GPU KV and model weights. Background tasks can still require new work when no compatible checkpoint exists. |
+| Tool grammar and parsing | `common/chat-auto-parser-generator.cpp`, chat, PEG, JSON-Schema grammar and sampling files | Tagged tool routing, string/JSON escape and bounded-length handling, delimiter/whitespace behavior, prefill/finish/error diagnostics. Grammar syntax constraints do not prove correct tool semantics or delivery. |
+| Native Flash Attention routing | `fattn.cu`, `fattn-common.cuh` | Retains integrated routing adjustments for the older fork, including the supported gfx1201/head256 path. Not every attention shape takes that route. |
+| RDNA4 hardware transpose | `fattn-mma-f16.cuh` | Uses guarded gfx1201 F16 transposed V loads with alignment, shape and existing LDS-capacity checks. It is not universal FP8 KV or native W4A4 acceleration. |
+| Projector lifecycle | server projector/queue/context files | Retains installed idle/unload and active-work coordination. No vision model is included or loaded by these scripts. |
+| Optional offline prefix probe | `tools/cli/prefix-offline-probe.cpp` and target | A measurement helper, not an extra inference optimization. Its own usage requires controlled one-slot/checkpoint inputs and no RAM cache; do not treat it as the full application/HTTP playback path. |
+| Experimental CK bridge | guarded dispatch and matching import library | Present for binary compatibility. The environment preset leaves CK off unless explicitly opted in; prior input-dependent output differences remain documented. No fix for all CK numerical differences is claimed. |
+
+These are retained integration changes, not a new benchmark. The add-on does not add rejected RDNA4 attention v2, research GDN, norm/SiLU fusion or experimental Q8/FP8 attention kernels. Disabled environment toggles do not prove such kernels exist in the binary.
+
+## Option A — use the existing runtime / 使用現有後端
+
+Extract the existing Runtime ZIP, preserving `runtime/` and `.kpack/`. Extract the add-on separately. In **Command Prompt**, call the preset with the Runtime package root:
+
+```bat
+call "D:\FPX-Fixes\configure-fpx-env-20261021-01.cmd" "D:\FPX-Runtime"
+"D:\FPX-Runtime\runtime\llama-server.exe" --model "D:\Models\your-model.gguf" --host 127.0.0.1 --port 8080 --parallel 1 --ctx-size 8192 --n-gpu-layers 99 --flash-attn on --cache-type-k q8_0 --cache-type-v q8_0 --cache-ram 1024
+```
+
+Paths and context are examples. The script only sets the current CMD environment; it does not launch, stop or replace a service, modify persistent AiTable settings, change model temperature/top-p/token limits or force a model's thinking mode. Keep the model/application's intended sampling and thinking configuration. Do not add thinking-disable flags that the selected model/template does not support.
+
+Preset contents:
+
+- `FPX_SHORT_PREFIX_CHECKPOINT=1` for the installed short-prefix policy.
+- `LLAMA_ARG_CACHE_RAM=1024`: a **1024 MiB host-RAM cap**, not reserved VRAM. Explicit CLI `--cache-ram` can set a different cap. `0` disables it; this preset does not use unlimited `-1`.
+- Package-local rocBLAS Tensile path and `ROCBLAS_USE_HIPBLASLT=0`, preserving this release's rocBLAS choice.
+- `GGML_CUDA_DISABLE_ROPE_FWHT_SET_ROWS=1` and the published release's retired research switches set to 0.
+- CK default `GGML_HIP_CK_ATTENTION=0` if the caller has not explicitly set it. Set 0 explicitly to clear a previous opt-in. Enabling CK is experimental and needs your own output checks.
+
+The previously published `run-server.cmd` example did not explicitly enable RAM cache. Use the example above or pass the cache cap in your own integration when RAM prompt caching is intended. Short checkpoint creation and RAM cache storage are different mechanisms. Observe actual server logs/cache-token/timing counters to verify reuse; an enabled setting alone is not acceleration evidence.
+
+## Option B — apply native fixes to source / 套用原生修正
+
+Requires Python 3.9+ and Git. Use a separate clean checkout of the exact upstream commit, preserving LF in both clone and checkout:
+
+```bat
+git -c core.autocrlf=false clone https://github.com/charlie12345/ROCmFPX.git ROCmFPX-work
+git -C ROCmFPX-work -c core.autocrlf=false checkout fb08d7cdb67082ddda15a25e64dc7ec413afc41c
+python apply-fpx-fixes-20261021-01.py --source ROCmFPX-work
+python apply-fpx-fixes-20261021-01.py --source ROCmFPX-work --apply
+```
+
+Default is check-only. The script validates patch SHA256 and each affected original file, rejects wrong/already-patched sources, backs up originals and lists new files, applies atomically with `git apply`, then checks all 23 resulting hashes. Git automatic newline conversion is disabled for patch application. Backups are placed beside the checkout with a unique name and `restore-manifest.json`. To restore, verify no later edits would be overwritten, restore backed-up originals and remove only manifest-listed originally absent files. The script does not silently merge changes, compile or install anything.
+
+### CK linking and source build boundary
+
+The native snapshot linked CK via an AiTable-specific sibling folder. **The only publication adjustment to that snapshot** is replacing that CMake path with the explicit cache parameter `FPX_CK_IMPORT_LIBRARY`. At configure time, supply the full path to the import library in the complete add-on ZIP, for example:
+
+```text
+-DFPX_CK_IMPORT_LIBRARY=D:/FPX-Fixes/dependencies/CK/ck-attention.lib
+```
+
+The import library is 1692 bytes; it contains Windows DLL import records, not the full CK implementation. The matching `ck-attention.dll` remains in the Runtime package; SHA256 `96ecc9b94fc7ecc4f2e2dcccc3230b8b5bb950c5774ea7a166665f07a4376ea2`. Keep its license and runtime dependencies. This add-on does not include CK kernel sources or reconstruct that DLL.
+
+Follow upstream's build instructions with your own Windows/MSVC/CMake/Ninja and HIP/ROCm setup; the published binary's provenance records LLVM24/ROCm10.1 and gfx1201. Dependencies, compiler/SDK compatibility and a full clean build must be validated separately. No retained object files, local absolute SDK paths or automated overwrite of an installed runtime are supplied. A fresh full build and bit-identical reproduction are **not** claimed by this release.
+
+## Checks performed for this add-on
+
+- Patch preflight on the fixed upstream checkout: passed.
+- Actual application plus all 23 output hashes: passed.
+- Repeat application, locally modified affected source, and tampered patch: rejected; the modified-source case made no writes.
+- Runtime environment variables and existing server `--version`: passed; no model or HTTP service started.
+- Complete add-on ZIP CRC and download checksums: checked separately at publication.
+- No new GPU, model-quality, tool-accuracy, prefill, decode, TTFT or VRAM benchmark was run. Earlier results remain workload-specific; no new performance gain is asserted.
+
+## License and release retention
+
+Native core changes retain upstream MIT terms (`LICENSE`); CK import support has its preserved CK license (`CK-LICENSE.txt`), and runtime/SDK dependencies retain their own notices. Model license terms are separate. Keep prior versioned archives/scripts/patches/manifests/notes/checksums unchanged. This add-on only adds new files and updates the mutable repository README index; it does not replace Runtime 20261021-01.

README.md

@@ -59,7 +59,7 @@ Short-prefix reuse is content-dependent. No universal decode speedup or full-wor
 
 `build-manifest-20261021-01.json` records the actual included binary hashes. `SHA256SUMS-20261021-01.txt` at the repository root covers downloadable files; the runtime ZIP contains a separate checksum list for extracted payloads. Native `--version` / `--help` launch checks are performed from the standalone package; no new GPU model benchmark or other-GPU validation is implied.
 
-The integrated C/C++ source snapshot is retained locally and is not uploaded in this binary release. Upstream source links and all bundled license notices remain provided. The installed build used retained build objects and locally prepared dependency wiring; a fully reproducible public source build is not claimed.
+The complete integrated C/C++ source snapshot is retained locally and is not uploaded. A curated native patch/application-script add-on is now available below. Upstream source links and all bundled license notices remain provided. The installed build used retained build objects and locally prepared dependency wiring; a fully reproducible public source build is not claimed.
 
 ## Versioned releases / 舊版本保留
 
@@ -72,3 +72,14 @@ Each release has an immutable version identifier. Future updates add a new ZIP,
 ## License
 
 llama.cpp/ROCmFPX core code is MIT; preserve `LICENSE` and `runtime/licenses/`. Bundled AMD and other dependencies retain their own licenses, including MIT and Apache-2.0 with LLVM exceptions where applicable. See `THIRD_PARTY_NOTICES.md`. The repository's MIT metadata does not relicense the paired model weights or every third-party binary. Model terms are documented in the model repository.
+
+## Optimization and fix scripts / 優化與修正腳本
+
+- [Complete add-on ZIP — 20261021-01](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/resolve/main/AiTable-ROCmFPX-Optimizations-Fixes-20261021-01.zip)
+- [Scope, configuration, source application and limits](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/FPX-Optimizations-20261021-01.md)
+- [Source application script](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/apply-fpx-fixes-20261021-01.py)
+- [Runtime environment preset](https://huggingface.co/Roland2531/llama.cpp-ROCmFPX-gfx1201-Windows-x64/blob/main/configure-fpx-env-20261021-01.cmd)
+
+Includes installed short-prefix/RAM restore, tool grammar/parser/sampler fixes and guarded Flash Attention/transpose integration. The environment preset enables a bounded 1024 MiB host-RAM prompt cache; the earlier standalone run-server example did not explicitly enable it. CK remains off by default in the standalone preset. No new inference binary, model weights or fresh GPU speedup is claimed. The source script was checked/applied on the pinned upstream commit and verifies all 23 changed files; full rebuild reproducibility remains unverified.
+
+The existing Runtime ZIP, version-specific notes, manifest and checksum files remain unchanged. Initial Runtime checksums describe its original release revision; this root README is a mutable index.

aitable-fpx-native-fixes-20261021-01.patch

@@ -0,0 +1,1959 @@
+--- a/common/chat-auto-parser-generator.cpp
++++ b/common/chat-auto-parser-generator.cpp
+@@ -94,6 +94,26 @@
+     return data;
+ }
+ 
++// A required tagged call with thinking explicitly off has no generated prose
++// channel. Preserve templates that prefill an open reasoning block: those need
++// their existing parser to finish it before the call can begin.
++static bool require_tagged_call_only(const generation_params & inputs, const analyze_reasoning & reasoning) {
++    if (inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_REQUIRED || inputs.enable_thinking) {
++        return false;
++    }
++    if (reasoning.mode == reasoning_mode::NONE) {
++        return true;
++    }
++    const auto start = trim_whitespace(reasoning.start);
++    const auto end = trim_whitespace(reasoning.end);
++    if (start.empty() || end.empty()) {
++        return false;
++    }
++    const auto opened = inputs.generation_prompt.rfind(start);
++    const auto closed = inputs.generation_prompt.rfind(end);
++    return opened == std::string::npos || (closed != std::string::npos && closed > opened);
++}
++
+ common_peg_arena autoparser::build_parser(const generation_params & inputs) const {
+     if (!analysis_complete) {
+         throw std::invalid_argument("Cannot call build_parser on autoparser without performing analysis first, call analyze_template(...)");
+@@ -108,12 +128,17 @@
+ 
+         // Build reasoning parser
+         ctx.reasoning_parser = reasoning.build_parser(ctx);
+-
+         auto parser = p.eps();
+ 
+         bool has_tools           = inputs.tools.is_array() && !inputs.tools.empty();
+         bool has_response_format = inputs.json_schema.is_object() && !inputs.json_schema.empty();
+         bool pure_content        = reasoning.mode == reasoning_mode::NONE;
++        const bool direct_tagged = !has_response_format && has_tools && jinja_caps.supports_tool_calls &&
++                                   tools.format.mode == tool_format::TAG_WITH_TAGGED &&
++                                   require_tagged_call_only(inputs, reasoning);
++        if (direct_tagged) {
++            ctx.reasoning_parser = p.eps();
++        }
+ 
+         if (has_response_format) {
+             auto response_format = p.rule("response-format", p.content(p.schema(p.json(), "response-format-schema", inputs.json_schema)));
+@@ -128,7 +153,10 @@
+         } else {
+             parser = content.build_parser(ctx);
+         }
+-        return pure_content ? p.prefix(inputs.generation_prompt, reasoning.start) + parser : p.prefix(inputs.generation_prompt, reasoning.start) << parser;
++        // A closed no-thinking prefill is fixed template text, not permission
++        // to start another reasoning block in the generated suffix.
++        const auto prefix = p.prefix(inputs.generation_prompt, direct_tagged ? "" : reasoning.start);
++        return pure_content ? prefix + parser : prefix << parser;
+     });
+ }
+ 
+@@ -357,7 +385,20 @@
+     auto &       p           = ctx.p;
+     const auto & inputs      = ctx.inputs;
+ 
+-    auto until_suffix = p.rule("until-suffix", p.until(arguments.value_suffix));
++    // Tagged empty values may share whitespace between their opening and
++    // closing delimiters. Only an exact whitespace overlap can be reused.
++    const auto value_open = arguments.name_suffix + arguments.value_prefix;
++    size_t shared_space = 0;
++    for (size_t n = 1; n < arguments.value_suffix.size() && n <= value_open.size(); ++n) {
++        const auto prefix = arguments.value_suffix.substr(0, n);
++        if (prefix.find_first_not_of(" \t\r\n") != std::string::npos) { break; }
++        if (value_open.compare(value_open.size() - n, n, prefix) == 0) { shared_space = n; }
++    }
++    const auto empty_close = arguments.value_suffix.substr(shared_space);
++    // Stop at either close, so the ordinary string branch cannot consume an
++    // empty value's closing tag and reinterpret later tools as its contents.
++    auto until_suffix = p.rule("until-suffix", shared_space ?
++        p.until_one_of({arguments.value_suffix, empty_close}) : p.until(arguments.value_suffix));
+ 
+     common_peg_parser tool_choice = p.choice();
+ 
+@@ -381,15 +422,20 @@
+         for (const auto & [param_name, param_schema] : properties.items()) {
+             bool is_required = required.find(param_name) != required.end();
+ 
+-            auto arg =
+-                p.tool_arg(p.tool_arg_open(arguments.name_prefix + p.tool_arg_name(p.literal(param_name)) +
+-                                           arguments.name_suffix) +
+-                           arguments.value_prefix +
+-                           (schema_info.resolves_to_string(param_schema) ?
+-                                p.tool_arg_string_value(until_suffix) :
+-                                p.tool_arg_json_value(p.schema(
+-                                    p.json(), "tool-" + name + "-arg-" + param_name + "-schema", param_schema, false))) +
+-                           p.tool_arg_close(p.literal(arguments.value_suffix)));
++            const bool is_string = schema_info.resolves_to_string(param_schema);
++            auto value_and_close =
++                (is_string ? p.tool_arg_string_value(until_suffix) :
++                    p.tool_arg_json_value(p.schema(
++                        p.json(), "tool-" + name + "-arg-" + param_name + "-schema", param_schema, false))) +
++                p.tool_arg_close(p.literal(arguments.value_suffix));
++            if (is_string && shared_space) {
++                value_and_close = p.choice({
++                    p.tool_arg_string_value(p.eps()) + p.tool_arg_close(p.literal(empty_close)),
++                    value_and_close});
++            }
++            auto arg = p.tool_arg(
++                p.tool_arg_open(arguments.name_prefix + p.tool_arg_name(p.literal(param_name)) + arguments.name_suffix) +
++                arguments.value_prefix + value_and_close);
+ 
+             auto named_arg = p.rule("tool-" + name + "-arg-" + param_name, arg);
+             if (is_required) {
+@@ -478,6 +524,9 @@
+         tool_calls = p.optional(tool_calls);
+     }
+ 
++    if (require_tagged_call_only(inputs, *ctx.reasoning)) {
++        return p.space() + tool_calls + p.end();
++    }
+     std::string trigger_marker       = !format.section_start.empty() ? format.section_start : format.per_call_start;
+     auto        content_before_tools = trigger_marker.empty() ? p.eps() : p.until(trigger_marker);
+     return ctx.reasoning_parser + p.optional(p.content(content_before_tools)) + tool_calls + p.end();
+--- a/common/chat.cpp
++++ b/common/chat.cpp
+@@ -1,4 +1,6 @@
+ #include "chat.h"
++#include "jinja/utils.h"
++#include <iomanip>
+ 
+ #include "chat-auto-parser-helpers.h"
+ #include "chat-auto-parser.h"
+@@ -2364,6 +2366,66 @@
+     return common_chat_peg_parse(params.parser, input, is_partial, params);
+ }
+ 
++
++static int64_t peg_diag_bound(size_t value, size_t limit = 2147483647) {
++    return static_cast<int64_t>(std::min(value, limit));
++}
++
++static bool peg_diag_utf8(const std::string & s) {
++    for (size_t i = 0; i < s.size();) {
++        unsigned char c = static_cast<unsigned char>(s[i++]);
++        if (c < 0x80) continue;
++        int n = c >= 0xc2 && c <= 0xdf ? 1 : c >= 0xe0 && c <= 0xef ? 2 : c >= 0xf0 && c <= 0xf4 ? 3 : -1;
++        if (n < 0 || i + n > s.size()) return false;
++        uint32_t cp = c & (n == 1 ? 0x1f : n == 2 ? 0x0f : 0x07);
++        for (int j = 0; j < n; ++j) {
++            unsigned char b = static_cast<unsigned char>(s[i++]);
++            if ((b & 0xc0) != 0x80) return false;
++            cp = (cp << 6) | (b & 0x3f);
++        }
++        if ((n == 1 && cp < 0x80) || (n == 2 && cp < 0x800) || (n == 3 && cp < 0x10000)
++                || cp > 0x10ffff || (cp >= 0xd800 && cp <= 0xdfff)) return false;
++    }
++    return true;
++}
++
++static json peg_diag_structure(const common_peg_arena & parser, const std::string & input,
++        const std::string & effective, const std::string & prefix, bool partial, size_t end) {
++    json data={{"schema","aitable.native_peg@1"},{"phase",partial?"partial":"final"},
++        {"result_end",peg_diag_bound(end)},{"input_bytes",peg_diag_bound(input.size())},
++        {"effective_bytes",peg_diag_bound(effective.size())},{"generation_prompt_bytes",peg_diag_bound(prefix.size())},
++        {"utf8_valid",peg_diag_utf8(effective)}};
++    const std::pair<const char *, const char *> markers[]={
++        {"tool_call_open_count","<tool_call>"},{"tool_call_close_count","</tool_call>"},
++        {"function_open_count","<function="},{"function_close_count","</function>"},
++        {"parameter_open_count","<parameter="},{"parameter_close_count","</parameter>"},
++        {"think_open_count","<think>"},{"think_close_count","</think>"}};
++    for (const auto & marker:markers) {
++        size_t count=0,pos=0;
++        while ((pos=effective.find(marker.second,pos))!=std::string::npos) {
++            ++count;pos+=std::char_traits<char>::length(marker.second);
++            if (count==65535) break;
++        }
++        data[marker.first]=peg_diag_bound(count,65535);
++    }
++    const auto last=effective.find_last_not_of(" \t\r\n");
++    std::string tail=effective.empty()?"empty":last==std::string::npos?"whitespace":"other";
++    if (last!=std::string::npos) {
++        const std::pair<const char *,const char *> endings[]={
++            {"tool_call_close","</tool_call>"},{"function_close","</function>"},
++            {"parameter_close","</parameter>"},{"json_object_close","}"},{"json_array_close","]"}};
++        for (const auto & ending:endings) {
++            size_t n=std::char_traits<char>::length(ending.second);
++            if (last+1>=n && effective.compare(last+1-n,n,ending.second)==0) {tail=ending.first;break;}
++        }
++    }
++    data["tail"]=tail;
++    jinja::hasher hash;hash.update(parser.save());
++    std::ostringstream fingerprint;fingerprint<<std::hex<<std::setfill('0')<<std::setw(16)<<hash.digest();
++    data["parser_hash"]=fingerprint.str();data["parser_hash_alg"]="jinja-hasher64";
++    return data;
++}
++
+ common_chat_msg common_chat_peg_parse(const common_peg_arena &          src_parser,
+                                       const std::string &               input,
+                                       bool                              is_partial,
+@@ -2411,10 +2473,11 @@
+             }
+             return msg;
+         }
+-        LOG_WRN("%s: unparsed %s output: %s\n", __func__, common_chat_format_name(params.format),
+-                effective_input.substr(result.end).c_str());
+-        throw std::runtime_error(std::string("The model produced output that does not match the expected ") +
+-                                 common_chat_format_name(params.format) + " format");
++        const auto diagnostics=peg_diag_structure(parser,input,effective_input,params.generation_prompt,is_partial,result.end);
++        LOG_WRN("%s: unparsed %s output (phase=%s, bytes=%zu, result_end=%zu)\n", __func__,
++                common_chat_format_name(params.format), is_partial?"partial":"final", effective_input.size(), result.end);
++        throw common_chat_peg_error(std::string("The model produced output that does not match the expected ") +
++                                   common_chat_format_name(params.format) + " format",diagnostics.dump());
+     }
+ 
+     common_chat_msg msg;
+--- a/common/chat.h
++++ b/common/chat.h
+@@ -11,6 +11,7 @@
+ #include "nlohmann/json_fwd.hpp"
+ 
+ #include <chrono>
++#include <stdexcept>
+ #include <functional>
+ #include <map>
+ #include <string>
+@@ -18,6 +19,15 @@
+ 
+ using chat_template_caps = jinja::caps;
+ using json = nlohmann::ordered_json;
++
++// Metadata contains only the finite native PEG diagnostic contract. Keep the
++// original exception message intact across the common DLL/server boundary.
++class common_chat_peg_error : public std::runtime_error {
++public:
++    std::string native_peg_json;
++    common_chat_peg_error(const std::string & message, const std::string & diagnostics)
++        : std::runtime_error(message), native_peg_json(diagnostics) {}
++};
+ 
+ struct common_chat_templates;
+ 
+--- a/common/common.h
++++ b/common/common.h
+@@ -609,9 +609,9 @@
+     int32_t n_cache_reuse       = 0;     // min chunk size to reuse from the cache via KV shifting
+     bool    cache_prompt        = true;  // whether to enable prompt caching
+     bool    cache_idle_slots    = true;  // save and clear idle slots upon starting a new task
+-    int32_t n_ctx_checkpoints   = 32;    // max number of context checkpoints per slot
+-    int32_t checkpoint_every_nt = 8192;  // make a checkpoint every n tokens during prefill
+-    int32_t cache_ram_mib       = 8192;  // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.
++    int32_t n_ctx_checkpoints   = 4;    // max number of context checkpoints per slot
++    int32_t checkpoint_every_nt = 2048;  // make a checkpoint every n tokens during prefill
++    int32_t cache_ram_mib       = 0;  // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.
+     int32_t cache_disk_limit_mib = 8192; // disk prompt-cache limit when cache_disk_path is set
+ 
+     std::string hostname      = "127.0.0.1";
+--- a/common/json-schema-to-grammar.cpp
++++ b/common/json-schema-to-grammar.cpp
+@@ -343,6 +343,115 @@
+             rules.push_back(visit(alt_schemas[i], name + (name.empty() ? "alternative-" : "-") + std::to_string(i)));
+         }
+         return string_join(rules, " | ");
++    }
++
++    // A repeated ASCII character class is a language of decoded Unicode scalars,
++    // not JSON source characters. Lower it through JSON encoding before applying
++    // length bounds. Other regex forms continue through the existing converter.
++    bool _visit_encoded_charclass(const json & schema, const std::string & name, std::string & rule) {
++        const std::string pattern = schema.at("pattern");
++        if (pattern.size() < 6 || pattern[0] != '^' || pattern[1] != '[' ||
++                pattern.substr(pattern.size()-3) != "]*$") { return false; }
++        size_t pos = 2, end = pattern.size()-3;
++        bool negated = pos < end && pattern[pos] == '^';
++        if (negated) { ++pos; }
++        bool included[128] = {};
++        auto hex_value = [](char c) -> int {
++            if (c >= '0' && c <= '9') return c-'0';
++            if (c >= 'a' && c <= 'f') return c-'a'+10;
++            if (c >= 'A' && c <= 'F') return c-'A'+10;
++            return -1;
++        };
++        auto atom = [&](int & value) -> bool {
++            if (pos >= end) return false;
++            unsigned char c = pattern[pos++];
++            if (c >= 128 || c == ']') return false;
++            if (c != '\\') { value=c; return true; }
++            if (pos >= end) return false;
++            char escape=pattern[pos++];
++            if (escape=='x' || escape=='u') {
++                const size_t count=escape=='x'?2:4;
++                if (pos+count>end) return false;
++                value=0;
++                for(size_t i=0;i<count;++i){int digit=hex_value(pattern[pos++]);if(digit<0)return false;value=value*16+digit;}
++                return value<128;
++            }
++            switch(escape){
++                case 'n':value=10;return true; case 'r':value=13;return true;
++                case 't':value=9;return true; case 'b':value=8;return true;
++                case 'f':value=12;return true;
++                default:
++                    if(std::string("\\][^-/.+*?(){}|$\"").find(escape)==std::string::npos)return false;
++                    value=static_cast<unsigned char>(escape);return true;
++            }
++        };
++        bool any=false;
++        while(pos<end){
++            int lo,hi;if(!atom(lo))return false;hi=lo;
++            if(pos+1<end && pattern[pos]=='-'){++pos;if(!atom(hi)||hi<lo)return false;}
++            for(int c=lo;c<=hi;++c)included[c]=true;
++            any=true;
++        }
++        if(!any)return false;
++        bool allowed[128];for(int c=0;c<128;++c)allowed[c]=negated?!included[c]:included[c];
++        int min_len=schema.value("minLength",0),max_len=schema.value("maxLength",std::numeric_limits<int>::max());
++        if(min_len<0||max_len<min_len){_errors.push_back("Invalid string length bounds");rule="";return true;}
++        std::vector<std::string> alternatives;
++        std::string raw="[";
++        for(int c=32;c<128;++c){
++            if(!allowed[c]||c=='"'||c=='\\')continue;
++            int lo=c;while(c+1<128&&allowed[c+1]&&c+1!='"'&&c+1!='\\')++c;
++            char buf[12];snprintf(buf,sizeof(buf),"\\x%02X",lo);raw+=buf;
++            if(c!=lo){snprintf(buf,sizeof(buf),"-\\x%02X",c);raw+=buf;}
++        }
++        if(raw.size()>1)alternatives.push_back(raw+"]");
++        if(negated)alternatives.push_back("[\\u0080-\\uD7FF\\uE000-\\U0010FFFF]");
++        for(const auto & item:std::vector<std::pair<int,std::string>>{{34,"\\\""},{92,"\\\\"},{47,"\\/"},{8,"\\b"},{12,"\\f"},{10,"\\n"},{13,"\\r"},{9,"\\t"}}){
++            if(allowed[item.first])alternatives.push_back(format_literal(item.second));
++        }
++        const auto hex=_add_rule("json-pattern-hex","[0-9a-fA-F]");
++        auto hex_class=[](unsigned mask){
++            std::string result="[";
++            for(unsigned i=0;i<16;++i)if(mask&(1u<<i)){
++                result+=char(i<10?'0'+i:'a'+i-10);if(i>=10)result+=char('A'+i-10);
++            }
++            return result+"]";
++        };
++        std::map<unsigned,unsigned> groups;
++        for(unsigned high=0;high<8;++high){unsigned lowmask=0;for(unsigned low=0;low<16;++low)if(allowed[high*16+low])lowmask|=1u<<low;if(lowmask)groups[lowmask]|=1u<<high;}
++        std::vector<std::string> ascii_unicode;
++        for(auto group:groups)ascii_unicode.push_back(hex_class(group.second)+" "+hex_class(group.first));
++        if(!ascii_unicode.empty())alternatives.push_back(format_literal("\\u00")+" ("+string_join(ascii_unicode," | ")+")");
++        if(negated){
++            const auto u=format_literal("\\u");
++            alternatives.push_back(u+" (\"00\" [89aAbBcCdDeEfF] "+hex+" | \"0\" [1-9a-fA-F] "+hex+"{2} | [1-9a-cA-CeEfF] "+hex+"{3} | [dD] [0-7] "+hex+"{2})");
++            // One surrogate pair encodes one scalar for min/maxLength.
++            alternatives.push_back(u+" [dD] [89aAbB] "+hex+"{2} "+u+" [dD] [c-fC-F] "+hex+"{2}");
++        }
++        if(alternatives.empty()){
++            if(min_len){_errors.push_back("Pattern has no allowed JSON scalar");rule="";return true;}
++            rule=_add_rule(name,format_literal("\"\"")+" space");return true;
++        }
++        const auto encoded_char=_add_rule("json-pattern-char",string_join(alternatives," | "));
++        // Reuse powers of two instead of expanding a {0,4096} quantifier beyond
++        // the native grammar parser's repetition guard. This preserves the bound.
++        std::function<std::string(int)> block = [&](int count) -> std::string {
++            if(count==1)return encoded_char;
++            const auto half=block(count/2);
++            return _add_rule(encoded_char+"-block-"+std::to_string(count),half+" "+half);
++        };
++        std::function<std::string(int)> upto = [&](int count) -> std::string {
++            if(!count)return format_literal("");
++            if(count==1)return _add_rule(encoded_char+"-upto-1",encoded_char+"?");
++            int power=1;while(power<=count/2)power*=2;
++            return _add_rule(encoded_char+"-upto-"+std::to_string(count),upto(power-1)+" | "+block(power)+" "+upto(count-power));
++        };
++        std::string repetition;
++        int remaining=min_len,power=1;
++        while(remaining){if(remaining&1)repetition+=block(power)+" ";remaining/=2;if(remaining)power*=2;}
++        repetition+=schema.contains("maxLength")?upto(max_len-min_len):encoded_char+"*";
++        rule=_add_rule(name,format_literal("\"")+" "+repetition+" "+format_literal("\"")+" space");
++        return true;
+     }
+ 
+     std::string _visit_pattern(const std::string & pattern, const std::string & name) {
+@@ -959,6 +1068,8 @@
+             return _add_rule(rule_name, "\"[\" space " + build_repetition(item_rule_name, min_items, max_items, "\",\" space") + " \"]\" space");
+         }
+         if ((schema_type.is_null() || schema_type == "string") && schema.contains("pattern")) {
++            std::string encoded;
++            if (_visit_encoded_charclass(schema, rule_name, encoded)) return encoded;
+             return _visit_pattern(schema["pattern"], rule_name);
+         }
+         if ((schema_type.is_null() || schema_type == "string") && std::regex_match(schema_format, std::regex("^uuid[1-5]?$"))) {
+--- a/common/peg-parser.cpp
++++ b/common/peg-parser.cpp
+@@ -8,6 +8,7 @@
+ #include <algorithm>
+ #include <initializer_list>
+ #include <map>
++#include <set>
+ #include <memory>
+ #include <nlohmann/json.hpp>
+ #include <regex>
+@@ -93,6 +94,10 @@
+         std::vector<uint32_t> next_chars;
+     };
+ 
++    // Keep the existing exported trie instantiations in the frozen Windows ABI.
++#if defined(__clang__) || defined(__GNUC__)
++    __attribute__((used, noinline))
++#endif
+     std::vector<prefix_and_next> collect_prefix_and_next() {
+         std::vector<uint32_t>        prefix;
+         std::vector<prefix_and_next> result;
+@@ -1342,7 +1347,7 @@
+ common_peg_parser common_peg_parser_builder::json_array() {
+     return rule("json-array", [this]() {
+         auto ws = space();
+-        auto elements = sequence({json(), zero_or_more(sequence({literal(","), ws, json()}))});
++        auto elements = sequence({json(), zero_or_more(sequence({ws, literal(","), ws, json()}))});
+         return sequence({
+             literal("["),
+             ws,
+@@ -1502,54 +1507,114 @@
+     return std::string(buf);
+ }
+ 
+-static std::string gbnf_excluding_pattern(const std::vector<std::string> & strings) {
+-    trie matcher(strings);
+-    auto pieces = matcher.collect_prefix_and_next();
+-
+-    std::string pattern;
+-    std::string trailing;  // optional proper-prefix of a delimiter, allowed only at the very end
+-    for (size_t i = 0; i < pieces.size(); ++i) {
+-        if (i > 0) {
+-            pattern += " | ";
+-        }
+-
+-        const auto & pre = pieces[i].prefix;
+-        const auto & chars = pieces[i].next_chars;
+-
+-        std::string cls;
+-        cls.reserve(chars.size());
+-        for (uint32_t ch : chars) {
+-            cls += gbnf_escape_char_class(ch);
+-        }
+-
+-        if (!pre.empty()) {
+-            std::string pre_literal = gbnf_format_literal(common_unicode_cpts_to_utf8(pre));
+-            pattern += pre_literal + " [^" + cls + "]";
+-            // Each interior alternative consumes a delimiter-prefix plus a disambiguating
+-            // char, so the repetition alone cannot match a value that *ends* on a proper
+-            // prefix of a delimiter (e.g. a trailing "\n" when the delimiter is
+-            // "\n</parameter>\n"). The runtime until() (greedy first-match) accepts such
+-            // values, so without this the grammar would reject input the parser accepts.
+-            // Allow the value to terminate on any proper prefix as an optional tail.
+-            // This makes the grammar a slight superset of the runtime language (a value
+-            // may end on the longest prefix, which greedy first-match would not itself
+-            // produce); harmless for constrained generation, which only needs to admit
+-            // every runtime-valid string.
+-            if (!trailing.empty()) {
+-                trailing += " | ";
+-            }
+-            trailing += pre_literal;
+-        } else {
+-            pattern += "[^" + cls + "]";
+-        }
+-    }
+-
+-    std::string result = "(" + pattern + ")*";
+-    if (!trailing.empty()) {
+-        result += " (" + trailing + ")?";
++// Track the longest suffix that is also a delimiter prefix. Resetting at each
++// prefix/mismatch fragment loses overlaps such as "<<" before "</...>".
++static std::vector<uint32_t> until_codepoints(const std::string & text) {
++    std::vector<uint32_t> result;
++    for (size_t pos = 0; pos < text.size();) {
++        const auto decoded = common_parse_utf8_codepoint(text, pos);
++        if (decoded.status != utf8_parse_result::SUCCESS) {
++            throw std::invalid_argument("Invalid UTF-8 in PEG delimiter");
++        }
++        result.push_back(decoded.codepoint);
++        pos += decoded.bytes_consumed;
+     }
+     return result;
+ }
++
++using until_word = std::vector<uint32_t>;
++
++static bool until_has_suffix(const until_word & value, const until_word & suffix) {
++    return value.size() >= suffix.size() &&
++           std::equal(suffix.rbegin(), suffix.rend(), value.rbegin());
++}
++
++static bool until_can_join(const std::vector<std::string> & delimiters, const std::string & close) {
++    if (close.empty() || std::find(delimiters.begin(), delimiters.end(), close) == delimiters.end()) {
++        return false;
++    }
++    const auto close_length = until_codepoints(close).size();
++    // This covers every pending earlier start before this continuation ends.
++    // Longer lookahead-dependent continuations retain their separate grammar;
++    // they cannot be fused by inspecting only these two PEG nodes.
++    return std::all_of(delimiters.begin(), delimiters.end(), [&](const std::string & delimiter) {
++        return until_codepoints(delimiter).size() <= close_length + 1;
++    });
++}
++
++static std::string gbnf_excluding_pattern(const std::vector<std::string> & strings,
++                                          const common_grammar_builder & builder,
++                                          const std::string & base,
++                                          const std::string * close = nullptr) {
++    std::vector<until_word> delimiters;
++    std::set<until_word> prefixes{until_word{}};
++    std::set<uint32_t> alphabet;
++    for (const auto & text : strings) {
++        if (text.empty()) { continue; } // runtime trie does not match an empty word
++        auto word = until_codepoints(text);
++        alphabet.insert(word.begin(), word.end());
++        for (size_t size = 1; size < word.size(); ++size) {
++            prefixes.emplace(word.begin(), word.begin() + size);
++        }
++        delimiters.push_back(std::move(word));
++    }
++    if (delimiters.empty()) { return ".*"; }
++
++    const std::vector<until_word> states(prefixes.begin(), prefixes.end());
++    auto rule_name = [&](size_t state) { return base + "-q" + std::to_string(state); };
++    std::string excluded;
++    for (auto cp : alphabet) { excluded += gbnf_escape_char_class(cp); }
++    const auto continuation = close ? until_codepoints(*close) : until_word{};
++
++    for (size_t i = 0; i < states.size(); ++i) {
++        std::vector<std::string> alternatives;
++        bool may_close = true;
++        if (close) {
++            auto joined = states[i];
++            joined.insert(joined.end(), continuation.begin(), continuation.end());
++            // No delimiter may start in the value and complete across its end.
++            for (size_t start = 0; start < states[i].size() && may_close; ++start) {
++                for (const auto & delimiter : delimiters) {
++                    if (start + delimiter.size() <= joined.size() &&
++                        std::equal(delimiter.begin(), delimiter.end(), joined.begin() + start)) {
++                        may_close = false;
++                        break;
++                    }
++                }
++            }
++        }
++        if (may_close) { alternatives.push_back(close ? gbnf_format_literal(*close) : "\"\""); }
++
++        std::map<size_t, std::vector<uint32_t>> transitions;
++        for (auto cp : alphabet) {
++            auto next = states[i];
++            next.push_back(cp);
++            if (std::any_of(delimiters.begin(), delimiters.end(), [&](const until_word & delimiter) {
++                    return until_has_suffix(next, delimiter);
++                })) { continue; }
++            size_t target = 0, longest = 0;
++            for (size_t j = 1; j < states.size(); ++j) {
++                if (states[j].size() > longest && until_has_suffix(next, states[j])) {
++                    target = j;
++                    longest = states[j].size();
++                }
++            }
++            transitions[target].push_back(cp);
++        }
++        for (const auto & [target, chars] : transitions) {
++            std::string cls;
++            for (auto cp : chars) { cls += gbnf_escape_char_class(cp); }
++            alternatives.push_back("[" + cls + "] " + rule_name(target));
++        }
++        alternatives.push_back("[^" + excluded + "] " + rule_name(0));
++        const auto name = rule_name(i);
++        if (builder.add_rule(name, string_join(alternatives, " | ")) != name) {
++            throw std::runtime_error("PEG delimiter grammar rule name collision");
++        }
++    }
++    return rule_name(0);
++}
++
+ 
+ static std::unordered_set<std::string> collect_reachable_rules(
+     const common_peg_arena & arena,
+@@ -1658,6 +1723,27 @@
+         }
+     };
+ 
++    // A rule/tag is transparent for matching an until + literal pair, while
++    // its PEG AST and serialized parser stay unchanged.
++    auto terminal_parser = [&](common_peg_parser_id id) -> const common_peg_parser_variant & {
++        while (true) {
++            const auto & parser = effective_parser(id);
++            if (const auto * rule = std::get_if<common_peg_rule_parser>(&parser)) { id = rule->child; }
++            else if (const auto * ref = std::get_if<common_peg_ref_parser>(&parser)) { id = get_rule(ref->name); }
++            else { return parser; }
++        }
++    };
++    std::map<std::string, std::string> until_grammars;
++    auto until_grammar = [&](const common_peg_until_parser & until, const std::string * close = nullptr) {
++        const auto key = nlohmann::json({until.delimiters, close ? nlohmann::json(*close) : nlohmann::json(nullptr)}).dump();
++        auto found = until_grammars.find(key);
++        if (found != until_grammars.end()) { return found->second; }
++        const auto rule = gbnf_excluding_pattern(until.delimiters, builder,
++            "aitable-until-" + std::to_string(until_grammars.size()), close);
++        until_grammars.emplace(key, rule);
++        return rule;
++    };
++
+     // Generate GBNF for a parser
+     std::function<std::string(common_peg_parser_id)> to_gbnf = [&](common_peg_parser_id id) -> std::string {
+         const auto & parser = parsers_.at(id);
+@@ -1673,8 +1759,22 @@
+                 return gbnf_format_literal(p.literal);
+             } else if constexpr (std::is_same_v<T, common_peg_sequence_parser>) {
+                 std::string s;
+-                for (const auto & child : p.children) {
+-                    auto child_gbnf = to_gbnf(child);
++                for (size_t i = 0; i < p.children.size(); ++i) {
++                    const auto child = p.children[i];
++                    std::string child_gbnf;
++                    bool joined = false;
++                    if (i + 1 < p.children.size()) {
++                        const auto * until = std::get_if<common_peg_until_parser>(&terminal_parser(child));
++                        const auto * close = std::get_if<common_peg_literal_parser>(&terminal_parser(p.children[i + 1]));
++                        if (until && close && until_can_join(until->delimiters, close->literal)) {
++                            // Compile the terminator with the scan: a value suffix
++                            // must not create an earlier delimiter across this join.
++                            child_gbnf = until_grammar(*until, &close->literal);
++                            joined = true;
++                            ++i;
++                        }
++                    }
++                    if (!joined) { child_gbnf = to_gbnf(child); }
+                     if (child_gbnf.empty()) {
+                         continue;
+                     }
+@@ -1682,8 +1782,8 @@
+                         s += " ";
+                     }
+                     const auto & child_parser = effective_parser(child);
+-                    if (std::holds_alternative<common_peg_choice_parser>(child_parser) ||
+-                        std::holds_alternative<common_peg_sequence_parser>(child_parser)) {
++                    if (!joined && (std::holds_alternative<common_peg_choice_parser>(child_parser) ||
++                        std::holds_alternative<common_peg_sequence_parser>(child_parser))) {
+                         s += "(" + child_gbnf + ")";
+                     } else {
+                         s += child_gbnf;
+@@ -1765,7 +1865,7 @@
+                 if (p.delimiters.empty()) {
+                     return ".*";
+                 }
+-                return gbnf_excluding_pattern(p.delimiters);
++                return until_grammar(p);
+             } else if constexpr (std::is_same_v<T, common_peg_schema_parser>) {
+                 if (schema_delegates(p)) {
+                     return to_gbnf(p.child);
+--- a/common/sampling.cpp
++++ b/common/sampling.cpp
+@@ -1,4 +1,5 @@
+ #include "sampling.h"
++#include "nlohmann/json.hpp"
+ 
+ #include "common.h"
+ #include "fit.h"
+@@ -109,6 +110,17 @@
+     std::vector<T> data;
+ };
+ 
++struct sampling_diagnostic_state {
++    uint64_t sample_steps = 0, grammar_mask_calls = 0, grammar_single_checks = 0;
++    uint64_t skip_no_grammar = 0, skip_reasoning = 0, grammar_accepts = 0, generated_accepts = 0;
++    uint64_t grammar_checked_steps = 0, skip_sampler_error = 0, skip_other = 0;
++    uint64_t invalid_selections = 0;
++    size_t prefill_tokens = 0;
++    bool prefill_applied = false;
++    std::string failure_reason;
++    std::vector<llama_token> eog_probes;
++};
++
+ struct common_sampler {
+     common_params_sampling params;
+ 
+@@ -121,6 +133,7 @@
+     std::vector<llama_token_data> cur;
+ 
+     llama_token_data_array cur_p;
++    sampling_diagnostic_state diagnostic;
+ 
+     void reset() {
+         prev.clear();
+@@ -409,6 +422,13 @@
+         /* .cur_p   = */ {},
+     };
+ 
++    result->diagnostic.prefill_tokens = prefill_tokens.size();
++    result->diagnostic.prefill_applied = grmr && !params.grammar_lazy && common_grammar_needs_prefill(params.grammar);
++    for (const auto id : {llama_vocab_eos(vocab), llama_vocab_eot(vocab)}) {
++        if (id >= 0 && std::find(result->diagnostic.eog_probes.begin(), result->diagnostic.eog_probes.end(), id) == result->diagnostic.eog_probes.end()) {
++            result->diagnostic.eog_probes.push_back(id);
++        }
++    }
+     return result;
+ }
+ 
+@@ -448,6 +468,8 @@
+ 
+     // grammar_should_apply() checks the reasoning budget state, so calculate this before we accept
+     const auto accept_grammar = is_generated && grammar_should_apply(gsmpl);
++    if (is_generated) { ++gsmpl->diagnostic.generated_accepts; }
++    if (accept_grammar) { ++gsmpl->diagnostic.grammar_accepts; }
+ 
+     if (gsmpl->rbudget && is_generated) {
+         llama_sampler_accept(gsmpl->rbudget, token);
+@@ -479,6 +501,7 @@
+         /* .prev    = */ gsmpl->prev,
+         /* .cur     = */ gsmpl->cur,
+         /* .cur_p   = */ gsmpl->cur_p,
++        /* .diagnostic = */ gsmpl->diagnostic,
+     };
+ }
+ 
+@@ -536,41 +559,66 @@
+     return gsmpl->chain;
+ }
+ 
+-static llama_token common_sampler_validate_or_select_fallback(
+-        const char *                  func,
+-        struct llama_context *        ctx,
+-        llama_token_data_array &      cur_p,
+-        llama_token                   id) {
+-    const llama_vocab * vocab = llama_model_get_vocab(llama_get_model(ctx));
+-    const int32_t n_vocab = llama_vocab_n_tokens(vocab);
+-
+-    if (id >= 0 && id < n_vocab) {
++static llama_token common_sampler_validate_selected(
++        const char *, struct llama_context * ctx, llama_token_data_array & cur_p, llama_token id, common_sampler * gsmpl) {
++    const int32_t n_vocab = llama_vocab_n_tokens(llama_model_get_vocab(llama_get_model(ctx)));
++    const bool selected_in_range = cur_p.data != nullptr && cur_p.selected >= 0 &&
++                                   static_cast<size_t>(cur_p.selected) < cur_p.size;
++    // Normal selection remains O(1). Masked alternatives are expected and do not
++    // invalidate a finite selected token. Failure classification alone scans candidates.
++    if (selected_in_range && id >= 0 && id < n_vocab &&
++            cur_p.data[cur_p.selected].id == id && std::isfinite(cur_p.data[cur_p.selected].logit)) {
+         return id;
+     }
+-
+     static std::atomic<int> n_warn_invalid_sample { 0 };
+     const int n_warn = n_warn_invalid_sample.fetch_add(1, std::memory_order_relaxed);
+     if (n_warn < 16) {
+-        LOG_WRN("%s: sampler selected invalid token %d (n_vocab=%d); selecting a valid fallback candidate\n",
+-                func, id, n_vocab);
+-    } else {
+-        LOG_DBG("%s: sampler selected invalid token %d (n_vocab=%d); selecting a valid fallback candidate\n",
+-                func, id, n_vocab);
+-    }
+-
+-    for (size_t i = 0; i < cur_p.size; ++i) {
+-        const llama_token id_fallback = cur_p.data[i].id;
+-        if (id_fallback >= 0 && id_fallback < n_vocab) {
+-            cur_p.selected = i;
+-            return id_fallback;
+-        }
+-    }
+-
+-    cur_p.selected = -1;
+-    return LLAMA_TOKEN_NULL;
++        LOG_WRN("sampling_error: refusing invalid selection (token=%d, n_vocab=%d)\n", id, n_vocab);
++    }
++    auto fail = [&](const char * reason) -> llama_token {
++        ++gsmpl->diagnostic.invalid_selections;
++        gsmpl->diagnostic.failure_reason = reason;
++        throw common_sampler_error(reason);
++    };
++    bool available = false;
++    if (cur_p.data != nullptr) {
++        for (size_t i = 0; i < cur_p.size; ++i) {
++            const auto & item = cur_p.data[i];
++            if (std::isnan(item.logit) || item.logit == INFINITY) {
++                return fail("nonfinite_candidate_logits");
++            }
++            available |= item.id >= 0 && item.id < n_vocab && std::isfinite(item.logit);
++        }
++    }
++    if (!available) { return fail("no_available_token"); }
++    if (!selected_in_range) { return fail("invalid_selected_index"); }
++    if (id < 0 || id >= n_vocab) { return fail("invalid_selected_token"); }
++    if (cur_p.data[cur_p.selected].id != id) { return fail("selected_token_mismatch"); }
++    return fail("masked_selected_token");
++}
++
++static llama_token common_sampler_selected_id(const llama_token_data_array & cur_p) {
++    if (cur_p.data == nullptr || cur_p.selected < 0 || static_cast<size_t>(cur_p.selected) >= cur_p.size) {
++        return LLAMA_TOKEN_NULL;
++    }
++    return cur_p.data[cur_p.selected].id;
+ }
+ 
+ llama_token common_sampler_sample(struct common_sampler * gsmpl, struct llama_context * ctx, int idx, bool grammar_first) {
++    ++gsmpl->diagnostic.sample_steps;
++    struct diagnostic_step_scope {
++        common_sampler * smpl;
++        uint64_t calls, failures;
++        ~diagnostic_step_scope() {
++            auto & d = smpl->diagnostic;
++            if (d.grammar_mask_calls + d.grammar_single_checks != calls) { ++d.grammar_checked_steps; }
++            else if (!smpl->grmr) { ++d.skip_no_grammar; }
++            else if (d.invalid_selections != failures) { ++d.skip_sampler_error; }
++            else if (!grammar_should_apply(smpl)) { ++d.skip_reasoning; }
++            else { ++d.skip_other; }
++        }
++    } diagnostic_scope {gsmpl,gsmpl->diagnostic.grammar_mask_calls + gsmpl->diagnostic.grammar_single_checks,
++                        gsmpl->diagnostic.invalid_selections};
+     llama_synchronize(ctx);
+ 
+     // start measuring sampling time after the llama_context synchronization in order to not measure any ongoing async operations
+@@ -606,22 +654,8 @@
+                 }
+             }
+ 
+-            id = common_sampler_validate_or_select_fallback(__func__, ctx, cur_p, id);
+-            if (id == id_backend) {
+-                return id;
+-            }
+-
+-            static std::atomic<int> n_warn_backend_sample_invalid { 0 };
+-            const int n_warn = n_warn_backend_sample_invalid.fetch_add(1, std::memory_order_relaxed);
+-            if (n_warn < 16) {
+-                LOG_WRN("%s: backend sampler returned a token outside the current candidates (id=%d, found=%d); falling back to CPU sampler\n",
+-                        __func__, id_backend, (int) found);
+-            } else {
+-                LOG_DBG("%s: backend sampler returned a token outside the current candidates (id=%d, found=%d); falling back to CPU sampler\n",
+-                        __func__, id_backend, (int) found);
+-            }
+-            cur_p.selected = -1;
+-            gsmpl->set_logits(ctx, idx, false);
++            return common_sampler_validate_selected(__func__, ctx, cur_p, id, gsmpl);
++
+         }
+     }
+ 
+@@ -629,30 +663,14 @@
+     llama_sampler_apply(rbudget, &cur_p);
+ 
+     if (grammar_first && grammar_should_apply(gsmpl)) {
++        ++gsmpl->diagnostic.grammar_mask_calls;
+         llama_sampler_apply(grmr, &cur_p);
+     }
+ 
+     llama_sampler_apply(chain, &cur_p);
+ 
+-    id = cur_p.data[cur_p.selected].id;
+-    id = common_sampler_validate_or_select_fallback(__func__, ctx, cur_p, id);
+-    if (id == LLAMA_TOKEN_NULL) {
+-        gsmpl->set_logits(ctx, idx, false);
+-
+-        llama_sampler_apply(rbudget, &cur_p);
+-
+-        if (grammar_first && grammar_should_apply(gsmpl)) {
+-            llama_sampler_apply(grmr, &cur_p);
+-        }
+-
+-        llama_sampler_apply(chain, &cur_p);
+-
+-        GGML_ASSERT(cur_p.selected != -1 && "no selected token during sampling - check your sampling configuration");
+-
+-        id = cur_p.data[cur_p.selected].id;
+-        id = common_sampler_validate_or_select_fallback(__func__, ctx, cur_p, id);
+-        return id;
+-    }
++    id = common_sampler_selected_id(cur_p);
++    id = common_sampler_validate_selected(__func__, ctx, cur_p, id, gsmpl);
+ 
+     if (grammar_first || !grammar_should_apply(gsmpl)) {
+         return id;
+@@ -663,6 +681,7 @@
+         llama_token_data       single_token_data       = { id, 1.0f, 0.0f };
+         llama_token_data_array single_token_data_array = { &single_token_data, 1, -1, false };
+ 
++        ++gsmpl->diagnostic.grammar_single_checks;
+         llama_sampler_apply(grmr, &single_token_data_array);
+ 
+         const bool is_valid = single_token_data_array.data[0].logit != -INFINITY;
+@@ -678,6 +697,7 @@
+     llama_sampler_apply(rbudget,  &cur_p);
+ 
+     if (grammar_should_apply(gsmpl)) {
++        ++gsmpl->diagnostic.grammar_mask_calls;
+         llama_sampler_apply(grmr,  &cur_p);
+     }
+ 
+@@ -685,8 +705,8 @@
+ 
+     GGML_ASSERT(cur_p.selected != -1 && "no selected token during sampling - check your sampling configuration");
+ 
+-    id = cur_p.data[cur_p.selected].id;
+-    id = common_sampler_validate_or_select_fallback(__func__, ctx, cur_p, id);
++    id = common_sampler_selected_id(cur_p);
++    id = common_sampler_validate_selected(__func__, ctx, cur_p, id, gsmpl);
+ 
+     return id;
+ }
+@@ -965,3 +985,56 @@
+ 
+     return samplers;
+ }
++
++std::string common_sampler_diagnostics_json(struct common_sampler * gsmpl) {
++    using json = nlohmann::ordered_json;
++    if (!gsmpl) { return json({{"schema","aitable.native_sampling@1"},{"sampler_attached",false}}).dump(); }
++    const auto & d = gsmpl->diagnostic;
++    const auto & grammar = common_grammar_value(gsmpl->params.grammar);
++    uint64_t hash = 14695981039346656037ull;
++    for (unsigned char c : grammar) { hash = (hash ^ c) * 1099511628211ull; }
++    char hash_text[17]; snprintf(hash_text,sizeof(hash_text),"%016llx",static_cast<unsigned long long>(hash));
++    json out = {{"schema","aitable.native_sampling@1"},{"sampler_attached",true},
++        {"grammar_configured",!grammar.empty()},{"grammar_attached",gsmpl->grmr!=nullptr},
++        {"grammar_lazy",gsmpl->params.grammar_lazy},{"grammar_hash_alg","fnv1a64"},{"grammar_hash",hash_text},
++        {"generation_prompt_bytes",gsmpl->params.generation_prompt.size()},{"prefill_tokens",d.prefill_tokens},
++        {"prefill_applied",d.prefill_applied},{"ignore_eos",gsmpl->params.ignore_eos},
++        {"sample_steps",d.sample_steps},{"grammar_full_mask_calls",d.grammar_mask_calls},
++        {"grammar_single_checks",d.grammar_single_checks},{"skip_no_grammar_steps",d.skip_no_grammar},
++        {"skip_reasoning_steps",d.skip_reasoning},{"generated_accepts",d.generated_accepts},
++        {"grammar_checked_steps",d.grammar_checked_steps},{"skip_sampler_error_steps",d.skip_sampler_error},
++        {"skip_other_steps",d.skip_other},
++        {"grammar_accepts",d.grammar_accepts},{"invalid_selections",d.invalid_selections},
++        {"failure_reason",d.failure_reason.empty()?json(nullptr):json(d.failure_reason)},
++        {"grammar_should_apply_now",grammar_should_apply(gsmpl)},
++        {"eog_probe_scope","eos_eot_only"},{"eog_probe_count",d.eog_probes.size()},
++        {"grammar_allows_probe_eog",nullptr},{"last_chain_probe_eog_present",nullptr},
++        {"last_chain_probe_eog_finite",nullptr},{"lazy_trigger_state","not_observed"}};
++    if (!d.eog_probes.empty()) {
++        bool present=false,finite=false;
++        for (size_t i=0; gsmpl->cur_p.data && i<gsmpl->cur_p.size; ++i) {
++            for (auto id:d.eog_probes) if (gsmpl->cur_p.data[i].id==id) {
++                present=true;finite|=std::isfinite(gsmpl->cur_p.data[i].logit);
++            }
++        }
++        out["last_chain_probe_eog_present"]=present;
++        out["last_chain_probe_eog_finite"]=present?json(finite):json(nullptr);
++        if (gsmpl->grmr) {
++            std::vector<llama_token_data> probe;
++            for (auto id:d.eog_probes) probe.push_back({id,0.0f,0.0f});
++            llama_token_data_array a{probe.data(),probe.size(),-1,false};
++            llama_sampler_apply(gsmpl->grmr,&a);
++            bool allowed=false;for(const auto & item:probe) allowed|=std::isfinite(item.logit);
++            out["grammar_allows_probe_eog"]=allowed;
++        }
++    }
++    out["diagnostic_saturated"]=false;
++    for (const char * key : {"generation_prompt_bytes","prefill_tokens","sample_steps","grammar_full_mask_calls",
++            "grammar_single_checks","skip_no_grammar_steps","skip_reasoning_steps","generated_accepts",
++            "grammar_accepts","invalid_selections","grammar_checked_steps","skip_sampler_error_steps","skip_other_steps"}) {
++        if (out[key].get<uint64_t>() > 2147483647u) {
++            out[key]=2147483647u;out["diagnostic_saturated"]=true;
++        }
++    }
++    return out.dump();
++}
+--- a/common/sampling.h
++++ b/common/sampling.h
+@@ -5,6 +5,7 @@
+ #include "common.h"
+ 
+ #include <string>
++#include <stdexcept>
+ #include <vector>
+ 
+ // common_sampler extends llama_sampler with additional functionality:
+@@ -31,6 +32,14 @@
+ //
+ // TODO: measure grammar performance
+ //
++
++// Sampling failure is terminal for this generation. The finite code contains no model text.
++class common_sampler_error : public std::runtime_error {
++public:
++    explicit common_sampler_error(const char * reason)
++        : std::runtime_error(std::string("sampling_error: ") + reason), code(reason) {}
++    const std::string code;
++};
+ 
+ struct common_sampler;
+ 
+@@ -121,3 +130,6 @@
+ };
+ 
+ typedef std::unique_ptr<common_sampler, common_sampler_deleter> common_sampler_ptr;
++
++// Finite, no-text snapshot. Called only when a generation finishes.
++std::string common_sampler_diagnostics_json(struct common_sampler * gsmpl);
+--- a/ggml/src/ggml-cuda/fattn-common.cuh
++++ b/ggml/src/ggml-cuda/fattn-common.cuh
+@@ -1719,7 +1719,7 @@
+         const int tiles_nwaves = (ntiles_dst + max_blocks - 1) / max_blocks;
+         const int tiles_efficiency_percent = 100 * ntiles_dst / (max_blocks*tiles_nwaves);
+ 
+-        const bool use_stream_k = cc >= GGML_CUDA_CC_ADA_LOVELACE || amd_wmma_available(cc) || tiles_efficiency_percent < 75;
++        const bool use_stream_k = (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_ADA_LOVELACE) || (amd_wmma_available(cc) && Q->ne[0] == 64) || tiles_efficiency_percent < 75;
+ 
+         blocks_num.x = ntiles_dst;
+         blocks_num.y = 1;
+--- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh
++++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
+@@ -157,8 +157,8 @@
+ 
+     GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256,  8,  64, 2,  32, 128, 128, 128, 1, true);
+     GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 16,  64, 2,  32, 128, 128, 128, 1, true);
+-    GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 128, 2,  64, 128, 128,  64, 1, true);
+-    GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 128, 2,  64, 128, 128,  64, 1, true);
++    GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 256, 2,  64, 128, 128,  64, 1, true);
++    GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 256, 2,  64, 128, 128,  64, 1, true);
+ 
+     GGML_CUDA_FATTN_MMA_CONFIG_CASE(320, 256, 32, 128, 2,  32, 160, 128, 128, 1, true);
+     GGML_CUDA_FATTN_MMA_CONFIG_CASE(320, 256, 64, 128, 2,  32, 160, 128, 128, 1, true);
+@@ -359,6 +359,39 @@
+ }
+ 
+ // ------------------------------------------------------------------------------------------------------------------
++
++
++// F16 global V -> transposed LDS [D, K]. Only the guarded RDNA4 path calls this.
++// The original temporary dequantization and tensor strides remain unchanged.
++template<int KB, int DB, int NW, int LDS_STRIDE>
++static __device__ __forceinline__ void fpx_rdna4_load_v_transposed(
++        const half2 * src, half2 * dst, const int stride_src) {
++#if defined(__gfx1201__)
++    using half8 = __fp16 __attribute__((ext_vector_type(8)));
++    constexpr int stride_dst = LDS_STRIDE; // half2 units; caller proves pool capacity
++    constexpr int dtiles = (DB + 31)/32;
++    const int lane = threadIdx.x % 32;
++    for (int tile = threadIdx.y; tile < (KB/16)*dtiles; tile += NW) {
++        const int kd = (tile/dtiles)*16;
++        const int dd = (tile%dtiles)*32;
++#pragma unroll
++        for (int part = 0; part < 2; ++part) {
++            const half * address = (const half *)(src + (kd+part*8+lane/4)*stride_src) + dd+(lane%4)*8;
++            const half8 values = __builtin_amdgcn_global_load_tr_b128_v8f16(
++                (__attribute__((address_space(1))) half8 *)address);
++#pragma unroll
++            for (int e=0;e<8;++e) {
++                const int d = dd+(e%4)*8+lane%8;
++                const int k = kd+part*8+(lane/8)*2+e/4;
++                ((half *)dst)[d*(2*stride_dst)+k] = ((const half *)&values)[e];
++            }
++        }
++    }
++#else
++    GGML_UNUSED_VARS(src,dst,stride_src);
++    NO_DEVICE_CODE;
++#endif
++}
+ 
+ template<int stride_tile, int nwarps, int nbatch_fa, bool use_cp_async, bool oob_check>
+ static __device__ __forceinline__ void flash_attn_ext_f16_load_tile(
+@@ -949,6 +982,19 @@
+     }
+ 
+ 
++
++    // No extra LDS allocation: the transposed tile must fit the existing K/V pool.
++    constexpr int fpx_tr_capacity = nbatch_fa*(stride_tile_K>stride_tile_V ? stride_tile_K : stride_tile_V)/(2*nbatch_V2);
++    constexpr int fpx_tr_stride = fpx_tr_capacity < nbatch_fa/2+4 ? fpx_tr_capacity : nbatch_fa/2+4;
++#if defined(__gfx1201__)
++    constexpr bool fpx_tr_eligible = DKQ==256 && DV==256 && ncols2>1 && nstages==0 && !V_is_K_view && !oob_check &&
++        nbatch_fa%16==0 && nbatch_V2%16==0 &&
++        (2*nbatch_V2)*fpx_tr_stride <= nbatch_fa*(stride_tile_K>stride_tile_V ? stride_tile_K : stride_tile_V);
++    const bool fpx_use_tr_v = fpx_tr_eligible && stride_V%4==0 && (reinterpret_cast<uintptr_t>(V_h2)&15)==0;
++#else
++    constexpr bool fpx_use_tr_v = false;
++#endif
++
+     // Calculate VKQ tile, need to use logical rather than physical elements for i0 due to transposition of V:
+ #pragma unroll
+     for (int i0_start = 0; i0_start < DV; i0_start += 2*nbatch_V2) {
+@@ -959,8 +1005,13 @@
+         if constexpr (nstages <= 1) {
+             if (!V_is_K_view || i0_stop > 2*nbatch_K2) {
+                 constexpr bool use_cp_async = nstages == 1;
++                if (fpx_use_tr_v) {
++                    fpx_rdna4_load_v_transposed<nbatch_fa, 2*nbatch_V2, nwarps, fpx_tr_stride>(
++                        V_h2 + int64_t(k_VKQ_0)*stride_V + i0_start/2, tile_V, stride_V);
++                } else {
+                 flash_attn_ext_f16_load_tile<stride_tile_V, nwarps, nbatch_fa, use_cp_async, oob_check>
+                     (V_h2 + int64_t(k_VKQ_0)*stride_V + i0_start/2, tile_V, i0_diff/2, stride_V, k_VKQ_sup);
++                }
+                 if (use_cp_async) {
+                     cp_async_wait_all();
+                 }
+@@ -978,7 +1029,14 @@
+                 const int k0 = k00 + (threadIdx.y % np)*T_A_VKQ::J;
+ 
+                 T_A_VKQ A; // Transposed in SRAM but not in registers, gets transposed on load.
++                if (fpx_use_tr_v) {
++#pragma unroll
++                    for (int l=0;l<A.ne;++l) {
++                        A.x[l] = tile_V[(i_VKQ_0-i0_start+A.get_i(l))*fpx_tr_stride+k0+A.get_j(l)];
++                    }
++                } else {
+                 load_ldmatrix_trans(A, tile_V_i + 2*k0*stride_tile_V + (i_VKQ_0 - i0_start)/2, stride_tile_V);
++                }
+                 if constexpr (T_B_KQ::I == 8) {
+                     mma(VKQ_C[i_VKQ_0/T_A_VKQ::I], A, B[k00/(np*T_A_VKQ::J)]);
+                 } else {
+@@ -1750,7 +1808,7 @@
+ #endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING
+ 
+ #if defined(AMD_WMMA_AVAILABLE)
+-    if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 128) {
++    if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 256) {
+         NO_DEVICE_CODE;
+         return;
+     }
+--- a/ggml/src/ggml-cuda/fattn.cu
++++ b/ggml/src/ggml-cuda/fattn.cu
+@@ -7,6 +7,52 @@
+ #include "fattn.cuh"
+ #include "convert.cuh"
+ 
++#ifdef GGML_USE_HIP
++extern "C" __declspec(dllimport) float ck_attention_bias(void*,void*,void*,void*,void*,void*,int,int,int,int,int,float,void*,int);
++static bool ggml_try_ck_attention(ggml_backend_cuda_context &ctx, ggml_tensor *dst) {
++    const char *enable = getenv("GGML_HIP_CK_ATTENTION");
++    if (!enable || strcmp(enable,"1")) return false;
++    const auto *q=dst->src[0], *k=dst->src[1], *v=dst->src[2], *m=dst->src[3];
++    float scale,max_bias,softcap;
++    memcpy(&scale,dst->op_params,sizeof(float));
++    memcpy(&max_bias,(const float*)dst->op_params+1,sizeof(float));
++    memcpy(&softcap,(const float*)dst->op_params+2,sizeof(float));
++    // Experimental narrow boundary: prefill only, one sequence, standard additive mask.
++    if(q->type!=GGML_TYPE_F32 || q->ne[0]!=256 || q->ne[1]<64 || q->ne[3]!=1 ||
++       k->ne[0]!=256 || v->ne[0]!=256 || k->ne[3]!=1 || v->ne[3]!=1 ||
++       k->ne[1]!=v->ne[1] || k->ne[2]!=v->ne[2] || k->ne[1]<q->ne[1] ||
++       q->ne[2]%k->ne[2] || dst->src[4] || max_bias!=0 || softcap!=0 ||
++       !ggml_is_contiguous(dst) || dst->type!=GGML_TYPE_F32) return false;
++    if(k->type!=GGML_TYPE_Q8_0) return false;
++    if(v->type!=GGML_TYPE_Q8_0) return false;
++    if(m && (m->type!=GGML_TYPE_F16 || m->ne[2]!=1 || m->ne[3]!=1 ||
++       m->nb[0]!=sizeof(half) || m->ne[0]<k->ne[1] || m->ne[1]<q->ne[1])) return false;
++    for(auto *t:{q,k,v}) if(t->nb[0]!=ggml_type_size(t->type)) return false;
++    ggml_cuda_pool_alloc<half> qp(ctx.pool()),kp(ctx.pool()),vp(ctx.pool()),op(ctx.pool());
++    ggml_cuda_pool_alloc<float> lse(ctx.pool());
++    auto stream=ctx.stream();
++    auto pack=[&](const ggml_tensor *t,ggml_cuda_pool_alloc<half>&buf) {
++        buf.alloc(ggml_nelements(t));
++        auto convert=ggml_get_to_fp16_nc_cuda(t->type);
++        GGML_ASSERT(convert);
++        const size_t ts=ggml_type_size(t->type);
++        convert(t->data,buf.ptr,t->ne[0],t->ne[1],t->ne[2],t->ne[3],t->nb[1]/ts,t->nb[2]/ts,t->nb[3]/ts,stream);
++    };
++    pack(q,qp);pack(k,kp);pack(v,vp);
++    op.alloc(ggml_nelements(dst));lse.alloc(q->ne[1]*q->ne[2]);
++    static int reports=0;
++    int log=reports<8?1:0;
++    if(log) {fprintf(stderr,"\nCK_ATTENTION_DISPATCH q=%lld k=%lld h=%lld hk=%lld mask=%d\n",(long long)q->ne[1],(long long)k->ne[1],(long long)q->ne[2],(long long)k->ne[2],m!=nullptr);++reports;}
++    float ret=ck_attention_bias(qp.ptr,kp.ptr,vp.ptr,op.ptr,lse.ptr,m?m->data:nullptr,
++        q->ne[1],k->ne[1],q->ne[2],k->ne[2],m?m->nb[1]/sizeof(half):0,scale,stream,0);
++    if(ret<0) GGML_ABORT("CK attention rejected eligible FPX tensor: %f",ret);
++    CUDA_CHECK(cudaGetLastError());
++    ggml_get_to_fp32_cuda(GGML_TYPE_F16)(op.ptr,(float*)dst->data,ggml_nelements(dst),stream);
++    return true;
++}
++#endif
++
++
+ template <int DKQ, int DV, int ncols2>
+ static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+     const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
+@@ -90,6 +136,12 @@
+         }
+     }
+ 
++    // PR28102: RDNA GQA divisibility routing, preserving Volta above.
++    if (amd_wmma_available(cc)) {
++        if (use_gqa_opt && gqa_ratio % 8 == 0) { ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 8>(ctx,dst); return; }
++        if (use_gqa_opt && gqa_ratio % 4 == 0) { ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 4>(ctx,dst); return; }
++        if (use_gqa_opt && gqa_ratio % 2 == 0) { ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 2>(ctx,dst); return; }
++    }
+     if (use_gqa_opt && gqa_ratio > 4) {
+         ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 8>(ctx, dst);
+         return;
+@@ -682,6 +734,12 @@
+         return BEST_FATTN_KERNEL_MMA_F16;
+     }
+ 
++    // PR28102 compatibility: this older fork has a legacy rocWMMA early return.
++    // Let only the newly supported RDNA head256/GQA shapes reach native MMA first.
++    if (amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] == 256 && Q->ne[1] * gqa_ratio_eff > 16) {
++        return BEST_FATTN_KERNEL_MMA_F16;
++    }
++
+     // Use the WMMA kernel if possible:
+     if (ggml_cuda_should_use_wmma_fattn(cc) && K->ne[1] % FATTN_KQ_STRIDE == 0 && Q->ne[0] != 40 && Q->ne[0] != 72 && Q->ne[0] != 192 && Q->ne[0] != 512 && Q->ne[0] != 576) {
+         if (can_use_vector_kernel && Q->ne[1] <= 2) {
+@@ -704,7 +762,7 @@
+     }
+ 
+     // AMD WMMA is always faster than the tile kernel if the full tile width of 16 can be utilized.
+-    if ((amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] <= 128) && Q->ne[0] != 40 && Q->ne[0] != 72 && Q->ne[1] * gqa_ratio_eff > 8) {
++    if ((amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] <= 256) && Q->ne[0] != 40 && Q->ne[0] != 72 && Q->ne[1] * gqa_ratio_eff > (Q->ne[0] <= 128 ? 8 : 16)) {
+         return BEST_FATTN_KERNEL_MMA_F16;
+     }
+ 
+@@ -761,7 +819,18 @@
+ 
+ void ggml_cuda_flash_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+     ggml_cuda_set_device(ctx.device);
+-
++#ifdef GGML_USE_HIP
++    if (ggml_try_ck_attention(ctx,dst)) return;
++#endif
++
++#ifdef GGML_USE_HIP
++    static int dispatch_reports = 0;
++    if (dispatch_reports < 8) {
++        const auto selected = ggml_cuda_get_best_fattn_kernel(ggml_cuda_get_device(), dst);
++        fprintf(stderr,"FA28102_DISPATCH kernel=%d mma=%d q=%lld k=%lld head=%lld\n",int(selected),int(selected==BEST_FATTN_KERNEL_MMA_F16),(long long)dst->src[0]->ne[1],(long long)dst->src[1]->ne[1],(long long)dst->src[0]->ne[0]);
++        ++dispatch_reports;
++    }
++#endif
+     const ggml_tensor * K = dst->src[1];
+     const ggml_tensor * V = dst->src[2];
+ 
+--- a/ggml/src/ggml-hip/CMakeLists.txt
++++ b/ggml/src/ggml-hip/CMakeLists.txt
+@@ -180,3 +180,10 @@
+ # argmax on RDNA3.5, making MTP speculative decode diverge from the
+ # non-speculative baseline. gfx1151 is RDNA3.5, and -ffinite-math-only also
+ # breaks inf handling outright (EXPM1 returned NaN where the CPU returns inf).
++
++# CK forward bridge with explicit supported-shape dispatch.
++set(FPX_CK_IMPORT_LIBRARY "" CACHE FILEPATH "Import library for the Windows CK attention bridge")
++if (NOT EXISTS "${FPX_CK_IMPORT_LIBRARY}")
++    message(FATAL_ERROR "Set FPX_CK_IMPORT_LIBRARY to the matching ck-attention.lib from the AiTable fix bundle")
++endif()
++target_link_libraries(ggml-hip PRIVATE "${FPX_CK_IMPORT_LIBRARY}")
+--- a/src/llama-context.cpp
++++ b/src/llama-context.cpp
+@@ -2984,7 +2984,7 @@
+                 return false;
+             }
+ 
+-            for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
++            for (size_t i = 0; i < mbuf.org.size(); ++i) { // cloned storage owns cpy, not source views
+                 ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf.org[i]);
+             }
+         }
+--- a/tools/cli/CMakeLists.txt
++++ b/tools/cli/CMakeLists.txt
+@@ -8,3 +8,7 @@
+ if(LLAMA_TOOLS_INSTALL)
+     install(TARGETS ${TARGET} RUNTIME)
+ endif()
++
++add_executable(prefix-offline-probe prefix-offline-probe.cpp)
++target_link_libraries(prefix-offline-probe PRIVATE server-context PUBLIC llama-common ${CMAKE_THREAD_LIBS_INIT})
++target_compile_features(prefix-offline-probe PRIVATE cxx_std_17)
+--- /dev/null
++++ b/tools/cli/prefix-offline-probe.cpp
+@@ -0,0 +1,58 @@
++#include "arg.h"
++#include "common.h"
++#include "server-context.h"
++#include "server-task.h"
++#include <chrono>
++#include <fstream>
++#include <thread>
++#include <iostream>
++#include <exception>
++int main(int argc,char **argv) {
++ try {
++  if(argc<4)throw std::invalid_argument("manifest output then native model arguments required");
++  json cases;std::ifstream in(argv[1]);if(!in)throw std::runtime_error("manifest open failed");in>>cases;
++  std::ofstream out(argv[2]);if(!out)throw std::runtime_error("output open failed");
++  std::vector<char*> args{argv[0]};for(int i=3;i<argc;++i)args.push_back(argv[i]);
++  common_init();common_params params;
++  if(!common_params_parse((int)args.size(),args.data(),params,LLAMA_EXAMPLE_SERVER))return 2;
++  if(params.n_parallel!=1 || params.n_ctx_checkpoints!=4 || params.checkpoint_every_nt!=2048 || params.cache_ram_mib!=0)throw std::invalid_argument("probe requires one slot, four checkpoints, interval2048, no RAM cache");
++  llama_backend_init();
++  server_context ctx;ctx.set_projector_idle_seconds(0);
++  if(!ctx.load_model(params))throw std::runtime_error("model load failed");
++  std::exception_ptr worker_error;
++  std::thread worker([&]{try{ctx.start_loop();}catch(...){worker_error=std::current_exception();ctx.terminate();}});
++  try {
++   for(const auto & c:cases) {
++    const std::string path=c.at("prompt");std::ifstream f(path,std::ios::binary);if(!f)throw std::runtime_error("prompt open failed");
++    const std::string prompt((std::istreambuf_iterator<char>(f)),{});
++    auto rd=ctx.get_response_reader();server_task task(SERVER_TASK_TYPE_COMPLETION);
++    task.id=rd.get_new_id();task.id_slot=0;task.index=0;task.cli=true;task.cli_prompt=prompt;
++    task.params.sampling=params.sampling;task.params.n_predict=c.value("n_predict",64);
++    task.params.stream=true;task.params.return_tokens=true;task.params.timings_per_token=true;
++    task.params.cache_prompt=c.value("cache_prompt",true);
++    auto start=std::chrono::steady_clock::now();double first=-1;bool final_seen=false;std::string collected;llama_tokens collected_tokens;
++    rd.post_task({std::move(task)});
++    auto elapsed=[&]{return std::chrono::duration<double,std::milli>(std::chrono::steady_clock::now()-start).count();};
++    auto cancel=[&]{return elapsed()>60000;};
++    while(rd.has_next()) {
++     auto result=rd.next(cancel);if(!result)throw std::runtime_error("cancelled or missing result");
++     if(result->is_error())throw std::runtime_error(result->to_json().dump());
++     if(auto partial=dynamic_cast<server_task_result_cmpl_partial*>(result.get())) {
++      if(first<0 && (!partial->content.empty() || !partial->tokens.empty()))first=elapsed();
++      if(!partial->is_progress){collected+=partial->content;collected_tokens.insert(collected_tokens.end(),partial->tokens.begin(),partial->tokens.end());}
++     }
++     if(auto fin=dynamic_cast<server_task_result_cmpl_final*>(result.get())) {
++      json row={{"name",c.at("name")},{"timings",fin->timings.to_json()},{"first_native_chunk_ms",first},
++       {"wall_ms",elapsed()},{"stop_type",(int)fin->stop},{"truncated",fin->truncated},
++       {"prompt_tokens",fin->n_prompt_tokens},{"cache_tokens",fin->n_prompt_tokens_cache},
++       {"content",collected},{"tokens",collected_tokens},{"stream_final_content",fin->content},{"stream_final_tokens",fin->tokens},{"native",fin->to_json()}};
++      out<<row.dump()<<"\n";out.flush();if(!out)throw std::runtime_error("result write failed");final_seen=true;break;
++     }
++    }
++    if(!final_seen)throw std::runtime_error("missing native final result");
++   }
++  }catch(...){ctx.terminate();worker.join();if(worker_error)std::rethrow_exception(worker_error);throw;}
++  ctx.terminate();worker.join();if(worker_error)std::rethrow_exception(worker_error);
++ }catch(const std::exception&e){std::cerr<<e.what()<<"\n";return 1;}
++ return 0;
++}
+--- /dev/null
++++ b/tools/server/prefix-short-policy.h
+@@ -0,0 +1,21 @@
++#pragma once
++#include <cstdint>
++#include <cstdlib>
++#include <cstring>
++#include <stdexcept>
++namespace prefix_short {
++inline bool requested() {
++ const char *v=std::getenv("FPX_SHORT_PREFIX_CHECKPOINT");
++ if(!v || !*v || !std::strcmp(v,"0"))return false;
++ if(std::strcmp(v,"1"))throw std::invalid_argument("FPX_SHORT_PREFIX_CHECKPOINT must be 0 or 1");
++ return true;
++}
++inline bool boundary(int64_t tokens){return tokens==1024 || tokens==2048;}
++template<class Checkpoints> auto victim(Checkpoints & checkpoints) {
++ // Exactly two prefix snapshots are pinned; total capacity remains four.
++ // Source/token invalidation is still owned by the existing restore path.
++ for(auto it=checkpoints.begin();it!=checkpoints.end();++it)
++  if(!boundary(it->n_tokens))return it;
++ throw std::logic_error("No evictable checkpoint in four-entry short-prefix policy");
++}
++}
+--- a/tools/server/server-common.h
++++ b/tools/server/server-common.h
+@@ -202,6 +202,9 @@
+     size_t size() const { return tokens.size(); }
+ 
+     bool empty() const { return tokens.empty(); }
++
++    // has_mtmd enables media handling; this tests actual media in this prompt.
++    bool has_media_chunks() const { return !map_idx_to_media.empty(); }
+ 
+     void clear() {
+         map_idx_to_media.clear();
+--- a/tools/server/server-context.cpp
++++ b/tools/server/server-context.cpp
+@@ -1,5 +1,7 @@
+ 
+ #include "server-context.h"
++#include "prefix-short-policy.h"
++#include "server-projector.h"
+ #include "server-chat.h"
+ #include "server-common.h"
+ #include "server-http.h"
+@@ -7,6 +9,7 @@
+ #include "server-queue.h"
+ 
+ #include "build-info.h"
++#include "ggml-backend.h"
+ #include "common.h"
+ #include "diffusion.h"
+ #include "diffusion-output.h"
+@@ -682,6 +685,14 @@
+ // server_context_impl (private implementation)
+ //
+ 
++static bool projector_prompt_has_media(const json & value) {
++    if (value.is_object()) return value.contains("multimodal_data");
++    if (value.is_array()) {
++        for (const auto & part : value) if (projector_prompt_has_media(part)) return true;
++    }
++    return false;
++}
++
+ struct server_context_impl {
+     friend struct server_context;
+     friend struct server_routes;
+@@ -765,20 +776,53 @@
+     std::set<std::string> model_aliases; // additional names for the model
+     std::set<std::string> model_tags;    // informational tags
+ 
++    mutable server_projector_lifecycle projector;
++    std::atomic<int> llm_load_count{0};
++    std::atomic<bool> llm_loaded{false};
+     bool sleeping = false;
+ 
++    void load_projector() const {
++        if (params_base.mmproj.path.empty()) throw std::runtime_error("projector is not configured");
++        mtmd_context_params p = mtmd_context_params_default();
++        p.use_gpu = params_base.mmproj_use_gpu;
++        p.print_timings = false;
++        p.n_threads = params_base.cpuparams.n_threads;
++        p.flash_attn_type = params_base.flash_attn_type;
++        p.warmup = params_base.warmup;
++        p.image_min_tokens = params_base.image_min_tokens;
++        p.image_max_tokens = params_base.image_max_tokens;
++        p.media_marker = get_media_marker();
++        auto & self = const_cast<server_context_impl &>(*this);
++        self.mctx = mtmd_init_from_file(params_base.mmproj.path.c_str(), model_tgt, p);
++        if (!self.mctx) throw std::runtime_error("projector initialization failed");
++        for (auto & slot : self.slots) slot.mctx = self.mctx;
++    }
++
++    void release_projector() {
++        mtmd_free(mctx);
++        mctx = nullptr;
++        for (auto & slot : slots) slot.mctx = nullptr;
++    }
++
++    void expire_projector() {
++        projector.expire([this] {
++            if (queue_tasks.has_pending_tasks()) return false;
++            for (const auto & slot : slots) if (slot.is_processing()) return false;
++            return true;
++        }, [this] { release_projector(); });
++    }
++
+     void destroy() {
++        projector.release_all([this] { release_projector(); });
+         spec.reset();
+         ctx_dft.reset();
+         model_dft.reset();
+ 
++        llm_loaded.store(false);
+         llama_init.reset();
+ 
+         ctx_tgt = nullptr;
+         model_tgt = nullptr;
+-
+-        mtmd_free(mctx);
+-        mctx = nullptr;
+ 
+         llama_batch_free(batch);
+     }
+@@ -806,6 +850,7 @@
+ 
+         params_base = params;
+ 
++        llm_load_count.fetch_add(1);
+         llama_init = common_init_from_params(params_base);
+ 
+         model_tgt = llama_init->model();
+@@ -816,6 +861,7 @@
+             return false;
+         }
+ 
++        llm_loaded.store(true);
+         vocab = llama_model_get_vocab(model_tgt);
+ 
+         n_ctx = llama_n_ctx(ctx_tgt);
+@@ -961,20 +1007,11 @@
+ 
+         std::string & mmproj_path = params_base.mmproj.path;
+         if (!mmproj_path.empty()) {
+-            mtmd_context_params mparams = mtmd_context_params_default();
+-
+-            mparams.use_gpu          = params_base.mmproj_use_gpu;
+-            mparams.print_timings    = false;
+-            mparams.n_threads        = params_base.cpuparams.n_threads;
+-            mparams.flash_attn_type  = params_base.flash_attn_type;
+-            mparams.warmup           = params_base.warmup;
+-            mparams.image_min_tokens = params_base.image_min_tokens;
+-            mparams.image_max_tokens = params_base.image_max_tokens;
+-            mparams.media_marker     = get_media_marker();
+-
+-            mctx = mtmd_init_from_file(mmproj_path.c_str(), model_tgt, mparams);
+-            if (mctx == nullptr) {
+-                SRV_ERR("failed to load multimodal model, '%s'\n", mmproj_path.c_str());
++            try {
++                load_projector();
++                projector.adopt_loaded();
++            } catch (const std::exception & error) {
++                SRV_ERR("projector load failed: %s\n", error.what());
+                 return false;
+             }
+             SRV_INF("loaded multimodal model, '%s'\n", mmproj_path.c_str());
+@@ -1167,15 +1204,21 @@
+ 
+         // wiring up server queues
+         queue_tasks.on_new_task([this](server_task && task) {
+-            process_single_task(std::move(task));
++            auto lease = projector.acquire(false, {}, {});
++            if (lease) process_single_task(std::move(task));
+         });
+         queue_tasks.on_update_slots([this]() {
+-            update_slots();
++            {
++                auto lease = projector.acquire(false, {}, {});
++                if (lease) update_slots();
++            }
++            expire_projector();
+         });
+         queue_tasks.on_sleeping_state([this](bool sleeping) {
+             handle_sleeping_state(sleeping);
+         });
+ 
++        queue_tasks.on_idle([this] { expire_projector(); });
+         metrics.init();
+ 
+         if (params_base.cache_idle_slots) {
+@@ -2019,14 +2062,27 @@
+ 
+     // n_tokens_cur: the number of tokens added to the batch for the current slot
+     void create_checkpoint(server_slot & slot, const int64_t n_tokens_cur, llama_pos pos_min, llama_pos pos_max) {
++        const int64_t probe_save_started = ggml_time_us();
++        const bool short_prefix = prefix_short::requested() && params_base.n_parallel == 1 &&
++            params_base.n_ctx_checkpoints == 4 && params_base.checkpoint_every_nt == 2048 &&
++            !slot.prompt.tokens.has_media_chunks() && !slot.task->tokens.has_media_chunks() &&
++            !ctx_dft && slot.lora.empty();
+         while (slot.prompt.checkpoints.size() >= (size_t) params_base.n_ctx_checkpoints) {
+-            // make room for the new checkpoint, if needed
+-            const auto & cur = slot.prompt.checkpoints.front();
++            // Retain the initial aligned prefix when possible, while keeping
++            // the configured checkpoint count bounded. Later checkpoints remain
++            // oldest-first; ordinary source changes still invalidate checkpoints.
++            auto victim = slot.prompt.checkpoints.begin();
++            if (slot.prompt.checkpoints.size() > 1 &&
++                    victim->n_tokens <= params_base.checkpoint_every_nt) {
++                ++victim;
++            }
++            if (short_prefix) victim = prefix_short::victim(slot.prompt.checkpoints);
++            const auto & cur = *victim;
+ 
+             SLT_WRN(slot, "erasing old context checkpoint (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", size = %.3f MiB)\n",
+                     cur.pos_min, cur.pos_max, cur.n_tokens, (float) cur.size() / 1024 / 1024);
+ 
+-            slot.prompt.checkpoints.erase(slot.prompt.checkpoints.begin());
++            slot.prompt.checkpoints.erase(victim);
+         }
+ 
+         auto & cur = slot.prompt.checkpoints.emplace_back();
+@@ -2050,6 +2106,7 @@
+                 "created context checkpoint %d of %d (pos_min = %d, pos_max = %d, n_tokens = %" PRId64 ", size = %.3f MiB)\n",
+                 (int) slot.prompt.checkpoints.size(), params_base.n_ctx_checkpoints, cur.pos_min,
+                 cur.pos_max, cur.n_tokens, (float) cur.size() / 1024 / 1024);
++        SLT_INF(slot, "checkpoint_probe_save n_tokens=%" PRId64 " wall_ms=%.3f\n", cur.n_tokens, (ggml_time_us()-probe_save_started)/1e3);
+     }
+ 
+     void process_single_task(server_task && task) {
+@@ -2984,9 +3041,12 @@
+ 
+                                     if (!do_reset) {
+                                         // restore the context checkpoint
++                                        const int64_t restore_probe_started = ggml_time_us();
+ 
+                                         const bool restored_tgt = it->load_tgt(ctx_tgt,       slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
+                                         const bool restored_dft = it->load_dft(ctx_dft.get(), slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE);
++
++                                        SLT_INF(slot, "checkpoint_probe_restore n_tokens=%" PRId64 " wall_ms=%.3f target_ok=%d draft_ok=%d\n", it->n_tokens, (ggml_time_us()-restore_probe_started)/1e3, (int) restored_tgt, (int) restored_dft);
+ 
+                                         if (!restored_tgt || !restored_dft) {
+                                             SLT_WRN(slot,
+@@ -3146,6 +3206,11 @@
+                         has_mtmd = true;
+                     }
+ 
++                    const bool short_prefix = prefix_short::requested() && params_base.n_parallel == 1 &&
++                        params_base.n_ctx_checkpoints == 4 && params_base.checkpoint_every_nt == 2048 &&
++                        !slot.prompt.tokens.has_media_chunks() && !slot.task->tokens.has_media_chunks() &&
++            !ctx_dft && slot.lora.empty();
++
+                     // add prompt tokens for processing in the current batch
+                     while (slot.prompt.n_tokens() < slot.task->n_tokens() && batch.n_tokens < n_batch) {
+                         // get next token to process
+@@ -3172,6 +3237,9 @@
+                         slot.prompt.tokens.push_back(cur_tok);
+ 
+                         slot.n_prompt_tokens_processed++;
++
++                        // Split before decode so the next iteration can snapshot actual state at 1024/2048.
++                        if (short_prefix && do_checkpoint && prefix_short::boundary(slot.prompt.n_tokens())) break;
+ 
+                         // process the last few tokens of the prompt separately in order to allow for a checkpoint to be created.
+                         // create checkpoints that many tokens before the end of the prompt:
+@@ -3225,7 +3293,10 @@
+                                     last_checkpoint = slot.prompt.checkpoints.back().n_tokens;
+                                 }
+ 
+-                                do_checkpoint = do_checkpoint && slot.prompt.n_tokens() - batch.n_tokens - last_checkpoint >= params_base.checkpoint_every_nt;
++                                const int64_t processed = slot.prompt.n_tokens() - n_tokens_cur;
++                                do_checkpoint = do_checkpoint && (
++                                    (short_prefix && prefix_short::boundary(processed)) ||
++                                    slot.prompt.n_tokens() - batch.n_tokens - last_checkpoint >= params_base.checkpoint_every_nt);
+ 
+                                 if (do_checkpoint) {
+                                     SLT_INF(slot, "%d tokens since last checkpoint at %d, creating new checkpoint during processing at position %d\n", params_base.checkpoint_every_nt, last_checkpoint, slot.prompt.n_tokens());
+@@ -3973,12 +4044,17 @@
+     return impl->load_model(params);
+ }
+ 
++void server_context::set_projector_idle_seconds(int seconds) {
++    impl->projector.configure(seconds);
++}
++
+ void server_context::start_loop() {
+     auto & params = impl->params_base;
+     impl->queue_tasks.start_loop(params.sleep_idle_seconds * 1000);
+ }
+ 
+ void server_context::terminate() {
++    impl->projector.begin_close();
+     impl->queue_tasks.terminate();
+ }
+ 
+@@ -4037,6 +4113,7 @@
+ // generator-like API for HTTP response generation
+ // may have bypass_sleep = true if the task does not use ctx_server
+ struct server_res_generator : server_http_res {
++    std::unique_ptr<server_projector_lifecycle::lease> projector_lease;
+     server_response_reader rd;
+     server_res_generator(server_queue & queue_tasks, server_response & queue_results, int sleep_idle_seconds, bool bypass_sleep = false)
+             : rd(queue_tasks, queue_results, HTTP_POLLING_SECONDS) {
+@@ -4073,7 +4150,7 @@
+             task_response_type res_type) {
+     GGML_ASSERT(type == SERVER_TASK_TYPE_COMPLETION || type == SERVER_TASK_TYPE_INFILL);
+ 
+-    auto res = create_response();
++    auto res = create_response(false, !files.empty() || projector_prompt_has_media(data.at("prompt")), req.should_stop);
+     auto completion_id = gen_chatcmplid();
+     auto & rd = res->rd;
+ 
+@@ -4294,8 +4371,14 @@
+     return res;
+ }
+ 
+-std::unique_ptr<server_res_generator> server_routes::create_response(bool bypass_sleep) {
+-    return std::make_unique<server_res_generator>(queue_tasks, queue_results, params.sleep_idle_seconds, bypass_sleep);
++std::unique_ptr<server_res_generator> server_routes::create_response(bool bypass_sleep, bool visual,
++        const std::function<bool()> & cancelled) {
++    auto response = std::make_unique<server_res_generator>(queue_tasks, queue_results, params.sleep_idle_seconds, bypass_sleep);
++    if (!bypass_sleep && !params.mmproj.path.empty()) {
++        response->projector_lease = ctx_server.projector.acquire(visual, cancelled, [this] { ctx_server.load_projector(); });
++        if (!response->projector_lease) throw std::runtime_error("projector request cancelled or server closing");
++    }
++    return response;
+ }
+ 
+ server_routes::server_routes(const common_params & params, server_context & ctx_server)
+@@ -4307,6 +4390,24 @@
+ }
+ 
+ void server_routes::init_routes() {
++    this->get_projector_status = [this](const server_http_req &) {
++        auto res = create_response(true);
++        const auto p = ctx_server.projector.status();
++        json memory = json::array();
++        for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
++            auto dev = ggml_backend_dev_get(i);
++            if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_GPU) continue;
++            size_t available = 0, total = 0;
++            ggml_backend_dev_memory(dev, &available, &total);
++            memory.push_back({{"device", ggml_backend_dev_name(dev)}, {"free_bytes", available}, {"total_bytes", total}});
++        }
++        res->ok({{"projector_loaded", p.loaded}, {"closing", p.closing},
++                 {"request_leases", p.readers}, {"visual_leases", p.visual_readers},
++                 {"load_count", p.loads}, {"unload_count", p.unloads},
++                 {"idle_seconds", p.idle_ms / 1000}, {"llm_loaded", ctx_server.llm_loaded.load()},
++                 {"llm_load_count", ctx_server.llm_load_count.load()}, {"gpu_memory", memory}});
++        return res;
++    };
+     // IMPORTANT: all lambda functions must start with create_response()
+     // this is to ensure that the server_res_generator can handle sleeping case correctly
+ 
+@@ -4635,6 +4736,7 @@
+         );
+ 
+         std::vector<raw_buffer> files; // dummy
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_INFILL,
+@@ -4647,6 +4749,7 @@
+         auto res = create_response();
+         std::vector<raw_buffer> files; // dummy
+         const json body = json::parse(req.body);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -4659,6 +4762,7 @@
+         auto res = create_response();
+         std::vector<raw_buffer> files; // dummy
+         const json body = json::parse(req.body);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -4675,6 +4779,7 @@
+             body,
+             meta->chat_params,
+             files);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -4693,6 +4798,7 @@
+             body,
+             meta->chat_params,
+             files);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -4721,6 +4827,7 @@
+             body,
+             meta->chat_params,
+             files);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -4739,6 +4846,7 @@
+             body,
+             meta->chat_params,
+             files);
++        res.reset(); // Hand preprocessing ownership to the inference response before a possible reload.
+         return handle_completions_impl(
+             req,
+             SERVER_TASK_TYPE_COMPLETION,
+@@ -5181,6 +5289,11 @@
+         }
+     }
+ 
++    if (projector_prompt_has_media(prompt)) {
++        res->projector_lease.reset();
++        res->projector_lease = ctx_server.projector.acquire(true, req.should_stop, [this] { ctx_server.load_projector(); });
++        if (!res->projector_lease) throw std::runtime_error("projector request cancelled or server closing");
++    }
+     auto tokenized_prompts = tokenize_input_prompts(ctx_server.vocab, ctx_server.mctx, prompt, true, true);
+     for (const auto & tokens : tokenized_prompts) {
+         // this check is necessary for models that do not add BOS token to the input
+--- a/tools/server/server-context.h
++++ b/tools/server/server-context.h
+@@ -67,6 +67,9 @@
+     // terminate main loop (will unblock start_loop)
+     void terminate();
+ 
++    // Set before load_model; 0 disables independent projector expiry.
++    void set_projector_idle_seconds(int seconds);
++
+     // get the underlaying llama_context, can return nullptr if sleeping
+     // not thread-safe, should only be used from the main thread
+     llama_context * get_llama_context() const;
+@@ -122,6 +125,7 @@
+     server_http_context::handler_t post_rerank;
+     server_http_context::handler_t get_lora_adapters;
+     server_http_context::handler_t post_lora_adapters;
++    server_http_context::handler_t get_projector_status;
+ 
+     // to be used in router mode
+     json get_model_info() const;
+@@ -146,5 +150,6 @@
+ 
+     server_queue & queue_tasks;
+     server_response & queue_results;
+-    std::unique_ptr<server_res_generator> create_response(bool bypass_sleep = false);
++    std::unique_ptr<server_res_generator> create_response(bool bypass_sleep = false, bool visual = false,
++            const std::function<bool()> & cancelled = {});
+ };
+--- /dev/null
++++ b/tools/server/server-projector.h
+@@ -0,0 +1,111 @@
++#pragma once
++
++#include <atomic>
++#include <chrono>
++#include <functional>
++#include <memory>
++#include <mutex>
++#include <shared_mutex>
++#include <stdexcept>
++#include <thread>
++
++// Protects projector ownership across HTTP preprocessing, queued requests and
++// native inference. Only the main loop invokes expire(); it checks slots/queue
++// before entering this boundary. No model or KV resource is owned here.
++class server_projector_lifecycle {
++    static int64_t now_ms() {
++        return std::chrono::duration_cast<std::chrono::milliseconds>(
++            std::chrono::steady_clock::now().time_since_epoch()).count();
++    }
++    mutable std::shared_mutex mutex;
++    std::atomic<bool> closing{false};
++    std::atomic<bool> loaded{false};
++    std::atomic<int64_t> last_visual_ms{0};
++    std::atomic<int> readers{0};
++    std::atomic<int> visual_readers{0};
++    std::atomic<int> loads{0};
++    std::atomic<int> unloads{0};
++    int64_t idle_ms = 120000;
++
++public:
++    struct snapshot {
++        bool loaded;
++        bool closing;
++        int readers;
++        int visual_readers;
++        int loads;
++        int unloads;
++        int64_t idle_ms;
++    };
++    class lease {
++        server_projector_lifecycle & owner;
++        std::shared_lock<std::shared_mutex> lock;
++        bool visual;
++    public:
++        lease(server_projector_lifecycle & owner, std::shared_lock<std::shared_mutex> lock, bool visual)
++            : owner(owner), lock(std::move(lock)), visual(visual) {
++            owner.readers.fetch_add(1);
++            if (visual) owner.visual_readers.fetch_add(1);
++        }
++        ~lease() {
++            if (visual) {
++                owner.last_visual_ms.store(now_ms());
++                owner.visual_readers.fetch_sub(1);
++            }
++            owner.readers.fetch_sub(1);
++        }
++        lease(const lease &) = delete;
++        lease & operator=(const lease &) = delete;
++    };
++
++    void configure(int seconds) {
++        if (seconds < 0 || seconds > 86400) throw std::invalid_argument("projector idle seconds must be 0..86400");
++        idle_ms = int64_t(seconds) * 1000; // 0 explicitly disables expiry.
++    }
++    void adopt_loaded() {
++        loaded.store(true); loads.fetch_add(1); last_visual_ms.store(now_ms());
++    }
++    std::unique_ptr<lease> acquire(bool visual, const std::function<bool()> & cancelled,
++                                 const std::function<void()> & load) {
++        auto stopped = [&] { return closing.load() || (cancelled && cancelled()); };
++        while (!stopped()) {
++            std::shared_lock<std::shared_mutex> read(mutex, std::try_to_lock);
++            if (read.owns_lock()) {
++                if (!visual || loaded.load()) {
++                    if (stopped()) return nullptr;
++                    return std::make_unique<lease>(*this, std::move(read), visual);
++                }
++                read.unlock();
++                std::unique_lock<std::shared_mutex> write(mutex, std::try_to_lock);
++                if (write.owns_lock() && !stopped() && !loaded.load()) {
++                    load(); // Failed load remains unavailable; no fake success.
++                    loaded.store(true); loads.fetch_add(1); last_visual_ms.store(now_ms());
++                }
++            }
++            std::this_thread::sleep_for(std::chrono::milliseconds(2));
++        }
++        return nullptr;
++    }
++    bool expire(const std::function<bool()> & safe, const std::function<void()> & release,
++                int64_t time_ms = -1) {
++        if (!idle_ms || !loaded.load() || closing.load()) return false;
++        const int64_t now = time_ms < 0 ? now_ms() : time_ms;
++        if (now - last_visual_ms.load() < idle_ms) return false;
++        std::unique_lock<std::shared_mutex> write(mutex, std::try_to_lock);
++        if (!write.owns_lock() || !loaded.load() || closing.load() || !safe()) return false;
++        release();
++        loaded.store(false); unloads.fetch_add(1);
++        return true;
++    }
++    void begin_close() { closing.store(true); }
++    void release_all(const std::function<void()> & release) {
++        std::unique_lock<std::shared_mutex> write(mutex);
++        if (loaded.load()) {
++            release(); loaded.store(false); unloads.fetch_add(1);
++        }
++    }
++    snapshot status() const {
++        return {loaded.load(), closing.load(), readers.load(), visual_readers.load(),
++                loads.load(), unloads.load(), idle_ms};
++    }
++};
+--- a/tools/server/server-queue.cpp
++++ b/tools/server/server-queue.cpp
+@@ -202,6 +202,8 @@
+                 if (res) {
+                     break; // new task arrived or terminate
+                 }
++                lock.unlock();
++                if (callback_idle) callback_idle();
+                 // otherwise, loop again to check sleeping condition
+             }
+         }
+--- a/tools/server/server-queue.h
++++ b/tools/server/server-queue.h
+@@ -29,6 +29,7 @@
+     std::function<void(server_task &&)> callback_new_task;
+     std::function<void(void)>           callback_update_slots;
+     std::function<void(bool)>           callback_sleeping_state;
++    std::function<void(void)>           callback_idle;
+ 
+ public:
+     // Add a new task to the end of the queue
+@@ -75,6 +76,12 @@
+      */
+     void start_loop(int64_t idle_sleep_ms = -1);
+ 
++    bool has_pending_tasks() {
++        std::unique_lock<std::mutex> lock(mutex_tasks);
++        return !queue_tasks.empty() || !queue_tasks_deferred.empty();
++    }
++    void on_idle(std::function<void(void)> callback) { callback_idle = std::move(callback); }
++
+     // for metrics
+     size_t queue_tasks_deferred_size() {
+         std::unique_lock<std::mutex> lock(mutex_tasks);
+--- a/tools/server/server.cpp
++++ b/tools/server/server.cpp
+@@ -102,6 +102,24 @@
+ int main(int argc, char ** argv) {
+     std::setlocale(LC_NUMERIC, "C");
+ 
++    // Candidate-only option; do not change common_params/DLL ABI.
++    int projector_idle_seconds = 120;
++    std::vector<char *> filtered_args;
++    try {
++        for (int i = 0; i < argc; ++i) {
++            if (std::string(argv[i]) == "--projector-idle-seconds") {
++                if (++i >= argc) throw std::invalid_argument("projector idle seconds missing");
++                size_t end = 0;
++                projector_idle_seconds = std::stoi(argv[i], &end);
++                if (argv[i][end] || projector_idle_seconds < 0 || projector_idle_seconds > 86400)
++                    throw std::invalid_argument("projector idle seconds must be 0..86400");
++            } else filtered_args.push_back(argv[i]);
++        }
++    } catch (const std::exception & error) {
++        fprintf(stderr, "%s\n", error.what()); return 1;
++    }
++    argc = int(filtered_args.size()); filtered_args.push_back(nullptr); argv = filtered_args.data();
++
+     // own arguments required by this example
+     common_params params;
+ 
+@@ -111,6 +129,10 @@
+         return 1;
+     }
+ 
++    if (!params.mmproj.path.empty() && params.sleep_idle_seconds >= 0) {
++        fprintf(stderr, "Independent projector lifecycle requires whole-model sleep disabled.\n");
++        return 1;
++    }
+     llama_backend_init();
+     llama_numa_init(params.numa);
+ 
+@@ -151,6 +173,7 @@
+ 
+     // struct that contains llama context and inference
+     server_context ctx_server;
++    ctx_server.set_projector_idle_seconds(projector_idle_seconds);
+ 
+     server_http_context ctx_http;
+     if (!ctx_http.init(params)) {
+@@ -210,6 +233,7 @@
+     ctx_http.get ("/health",                   ex_wrapper(routes.get_health)); // public endpoint (no API key check)
+     ctx_http.get ("/v1/health",                ex_wrapper(routes.get_health)); // public endpoint (no API key check)
+     ctx_http.get ("/metrics",                  ex_wrapper(routes.get_metrics));
++    if (!params.mmproj.path.empty()) ctx_http.get("/projector/status", ex_wrapper(routes.get_projector_status));
+     ctx_http.get ("/props",                    ex_wrapper(routes.get_props));
+     ctx_http.post("/props",                    ex_wrapper(routes.post_props));
+     ctx_http.get ("/models",                   ex_wrapper(routes.get_models)); // public endpoint (no API key check)

apply-fpx-fixes-20261021-01.py

@@ -0,0 +1,67 @@
+"""Check or apply the version-pinned AiTable ROCmFPX native patch. No build/start."""
+import argparse
+import hashlib
+import json
+from pathlib import Path, PurePosixPath
+import shutil
+import subprocess
+import uuid
+
+
+def digest(path):
+    return hashlib.sha256(path.read_bytes()).hexdigest() if path.is_file() else None
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument('--source', required=True, type=Path, help='Separate upstream source checkout; do not target installed runtimes')
+    parser.add_argument('--apply', action='store_true', help='Apply after checks; default checks only')
+    args = parser.parse_args()
+    bundle = Path(__file__).resolve().parent
+    manifest = json.loads((bundle / 'fpx-fixes-manifest-20261021-01.json').read_text(encoding='utf-8'))
+    patch = bundle / 'aitable-fpx-native-fixes-20261021-01.patch'
+    if digest(patch) != manifest['patch_sha256']:
+        raise RuntimeError('Patch checksum mismatch')
+    source = args.source.resolve(strict=True)
+    if not source.is_dir():
+        raise RuntimeError('Source must be a directory')
+    targets = {}
+    for name, hashes in manifest['files'].items():
+        relative = PurePosixPath(name)
+        if relative.is_absolute() or '..' in relative.parts or ':' in name or '\\' in name:
+            raise RuntimeError('Unsafe manifest path: ' + name)
+        target = source.joinpath(*relative.parts)
+        if not target.resolve().is_relative_to(source) or target.is_symlink():
+            raise RuntimeError('Target escapes the source tree: ' + name)
+        if target.exists() and not target.is_file():
+            raise RuntimeError('Target is not a regular file: ' + name)
+        if digest(target) != hashes['before']:
+            raise RuntimeError('Source differs from pinned upstream (or already patched): ' + name)
+        targets[name] = target
+    git_apply = ['git', '-c', 'core.autocrlf=false', '-c', 'apply.whitespace=nowarn', 'apply']
+    subprocess.run([*git_apply, '--check', str(patch)], cwd=source, check=True, timeout=30)
+    if not args.apply:
+        print('CHECK PASSED: upstream ' + manifest['upstream_commit'] + '; no files modified')
+        return
+    backup = source.parent / (source.name + '-aitable-backup-' + uuid.uuid4().hex)
+    backup.mkdir()
+    for name, target in targets.items():
+        if target.exists():
+            saved = backup / name
+            saved.parent.mkdir(parents=True, exist_ok=True)
+            shutil.copy2(target, saved)
+    (backup / 'restore-manifest.json').write_text(json.dumps({
+        'source': str(source), 'files': manifest['files'],
+        'restore': 'Stop builds first. Restore saved files; remove only files whose before hash is null. Refuse to overwrite later changes unless verified.'
+    }, indent=2), encoding='utf-8')
+    print('BACKUP: ' + str(backup), flush=True)
+    # Git apply is atomic by default; never pass --reject or allow partial application.
+    subprocess.run([*git_apply, str(patch)], cwd=source, check=True, timeout=30)
+    for name, target in targets.items():
+        if digest(target) != manifest['files'][name]['after']:
+            raise RuntimeError('Post-apply checksum mismatch; inspect retained backup: ' + name)
+    print('APPLIED AND VERIFIED. No compilation, model download or service start performed.')
+
+
+if __name__ == '__main__':
+    main()

configure-fpx-env-20261021-01.cmd

@@ -0,0 +1,23 @@
+@echo off
+if "%~1"=="" (
+ echo Usage: call configure-fpx-env-20261021-01.cmd "extracted-runtime-package-root"
+ exit /b 2
+)
+if not exist "%~1\runtime\llama-server.exe" (
+ echo Runtime package root is invalid.
+ exit /b 2
+)
+set "PATH=%~1\runtime;%PATH%"
+set "ROCBLAS_TENSILE_LIBPATH=%~1\runtime\rocblas\library"
+set "ROCBLAS_USE_HIPBLASLT=0"
+set "FPX_SHORT_PREFIX_CHECKPOINT=1"
+set "GGML_FPX_NORM_SILU_FUSION=0"
+set "GGML_CUDA_GDN_CHUNKED=0"
+set "GGML_CUDA_GDN_CHUNKED_BF16=0"
+set "GGML_FPX_Q8_ATTENTION=0"
+set "GGML_FA_RDNA4_RESEARCH=0"
+set "GGML_CUDA_DISABLE_ROPE_FWHT_SET_ROWS=1"
+if not defined GGML_HIP_CK_ATTENTION set "GGML_HIP_CK_ATTENTION=0"
+rem RAM checkpoint cache is a separate, bounded host-memory setting in MiB.
+set "LLAMA_ARG_CACHE_RAM=1024"
+exit /b 0

fpx-fixes-manifest-20261021-01.json

@@ -0,0 +1,138 @@
+{
+  "upstream_commit": "fb08d7cdb67082ddda15a25e64dc7ec413afc41c",
+  "paired_runtime_release": "20261021-01",
+  "patch_sha256": "dcdd6a8ebca68467f3faa9fd07e2f69f11e25446f100b4c214cc2619fc681f6b",
+  "files": {
+    "common/chat-auto-parser-generator.cpp": {
+      "before": "fbb122eb03840a8b2cf1b55acf1d5a0ee9c2f51de60a50a387ae5c55e5f29a18",
+      "after": "7448939113929b3e1b904b87fbb9fbc076fce84fbc486f3fd1cb0e25482bd7d4",
+      "frozen_native": "7448939113929b3e1b904b87fbb9fbc076fce84fbc486f3fd1cb0e25482bd7d4"
+    },
+    "common/chat.cpp": {
+      "before": "7569be8fe13a892ec5c629f0198109ac5efcb9f01276ad718c8ab614b707850b",
+      "after": "1a296ae1b0af51b1d847639cd1c99cd2700863d42e9286c62f2effbbbff15b81",
+      "frozen_native": "1a296ae1b0af51b1d847639cd1c99cd2700863d42e9286c62f2effbbbff15b81"
+    },
+    "common/chat.h": {
+      "before": "0eda44d1f9c2f0e456aaa11730d3119f99258ded03ca76d1473769f864610782",
+      "after": "0625c27306172d7cddadc227780a4e3a11b1f8e39872e9f6421097b40006681c",
+      "frozen_native": "0625c27306172d7cddadc227780a4e3a11b1f8e39872e9f6421097b40006681c"
+    },
+    "common/common.h": {
+      "before": "91e3190fd56820bad0bdff38b96c5bf5b3ef5d2e616a87a86596880f9135308e",
+      "after": "96c417eafc179c159f7fcef5b28e69f62396d3c15bd953584199b5299e3a0de3",
+      "frozen_native": "96c417eafc179c159f7fcef5b28e69f62396d3c15bd953584199b5299e3a0de3"
+    },
+    "common/json-schema-to-grammar.cpp": {
+      "before": "589c76fbc72159a28a9b2fd6a46248ce570960daeb688a92732315633069703d",
+      "after": "b3b4f0139cf41cdfa695b8d6c9c78c8a08f8e6f5fc7bfe0f85684b52b549f67b",
+      "frozen_native": "b3b4f0139cf41cdfa695b8d6c9c78c8a08f8e6f5fc7bfe0f85684b52b549f67b"
+    },
+    "common/peg-parser.cpp": {
+      "before": "71c7e355f89d99ff4c00b4142932b1a8033a07f2aa9db7b156af5ed479c4a4eb",
+      "after": "c3d0597d8bc219cde799ac5ae9ab52fb11414da2cec58bc4d613c2d4a1c17b2f",
+      "frozen_native": "c3d0597d8bc219cde799ac5ae9ab52fb11414da2cec58bc4d613c2d4a1c17b2f"
+    },
+    "common/sampling.cpp": {
+      "before": "9682ef2df8ae8ffb43d5dc561ed326a2e9e4b2e1a46388b59b9aee739ce01333",
+      "after": "bafabee9e05e436d960bd00d5fd24ea8d013899387fa616d2a3f5f6ceb412913",
+      "frozen_native": "bafabee9e05e436d960bd00d5fd24ea8d013899387fa616d2a3f5f6ceb412913"
+    },
+    "common/sampling.h": {
+      "before": "477c5336f273c0f584c71391ef82f459f8ad4f505d2343b68de87765c29028e0",
+      "after": "35b21345e60fccaee2b45159352a1ca20bd5ab57aee534edf62eea9359abe7bc",
+      "frozen_native": "35b21345e60fccaee2b45159352a1ca20bd5ab57aee534edf62eea9359abe7bc"
+    },
+    "ggml/src/ggml-cuda/fattn-common.cuh": {
+      "before": "8345bbd1f857d1b496e9375c3dd01d6c503e01dace76f2759ebb6fb8e25584d0",
+      "after": "74cfd9e838da8ddf71537c1e8e98bb11c2dd8625c7bd2242546ec7934db29ed7",
+      "frozen_native": "74cfd9e838da8ddf71537c1e8e98bb11c2dd8625c7bd2242546ec7934db29ed7"
+    },
+    "ggml/src/ggml-cuda/fattn-mma-f16.cuh": {
+      "before": "2523ee0bc7b00c62d3966c642b9a8804803183090161d0ffe92cbaedfd89983e",
+      "after": "c21bdec94bd3550ac69b730b621462a2604bc381fb227ef5aa2ec568e2855c8b",
+      "frozen_native": "c21bdec94bd3550ac69b730b621462a2604bc381fb227ef5aa2ec568e2855c8b"
+    },
+    "ggml/src/ggml-cuda/fattn.cu": {
+      "before": "16c3a9135b22f63e26653f5c639826b16c6884c64240d829d5d78641d1a15cfc",
+      "after": "60340eaee4b21f45fd02802c67db79f10842737cdf92bf289f5c54c4cddca4d3",
+      "frozen_native": "60340eaee4b21f45fd02802c67db79f10842737cdf92bf289f5c54c4cddca4d3"
+    },
+    "ggml/src/ggml-hip/CMakeLists.txt": {
+      "before": "dbee00eff8426df3d3f5f163636c48ab63ded77e8c868201cae441c0d27a9e02",
+      "after": "3470d7d8acf3fa09beb84a1e5b049afc840e1cb1483a9e881f1f75deea60fa0f",
+      "frozen_native": "9a9c2ee0c7844e44a47e4cd0b0bf82b53da5a195a83aa5dd744b62a92b143791"
+    },
+    "src/llama-context.cpp": {
+      "before": "49c33c06947c32390ebe81911ebfff25bf3b4803ca976ea825d2fcaa63696619",
+      "after": "2a1cc998d7051109a57d20b0795c617ab4ba3ac3f4fb0820907a33449bbb885c",
+      "frozen_native": "2a1cc998d7051109a57d20b0795c617ab4ba3ac3f4fb0820907a33449bbb885c"
+    },
+    "tools/cli/CMakeLists.txt": {
+      "before": "0af665c052ca6b89c3f6600b85ebee9183b928821f6ebec37dc2495cacd94917",
+      "after": "41f52f0f4f526fb835c65f2c2d1064daec059e735c2e164a53c91c2a1380d827",
+      "frozen_native": "41f52f0f4f526fb835c65f2c2d1064daec059e735c2e164a53c91c2a1380d827"
+    },
+    "tools/cli/prefix-offline-probe.cpp": {
+      "before": null,
+      "after": "00bb6eee48673acf0e8855f5368d096a3dd4ef707d4fb0e996d5ebc4d9437e9f",
+      "frozen_native": "00bb6eee48673acf0e8855f5368d096a3dd4ef707d4fb0e996d5ebc4d9437e9f"
+    },
+    "tools/server/prefix-short-policy.h": {
+      "before": null,
+      "after": "d7af7a6e5af8fa37e4ccddfba14068867f18e6c22194c9de93205cb3d95d32f6",
+      "frozen_native": "d7af7a6e5af8fa37e4ccddfba14068867f18e6c22194c9de93205cb3d95d32f6"
+    },
+    "tools/server/server-common.h": {
+      "before": "ba92ff90b21488b002d6d8d92a1d94abc5c39bc388d738939b3e4b32aabcb120",
+      "after": "4f99ce2a792d289f73caf6058e43285b950d994ae2151bf251bdef252ae2402a",
+      "frozen_native": "4f99ce2a792d289f73caf6058e43285b950d994ae2151bf251bdef252ae2402a"
+    },
+    "tools/server/server-context.cpp": {
+      "before": "154a0ee8158084109a12bdaaff8521ef1ccc829d48172beb7b816fd7011aafd6",
+      "after": "0cd397b0a02b784bf7ac382ede9ed9dd95ef18b3476250ee4ec1aa391ee33e0b",
+      "frozen_native": "0cd397b0a02b784bf7ac382ede9ed9dd95ef18b3476250ee4ec1aa391ee33e0b"
+    },
+    "tools/server/server-context.h": {
+      "before": "8f649d8246b40dd512e2cad39d8dc1ba59f73e23415a4a4c7fc190314fce211b",
+      "after": "10ddb15fbf1e72ce33baffece689c3b1f6c3fc6baf257efbf7ff1f5d22ee0415",
+      "frozen_native": "10ddb15fbf1e72ce33baffece689c3b1f6c3fc6baf257efbf7ff1f5d22ee0415"
+    },
+    "tools/server/server-projector.h": {
+      "before": null,
+      "after": "43293cbafc896484aeb3353cfc8afc67746f0699822678c13decd17c8ad7397d",
+      "frozen_native": "43293cbafc896484aeb3353cfc8afc67746f0699822678c13decd17c8ad7397d"
+    },
+    "tools/server/server-queue.cpp": {
+      "before": "2a4bff8757dd0dcb31c3633843271b0a26aaa8b36e07e809639ce19491f0fc64",
+      "after": "4e3269ae54b4a8fab30f86d4eaa74fb0930e131d49f1de94d0fad46d34b7e30d",
+      "frozen_native": "4e3269ae54b4a8fab30f86d4eaa74fb0930e131d49f1de94d0fad46d34b7e30d"
+    },
+    "tools/server/server-queue.h": {
+      "before": "da8c527f9d37d65f3457fe57c6042813b3581bbd36fcf82ab9dacd3f5fc74813",
+      "after": "6fbfcf333c8468ba6cfc67b21fb47d73427dec5d31dd3bad1f670b8d63b10e2c",
+      "frozen_native": "6fbfcf333c8468ba6cfc67b21fb47d73427dec5d31dd3bad1f670b8d63b10e2c"
+    },
+    "tools/server/server.cpp": {
+      "before": "67e8b3c9da2a5fd8443dfb1dd6e82b3b77a3581f632d119f194bbb103fc23364",
+      "after": "be77401cf215275b1216d005ebf0c0efe3c0441d18250be4276d9a176f9963df",
+      "frozen_native": "be77401cf215275b1216d005ebf0c0efe3c0441d18250be4276d9a176f9963df"
+    }
+  },
+  "publication_adjustment": "Only ggml-hip/CMakeLists.txt differs from frozen native code: local sibling CK import path is replaced by explicit FPX_CK_IMPORT_LIBRARY CMake parameter.",
+  "build_dependencies": {
+    "dependencies/CK/ck-attention.def": {
+      "bytes": 68,
+      "sha256": "ac30a0da25d6a9a9974ae8adf18382b6a13cca67821dacab5392fe9a983ed867"
+    },
+    "dependencies/CK/ck-attention.lib": {
+      "bytes": 1692,
+      "sha256": "3ae9dddd19e9b212d2ce401d47801e5f83745032f3c12f7b47bc6ac7eeecca17"
+    }
+  },
+  "ck_runtime_dll_sha256": "96ecc9b94fc7ecc4f2e2dcccc3230b8b5bb950c5774ea7a166665f07a4376ea2",
+  "no_models": true,
+  "full_source_archive": false,
+  "full_rebuild_verified": false,
+  "gpu_test_this_publication": false
+}