Release Ciru v4 (4.0.0): integrated runtime and patch notes

Publish the qualified combined runtime, matched source provenance, evaluation summaries and additive model-card updates. Existing target/draft/vision tensors remain unchanged.

INSTALL.md

@@ -1,14 +1,33 @@
 # Install and run Ornith1.5 Ciru Halo Agent
 
-## Updating to Ciru runtime 1.0.2
+## Updating to Ciru v4 (4.0.0)
 
-Download the current `bundle/plugin-site/`, `bundle/serve.sh` and
-`bundle/serve-vision.sh` contents and restart the model process. Remove older `ciru_ornith_g256-*.dist-info/` directories after the
-new `ciru_ornith_g256-1.0.2.dist-info/` directory is present, leaving only the
-current metadata. The pinned runtime wheels, native libraries and IU4 weights
-do not need reinstalling. Native schemas apply to automatic tool calls by
-default; an explicit `strict: false` opts a function out of schema enforcement.
+Stop the model process before updating. Download the complete plugin, native
+libraries and launchers together; a plugin-only update is insufficient. Existing
+target, DFlash2, vision and source-prefill weights and pinned engine wheels are
+unchanged. Keep a copy of the previous bundle for rollback.
 
+```bash
+uvx --from huggingface_hub hf download \
+  jcbtc/Ornith1.5-Ciru-Halo-Agent-vllm-strix-halo \
+  --include 'bundle/plugin-site/*' 'bundle/native/*' 'bundle/serve*.sh' \
+    'bundle/packaging/*' README.md INSTALL.md RUNTIME-FIXES.md RELEASE.json \
+    'evaluation/*' 'runtime/ciru-v4-ornith-source.tar.gz' \
+  --local-dir .
+```
+
+Once `bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/` is present, remove
+older `ciru_ornith_g256-*.dist-info/` directories so only the 4.0.0 metadata
+remains, then restart the usual launcher. Also update any separately copied
+plugin/native directories. Existing `bundle/paths.env` overrides remain valid;
+a fresh download defaults to `installed-runtime` and `bundle/models/`.
+
+Strict automatic tool schemas, the original sampler, 262,144-token context,
+eight active requests, 44 GiB shared cache, native vision and adaptive DFlash2
+policy are preserved. Explicit function `strict: false` retains its opt-out.
+For thinking/agent work, omit a short completion cap and leave room in the
+shared prompt/output context. Rollback requires restoring the previous complete
+plugin/native/launcher set and restarting, not mixing release files.
 
 This release includes the target model, trained DFlash2 drafter, native kernels, custom vLLM plugin, and exact vLLM/AITER runtime wheels and source archives. **Use this runtime; stock `pip install vllm` does not provide the custom quantization or serving path.**
 
@@ -74,12 +93,12 @@ curl http://127.0.0.1:8000/v1/chat/completions \
   -d '{"model":"ciru-halo-agent","messages":[{"role":"user","content":"Write a Python function that merges overlapping intervals."}],"temperature":0.6,"top_p":0.95,"max_tokens":4096,"chat_template_kwargs":{"enable_thinking":false}}'
 ```
 
-The default profile supplies **262,144 tokens of per-request context capacity**, **eight active sequences**, **44 GiB shared KV/state pool**, prefix caching and adaptive speculation. Input and output share the context window; your agent client must reserve output space and compact history before filling it. Eight independent, fully populated 256K histories are not promised. The default server profile is text-only. For optional image input, run `bash bundle/serve-vision.sh --host 127.0.0.1 --port 8000` after installation. The matching native BF16 vision encoder and projector are already included in `bundle/models/target/protected-00.safetensors`; no separate GGUF mmproj is required. This enables one image per request at a 1,048,576-pixel budget with up to eight active requests. See [the model card](README.md#optional-vision--image-input) and [vision validation limits](VISION.md).
+The default profile supplies **262,144 tokens of per-request context capacity**, **eight active sequences**, **44 GiB shared KV/state pool**, prefix caching and adaptive speculation. Input and output share the context window; your agent client must reserve output space and compact history before filling it. Eight independent, fully populated 256K histories are not promised. The default server profile enables images. To explicitly select the image launcher, run `bash bundle/serve-vision.sh --host 127.0.0.1 --port 8000` after installation. The matching native BF16 vision encoder and projector are already included in `bundle/models/target/protected-00.safetensors`; no separate GGUF mmproj is required. This enables one image per request at a 1,048,576-pixel budget with up to eight active requests. See [the model card](README.md#vision--image-input) and [vision validation limits](VISION.md).
 
 First startup compiles/loads GPU kernels and creates caches. Wait for `/health` before sending work. Keep `bundle/cache` writable. Change runtime/model locations with `ORNITH_RUNTIME_ROOT`, `ORNITH_MODEL`, and `ORNITH_DRAFT`. Ordinary users do not need to change quantization or draft-policy settings.
 
 ## Build from source
 
-The [Ciru source repository](https://github.com/ciru-ai/ornith-ciru-halo-agent) contains the model plugin, all eight native kernel sources, and the corresponding build script. See its `BUILD.md` for the native rebuild command. Runtime source archives are provided in this Hugging Face repository under `runtime/`; they include the matching vLLM/AITER source, licenses and release overlay notes. The binary installation above is the tested way to assemble the pinned engine; native source rebuilding is separate from retraining or requantizing the model.
+Use the [Ciru v4 source snapshot](runtime/ciru-v4-ornith-source.tar.gz) for the exact combined plugin, all eight native kernel sources and the build script. The [source repository](https://github.com/ciru-ai/ornith-ciru-halo-agent) retains the development history. See its `BUILD.md` for the native rebuild command. Runtime source archives are provided in this Hugging Face repository under `runtime/`; they include the matching vLLM/AITER source, licenses and release overlay notes. The binary installation above is the tested way to assemble the pinned engine; native source rebuilding is separate from retraining or requantizing the model.
 
 Pinned runtime: vLLM `0.1.0rc2.dev9+g9255fd9fb9.rocm100` (base `9255fd9fb9fedf4b29d574a8d8bb21d93892cc98` plus supplied cache overlay), AITER `0.1.0rc1`, Python 3.14.3, PyTorch `2.13.0+rocm10.0.0`, ROCm SDK 10.0.0 and Transformers 5.16.1. Preserve included third-party licenses when redistributing.

README.md

@@ -23,10 +23,11 @@ tags:
   - image-text-to-text
 ---
 
-**Ciru runtime 1.0.2 — September 15, 2026.** Automatic tool calls now use
-native schema constraints to prevent missing required arguments and malformed
-tool output. IU4 weights and inference kernels are unchanged.
-Image input is enabled using the upstream BF16 vision weights. See [patch notes](RUNTIME-FIXES.md).
+**Ciru v4 · package 4.0.0 — September 17, 2026.** This continues the Ciru v3
+release line with the combined prefill, attention, draft-path, prefix-cache
+and native tool-runtime improvements. Target, drafter and vision weights
+are unchanged. Update both native libraries and the complete plugin before
+restarting. See [patch notes](RUNTIME-FIXES.md) and [update instructions](INSTALL.md).
 
 # Ornith1.5 Ciru Halo Agent (vllm strix halo)
 
@@ -36,7 +37,62 @@ Image input is enabled using the upstream BF16 vision weights. See [patch notes]
 
 The design starts with the hardware: packed four-bit weights, four-bit activation paths, Strix Halo four-bit matrix instructions, specialized kernels, and adaptive DFlash2 speculative decoding. Prefix caching and a shared memory pool let agents return to long working histories.
 
-**Measured on AMD Ryzen AI Max+ 395 / Radeon 8060S (gfx1151):**
+## Ciru v4 results
+
+**Ciru v4 accelerates long-context ingestion and returning-agent workflows.**
+The new measurements below cover the complete combined runtime. Existing
+benchmark sections retain their **Ciru v3** version labels and comparison models.
+
+| Workload | Previous published Ciru | Ciru v4 | Change |
+| --- | ---: | ---: | ---: |
+| Cold prefill, 65,536 input tokens | 1,286.66 tok/s | 1,624.19 tok/s | +26.23% |
+| Cold prefill, 253,952 input tokens | 667.61 tok/s | 870.11 tok/s | +30.33% |
+| Cached 63K history, C1 mean decode | 123.34 tok/s | 147.69 tok/s | +19.74% |
+| Cached near-256K history, C1 mean decode | 94.00 tok/s | 118.81 tok/s | +26.40% |
+| Cached near-256K history, C8 mean decode | 11.10 tok/s | 19.86 tok/s | +78.93% |
+| Short coding batch, C8 aggregate | 294.89 tok/s | 304.80 tok/s | +3.36% |
+
+These final measurements used Ciru; the accepted published controls used Sozo
+and their original runtime. They compare complete products, including those
+host/runtime differences. They do not isolate a kernel's causal contribution.
+The 97 measured requests ended naturally without API errors; **90/90 coding
+outputs passed both native base and extended checks**. Cold-prefill requests
+produced a natural 17-token acknowledgement, not sustained reasoning output.
+
+**Tradeoffs:** short C2/C4 aggregate throughput fell 6.22%/8.84%; C4 mean decode
+fell 2.98%. Cached near-256K C8 decode rose 78.93%, but mean first-token latency
+increased from 18.03 s to 25.29 s. Its complete batch still fell from 53.64 s
+to 49.20 s. See the [complete final speed comparison](evaluation/CIRU-V4-PERFORMANCE.md).
+
+A separate same-host C1 HE0–9 test of the combined numerical parent measured
+decode gains of 19.81% and 16.15% in two comparisons, pooled +17.97%. Both were
+positive, but their spread exceeded the preset uncertainty rule. HF passed
+38/40 and final 36/40; all failures were the known HE1 case also seen at BF16.
+This is a familiar speed/health panel, not a full HumanEval score.
+
+### BF16 fidelity
+
+Short-prefix probability arrays were byte-identical to the previous HF model
+across 3,072 positions. On the reused 32K and 64K panels, each with 1,024 scored
+positions, the changes were mixed:
+
+| Metric | 32K: HF → final | 64K: HF → final |
+| --- | ---: | ---: |
+| Panel perplexity, lower is better | 7.96626 → 8.00191 | 5.12747 → 5.09702 |
+| KL distance from BF16, lower is closer | 0.222280 → 0.212698 | 0.116232 → 0.119387 |
+| BF16 top-token agreement | 80.9570% → 80.7617% | 88.2813% → 88.0859% |
+| BF16 top token in model top five | 97.3633% → 97.8516% | 98.9258% → 99.0234% |
+| Top-five overlap | 74.6289% → 74.9609% | 81.5625% → 81.6602% |
+
+At 32K, 38 top tokens were recovered and 40 lost; at 64K, 28 were recovered
+and 30 lost. One diagnosed 32K confident error was a real top-token recovery;
+another only crossed the probability threshold. These panels were reused during
+diagnosis and are not fresh held-out quality evidence. The measurements belong
+to the combined numerical parent, whose code and native libraries were preserved
+when the final R03/R04 tool changes were added. See [32K fidelity](evaluation/CIRU-V4-FIDELITY32.md),
+[64K fidelity](evaluation/CIRU-V4-FIDELITY64.md) and [composition proof](evaluation/CIRU-V4-COMPOSITION.json).
+
+**Historical Ciru v3 measurements on AMD Ryzen AI Max+ 395 / Radeon 8060S (gfx1151):**
 
 - **178 tok/s single-request decode** on the ten-question coding speed screen, with **166 ms mean time to first token**.
 - **295 tok/s aggregate at eight concurrent requests**, completing the ten-question batch in **5.52 seconds**.
@@ -46,11 +102,11 @@ The design starts with the hardware: packed four-bit weights, four-bit activatio
 
 These figures describe specific workloads, rather than a universal generation rate. The benchmark tables below include quality results and the workloads where other builds are faster.
 
-**[Full benchmarks, runner builds, serving settings, and limitations](https://llm.ciru.ai/research/ornith-strix/)** · **[Source and build instructions](https://github.com/ciru-ai/ornith-ciru-halo-agent)** · **[Credits](CREDITS.md)**
+**[Full benchmarks, runner builds, serving settings, and limitations](https://llm.ciru.ai/research/ornith-strix/)** · **[Ciru v4 source snapshot](runtime/ciru-v4-ornith-source.tar.gz)** · **[Source repository](https://github.com/ciru-ai/ornith-ciru-halo-agent)** · **[Credits](CREDITS.md)**
 
-## September 13 runtime fixes
+## Earlier September 13 runtime fixes
 
-**Existing installations should update the serving plugin and restart.** This release fixes a reproduced graph-replay cache-corruption crash and rejects malformed or ambiguous tool calls with explicit API errors. The weights, kernels and adaptive DFlash2 settings are unchanged. [Update instructions, root cause and validation scope](RUNTIME-FIXES.md).
+**Historical plugin-only update; Ciru v4 requires the broader update in [INSTALL.md](INSTALL.md).** This release fixes a reproduced graph-replay cache-corruption crash and rejects malformed or ambiguous tool calls with explicit API errors. The weights, kernels and adaptive DFlash2 settings are unchanged. [Update instructions, root cause and validation scope](RUNTIME-FIXES.md).
 
 ## Running Ciru Halo Agent
 
@@ -71,7 +127,7 @@ Read **[Installation and build instructions](INSTALL.md)** first for Linux prere
 
 ### Vision / image input
 
-Image input is enabled by default in runtime 1.0.2. The explicit image launcher is also available:
+Image input remains enabled by default in Ciru v4, as introduced in runtime 1.0.2. The explicit image launcher is also available:
 
 ```bash
 bash bundle/serve-vision.sh --host 127.0.0.1 --port 8000
@@ -143,7 +199,7 @@ The startup JSON reports the effective `ORNITH_C1_POLICY` (`auto`, `k0`, `k7`, o
 
 The pool is shared. **Eight active requests does not mean eight unrelated, fully populated 256K histories fit at once.** Output must fit in the remaining request context; there is no separate promise of 256K output after a 256K input. Preserve the supplied chat template and use the client’s intended thinking setting.
 
-## Coding speed and concurrency
+## Historical v3 coding speed and concurrency
 
 HumanEval 0–9, thinking off, greedy sampling, natural end-of-sequence, cold prompt salts. Each row runs all ten questions. These are **speed and health checks**, not evidence of general coding quality.
 
@@ -160,7 +216,7 @@ HumanEval 0–9, thinking off, greedy sampling, natural end-of-sequence, cold pr
 
 For comparison, the fresh **Q4_K_XL + DFlash2 (7-token)** C1 run at 256K capacity completed the same ten-question batch in **17.32 s**, at **131.65 tok/s mean request decode** and **0.351 s mean first-token latency**, with 10/10 health. Ciru Halo Agent completed that batch **1.63× faster**. The measured ROCmFP4 C1 MTP4 mode took **19.56 s**, at **113.31 tok/s**. Complete concurrency sweeps, alternative modes, and their health outcomes are on the research page.
 
-## Fast prompt ingestion
+## Historical v3 prompt ingestion
 
 Cold input, without a prefix-cache hit. These are backend prefill-counter rates; each request then generated a natural 17-token acknowledgement. The short acknowledgement is not a sustained decode or reasoning benchmark.
 
@@ -175,7 +231,7 @@ Cold input, without a prefix-cache hit. These are backend prefill-counter rates;
 
 The complete near-256K cold request took **380.83 s**, versus **1,032.06 s for Q4_K_XL** and **1,007.04 s for ROCmFP4**. ROCmFP4’s lower-depth rows come from its retained author-profile sweep; the near-256K point is the fresh cold capture. The research page identifies each source run.
 
-## Returning agents and large histories
+## Historical v3 returning agents and large histories
 
 All rows below use **all ten coding questions**, populated histories, and confirmed prefix-cache reuse. Timing includes queueing, suffix prefill, and generation; preparing the shared history is separate.
 
@@ -190,7 +246,7 @@ All rows below use **all ten coding questions**, populated histories, and confir
 
 The tested Q4_K_XL and ROCmFP4 configurations did not share the warmed prefix successfully across eight request slots. Their **cached C8 comparison remains unavailable**; this does not establish that the runners cannot support it.
 
-## Quality and actual agent work
+## Historical v3 quality and actual agent work
 
 The release was evaluated separately from the short coding speed screen. Full EvalScope runs use the native **541 IFEval, 1,319 GSM8K, and 164 HumanEval** tasks. Tool and Hermes scores use their own native graders.
 
@@ -207,7 +263,7 @@ Full EvalScope used C8, thinking off, greedy sampling, and a 32,768-token respon
 
 The difficult subset and tool suite are single blocks, not repeated confidence estimates. BF16 agreement uses seven held-out documents with eight adjacent prefix positions each: these 56 correlated anchors measure short-prefix fidelity, not task accuracy. The original BF16 source was fixed; the comparison quants’ exact source-weight ancestry is not established.
 
-### Hermes: three complete passes per model and concurrency
+### Historical v3 Hermes: three complete passes per model and concurrency
 
 Twenty native agent scenarios per pass, three seeds, 256K capacity, up to 64 turns, and output bounded by remaining context. Scores below are native scores, not percentages of tasks passed.
 
@@ -219,7 +275,7 @@ Twenty native agent scenarios per pass, three seeds, 256K capacity, up to 64 tur
 
 Ciru’s C8 mean workflow time was **1.49× faster than ROCmFP4**, while Q4 was slightly faster than Ciru in this suite. One interrupted Q4 transport run was retained separately and replaced once; it is excluded from the three complete passes. The research page includes every pass, score range, token timing, and failure count.
 
-## Where the speed advantage changes
+## Where the historical v3 speed advantage changes
 
 Ciru Halo Agent’s strongest results are speculative coding, prompt ingestion, and shared-history workloads. **It is not the fastest model on every task.** Four ordinary prose tasks with automatic speculation measured **49.90–63.27 tok/s**, compared with approximately **59.5–59.8 tok/s for Q4** and **73.9–74.1 tok/s for ROCmFP4** in the retained no-speculation prose captures. Ciru’s target-only diagnostic measured approximately **56–58 tok/s**. The diagnostic skips draft work; its allocations still include the drafter.
 
@@ -231,6 +287,7 @@ Ciru Halo Agent’s strongest results are speculative coding, prompt ingestion,
 
 - **[Ornith Team](https://huggingface.co/ornith-ai/Ornith-1.5-35B-A3B)** created the base Ornith model. Ciru’s work is its hardware-specific quantization, kernels, serving integration, and evaluation; we do not claim authorship of the base model’s training.
 - **[jzinno](https://huggingface.co/jzinno/Ornith-1.5-35B-A3B-DFlash2)** trained the Ornith DFlash2 drafter, initialized from **[z-lab’s Qwen3.5 DFlash](https://huggingface.co/z-lab/Qwen3.5-35B-A3B-DFlash)**. Credit also goes to the **[DFlash](https://arxiv.org/abs/2602.06036)** and **[DFlash 2](https://inco.ai/blog/dflash2/)** authors and the NVIDIA Nemotron dataset contributors whose work supports that drafter.
+- **[XGrammar](https://github.com/mlc-ai/xgrammar)** supplies the native grammar implementation; its license and pinned Ciru patch are included in the v4 source snapshot.
 - **[Qwen](https://github.com/QwenLM)** supplied the underlying Qwen model architecture used by this Ornith checkpoint.
 - **[vLLM](https://github.com/vllm-project/vllm)**, **[AMD ROCm](https://github.com/ROCm)**, **[AITER](https://github.com/ROCm/aiter)**, **[Composable Kernel](https://github.com/ROCm/composable_kernel)**, **[PyTorch](https://github.com/pytorch/pytorch)**, and **[Triton](https://github.com/triton-lang/triton)** provide the runtime and compiler foundations.
 - Thanks to **[peculiar-ragdoll](https://huggingface.co/peculiar-ragdoll/Unsloth-Ornith-1.5-35B-A3B)**, **[Daniel Han Chen](https://github.com/danielhanchen/llama.cpp)**, **[julianmb](https://huggingface.co/julianmb/Ornith-1.5-35B-A3B-ROCmFP4-GGUF)**, **[HaloFPX](https://github.com/julianmb/halofpx)**, and the llama.cpp/ROCmFP4 community for the comparison builds. Their artifacts are baselines, not components of Ciru’s target weights. The Q4 quant is a community Unsloth-style release, not an official Unsloth upload.

RELEASE.json

@@ -1,299 +1,141 @@
 {
-  "model": "Ornith1.5 Ciru Halo Agent",
-  "source_repository": "https://github.com/ciru-ai/ornith-ciru-halo-agent",
-  "source_commit": "65e32f00c8f13e5ec1f8b941639dc37fdc07424c",
-  "benchmarks": "https://llm.ciru.ai/research/ornith-strix/",
-  "target_revision": "10fbf86fed7ecee4a061f8b499a618f46001cac1",
-  "draft_revision": "9b4852c05fd00b672b7434b1bb105bc03c8682b0",
-  "files_bytes": {
-    "INSTALL.md": 5977,
-    "README.md": 18618,
-    "CREDITS.md": 6275,
-    ".gitattributes": 220,
-    "NOTICE": 1088,
-    "BUILD.md": 8422,
-    "NATIVE-BUILD-CHECK.json": 1700,
-    "bundle/paths.env": 268,
-    "bundle/serve.sh": 1070,
-    "runtime/INSTALL-ORNITH-RUNTIME.sh": 794,
-    "runtime/INSTALL-RUNTIME.sh": 4428,
-    "runtime/README.md": 1167,
-    "runtime/UPSTREAM-RUNTIME.md": 14159,
-    "runtime/nixos-module.nix": 349,
-    "runtime/requirements-runtime.lock": 2129,
-    "runtime/runtime-env.sh": 4973,
-    "runtime/shell.nix": 416,
-    "runtime/ciru-halo-agent-aiter-source.tar.gz": 55734195,
-    "runtime/ciru-halo-agent-vllm-source.tar.gz": 40520111,
-    "LICENSES/ornith-upstream-model-card.md": 29870,
-    "LICENSES/dflash2-upstream-model-card.md": 7675,
-    "LICENSES/dflash2-9b4852c05fd0-LICENSE": 11358,
-    "LICENSES/dflash2-main-LICENSE": 11358,
-    "LICENSES/sources.json": 1365,
-    "LICENSES/flash-linear-attention-MIT.txt": 1085,
-    "LICENSES/Apache-2.0.txt": 11358,
-    "runtime/aiter-jit-gfx1151/module_aiter_core.so": 567024,
-    "runtime/wheels/amd_aiter-0.1.0rc1-cp314-cp314-linux_x86_64.whl": 67412360,
-    "runtime/wheels/vllm-0.1.0rc2.dev9+g9255fd9fb9.rocm100-cp314-cp314-linux_x86_64.whl": 44318322,
-    "bundle/LICENSES/LICENSE-APACHE-2.0": 11358,
-    "bundle/models/source-prefill.safetensors": 1761622768,
-    "bundle/native/libornith_attention_iu4.so": 186640,
-    "bundle/native/libornith_dense_g256.so": 69824,
-    "bundle/native/libornith_dense_g256_n32.so": 76848,
-    "bundle/native/libornith_head_i8_tile.so": 128952,
-    "bundle/native/libornith_persistent_iu4.so": 103792,
-    "bundle/native/libornith_routed_direct.so": 156392,
-    "bundle/native/libornith_routed_n32.so": 151928,
-    "bundle/native/libornith_routed_storage_n32.so": 158536,
-    "bundle/packaging/serve.sh": 1474,
-    "bundle/cache/aiter/module_aiter_core.so": 567024,
-    "bundle/plugin-site/ornith_g256/__init__.py": 1368,
-    "bundle/plugin-site/ornith_g256/adaptive_c1.py": 14985,
-    "bundle/plugin-site/ornith_g256/attention.py": 772,
-    "bundle/plugin-site/ornith_g256/attention_compact.py": 10412,
-    "bundle/plugin-site/ornith_g256/attention_fast.py": 12274,
-    "bundle/plugin-site/ornith_g256/attention_folded.py": 9570,
-    "bundle/plugin-site/ornith_g256/attention_iu4.py": 2575,
-    "bundle/plugin-site/ornith_g256/attention_iu4_persistent.py": 12418,
-    "bundle/plugin-site/ornith_g256/attention_mixed.py": 4138,
-    "bundle/plugin-site/ornith_g256/attention_partition.py": 8642,
-    "bundle/plugin-site/ornith_g256/attention_storage.py": 4117,
-    "bundle/plugin-site/ornith_g256/attention_tile.py": 9641,
-    "bundle/plugin-site/ornith_g256/attention_verify.py": 3755,
-    "bundle/plugin-site/ornith_g256/attention_window.py": 3612,
-    "bundle/plugin-site/ornith_g256/cache_full1120.py": 3017,
-    "bundle/plugin-site/ornith_g256/column_backend.py": 5752,
-    "bundle/plugin-site/ornith_g256/column_kernel.py": 9422,
-    "bundle/plugin-site/ornith_g256/config.py": 2869,
-    "bundle/plugin-site/ornith_g256/dense_n32.py": 2449,
-    "bundle/plugin-site/ornith_g256/dense_source.py": 1818,
-    "bundle/plugin-site/ornith_g256/dflash_conv_boundary.py": 5508,
-    "bundle/plugin-site/ornith_g256/dflash_spec.py": 1393,
-    "bundle/plugin-site/ornith_g256/dynamic_graphs.py": 11779,
-    "bundle/plugin-site/ornith_g256/gdn_compact.py": 5202,
-    "bundle/plugin-site/ornith_g256/gdn_compact_kernel.py": 10042,
-    "bundle/plugin-site/ornith_g256/gdn_spec.py": 5696,
-    "bundle/plugin-site/ornith_g256/launch.py": 17783,
-    "bundle/plugin-site/ornith_g256/lifecycle.py": 8269,
-    "bundle/plugin-site/ornith_g256/loader.py": 2933,
-    "bundle/plugin-site/ornith_g256/method.py": 10049,
-    "bundle/plugin-site/ornith_g256/moe_base.py": 3683,
-    "bundle/plugin-site/ornith_g256/native.py": 12196,
-    "bundle/plugin-site/ornith_g256/no_spec.py": 2618,
-    "bundle/plugin-site/ornith_g256/phase_dispatch.py": 2184,
-    "bundle/plugin-site/ornith_g256/prefill_draft.py": 5630,
-    "bundle/plugin-site/ornith_g256/prefix_cache.py": 15261,
-    "bundle/plugin-site/ornith_g256/runtime.py": 6009,
-    "bundle/plugin-site/ornith_g256/worker.py": 13579,
-    "bundle/plugin-site/ornith_g256/worker_base.py": 4601,
-    "bundle/models/draft/LICENSE": 11358,
-    "bundle/models/draft/README.md": 7675,
-    "bundle/models/draft/config.json": 1311,
-    "bundle/models/draft/download-source.json": 1294,
-    "bundle/models/draft/evaluation.json": 8078,
-    "bundle/models/draft/manifest.json": 1599,
-    "bundle/models/draft/model.safetensors": 1052513488,
-    "bundle/models/target/.ciru-source-verified.json": 13078,
-    "bundle/models/target/chat_template.jinja": 7536,
-    "bundle/models/target/config.json": 3438,
-    "bundle/models/target/configuration.json": 58,
-    "bundle/models/target/g256-layer-00.safetensors": 434161416,
-    "bundle/models/target/g256-layer-01.safetensors": 434161416,
-    "bundle/models/target/g256-layer-02.safetensors": 434161416,
-    "bundle/models/target/g256-layer-03.safetensors": 430917376,
-    "bundle/models/target/g256-layer-04.safetensors": 434161416,
-    "bundle/models/target/g256-layer-05.safetensors": 434161416,
-    "bundle/models/target/g256-layer-06.safetensors": 434161416,
-    "bundle/models/target/g256-layer-07.safetensors": 430917376,
-    "bundle/models/target/g256-layer-08.safetensors": 434161416,
-    "bundle/models/target/g256-layer-09.safetensors": 434161416,
-    "bundle/models/target/g256-layer-10.safetensors": 434161432,
-    "bundle/models/target/g256-layer-11.safetensors": 430917392,
-    "bundle/models/target/g256-layer-12.safetensors": 434161432,
-    "bundle/models/target/g256-layer-13.safetensors": 434161432,
-    "bundle/models/target/g256-layer-14.safetensors": 434161432,
-    "bundle/models/target/g256-layer-15.safetensors": 430917392,
-    "bundle/models/target/g256-layer-16.safetensors": 434161432,
-    "bundle/models/target/g256-layer-17.safetensors": 434161432,
-    "bundle/models/target/g256-layer-18.safetensors": 434161432,
-    "bundle/models/target/g256-layer-19.safetensors": 430917392,
-    "bundle/models/target/g256-layer-20.safetensors": 434161432,
-    "bundle/models/target/g256-layer-21.safetensors": 434161432,
-    "bundle/models/target/g256-layer-22.safetensors": 434161432,
-    "bundle/models/target/g256-layer-23.safetensors": 430917392,
-    "bundle/models/target/g256-layer-24.safetensors": 434161432,
-    "bundle/models/target/g256-layer-25.safetensors": 434161432,
-    "bundle/models/target/g256-layer-26.safetensors": 434161432,
-    "bundle/models/target/g256-layer-27.safetensors": 430917392,
-    "bundle/models/target/g256-layer-28.safetensors": 434161432,
-    "bundle/models/target/g256-layer-29.safetensors": 434161432,
-    "bundle/models/target/g256-layer-30.safetensors": 434161432,
-    "bundle/models/target/g256-layer-31.safetensors": 430917392,
-    "bundle/models/target/g256-layer-32.safetensors": 434161432,
-    "bundle/models/target/g256-layer-33.safetensors": 434161432,
-    "bundle/models/target/g256-layer-34.safetensors": 434161432,
-    "bundle/models/target/g256-layer-35.safetensors": 430917392,
-    "bundle/models/target/g256-layer-36.safetensors": 434161432,
-    "bundle/models/target/g256-layer-37.safetensors": 434161432,
-    "bundle/models/target/g256-layer-38.safetensors": 434161432,
-    "bundle/models/target/g256-layer-39.safetensors": 430917392,
-    "bundle/models/target/generation_config.json": 202,
-    "bundle/models/target/head-w8.safetensors": 516505840,
-    "bundle/models/target/merges.txt": 3353259,
-    "bundle/models/target/model.safetensors.index.json": 171704,
-    "bundle/models/target/preprocessor_config.json": 390,
-    "bundle/models/target/processor_config.json": 1191,
-    "bundle/models/target/protected-00.safetensors": 1913075328,
-    "bundle/models/target/protected-01.safetensors": 3842944,
-    "bundle/models/target/protected-02.safetensors": 2784544,
-    "bundle/models/target/protected-03.safetensors": 3838744,
-    "bundle/models/target/protected-04.safetensors": 3842968,
-    "bundle/models/target/protected-05.safetensors": 2784560,
-    "bundle/models/target/protected-06.safetensors": 3838768,
-    "bundle/models/target/protected-07.safetensors": 3842968,
-    "bundle/models/target/protected-08.safetensors": 2784560,
-    "bundle/models/target/protected-09.safetensors": 3838768,
-    "bundle/models/target/protected-10.safetensors": 3842968,
-    "bundle/models/target/protected-11.safetensors": 2784560,
-    "bundle/models/target/protected-12.safetensors": 3838768,
-    "bundle/models/target/protected-13.safetensors": 3842968,
-    "bundle/models/target/protected-14.safetensors": 2784560,
-    "bundle/models/target/protected-15.safetensors": 1690442880,
-    "bundle/models/target/quantize_config.json": 107,
-    "bundle/models/target/tokenizer.json": 12807982,
-    "bundle/models/target/tokenizer_config.json": 16718,
-    "bundle/models/target/video_preprocessor_config.json": 385,
-    "bundle/models/target/vocab.json": 6722759,
-    "assets/ciru-halo-agent.png": 2407299,
-    "VISION.md": 3729,
-    "bundle/serve-vision.sh": 252,
-    "bundle/plugin-site/ornith_g256/graph_phase_guard.py": 6648,
-    "bundle/plugin-site/ornith_g256/parser_contract.py": 4203,
-    "bundle/plugin-site/ornith_g256/strict_qwen/__init__.py": 81,
-    "bundle/plugin-site/ornith_g256/strict_qwen/qwen3.py": 11658,
-    "bundle/plugin-site/ornith_g256/strict_qwen/qwen3_contract.py": 8719,
-    "RUNTIME-FIXES.md": 6843,
-    "ACCURACY-PATCH-1.0.1.json": 7126,
-    "bundle/plugin-site/ornith_g256/runtime_correctness.py": 2151,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/qwen_gdn_linear_attn.py": 78334,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/causal_conv1d.py": 51774,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/gdn_attn.py": 26224,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/gpu_model_runner.py": 345523,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/__init__.py": 79,
-    "bundle/plugin-site/ornith_g256/_vllm_correctness/manifest.json": 1607,
-    "TOOL-SCHEMA-PATCH-1.0.2.json": 7236,
-    "bundle/plugin-site/ornith_g256/native_tool_schema.py": 1897,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/INSTALLER": 2,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/METADATA": 254,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/RECORD": 5595,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/REQUESTED": 0,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/uv_build.json": 2,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/top_level.txt": 12,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/direct_url.json": 105,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/uv_cache.json": 137,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/WHEEL": 91,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/entry_points.txt": 121,
-    "bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/licenses/LICENSE-APACHE-2.0": 11358
+  "release": "Ciru v4",
+  "version": "4.0.0",
+  "runtime_version": "4.0.0",
+  "tool_component_version": "1.0.3",
+  "model": "Ornith1.5",
+  "status": "QUALIFIED_FOR_PUBLICATION",
+  "publication_ready": true,
+  "qualified_build_sha256": "c970d175624455ca30dd3efbc6f52225f6ad32595764232050342a23ced99d94",
+  "hf_parent_revision": "f6f9d19e120cc42a8680b1c1d99470e24a83a4fe",
+  "weights_modified": false,
+  "pinned_engine_wheels_modified": false,
+  "validation": {
+    "runtime_build": "c970d175624455ca30dd3efbc6f52225f6ad32595764232050342a23ced99d94",
+    "native_performance_requests": 97,
+    "coding_health_passes": 90,
+    "coding_health_tasks": 90,
+    "native_tool_calls": 18,
+    "complete_agent_episodes": 3,
+    "reports": [
+      "evaluation/CIRU-V4-PERFORMANCE.md",
+      "evaluation/CIRU-V4-FIDELITY32.md",
+      "evaluation/CIRU-V4-FIDELITY64.md"
+    ]
   },
-  "native_sha256": {
-    "libornith_attention_iu4.so": "2905806824bece62ce9d9140608859df869f3e3d3663c67bf79ad4ea495cd61f",
-    "libornith_dense_g256.so": "855838687084170807c5ba4556e9771f9b1b26f3556f48c0a16dbd8bbee41c9c",
-    "libornith_dense_g256_n32.so": "24dbf9b1eaff06ff6530e4f37292347bcf073b3c07f7e16ea6ab111711162a40",
-    "libornith_head_i8_tile.so": "2c03de5ed015a052b407436c7e42b50fa33ef76cb46dccab9b226546ca1d1e09",
-    "libornith_persistent_iu4.so": "40e5cf35157c3e716de8067f7eef1312a3b9e787dd52a6f2e232ccd2d07fa68c",
-    "libornith_routed_direct.so": "5b3ca6826547e0ff1c417e733229e104721cec403ccd8f88f7fcdb1726d6f4ed",
-    "libornith_routed_n32.so": "48415e71f6c85e3aefcecf99d72607bb05900244b4aaa9a7ba2c58378d92f5fb",
-    "libornith_routed_storage_n32.so": "3ea288cd0e37a4c373038744c3cc4cd4eac09675ac137e953d5ae7a6cd43dd3d"
-  },
-  "checks": {
-    "clean_runtime_install": "passed on Ciru",
-    "release_launcher_dry_run": "passed using new runtime",
-    "native_source_build": "8/8 compiled; distributed measured libraries unchanged",
-    "new_gpu_inference_run": false
-  },
-  "image_sha256": "ddbb4819065a652a6feb20ce760c8292b546800cd2e48d01c560809ad5b58765",
-  "optional_vision": {
-    "launcher": "bundle/serve-vision.sh",
-    "encoder_projector": "bundle/models/target/protected-00.safetensors",
-    "format": "native BF16 safetensors; no GGUF mmproj",
-    "vision_tensors": 333,
-    "images_per_request": 1,
-    "image_max_pixels": 1048576,
-    "encoder_attention_backend": "TRITON_ATTN",
-    "validation": "Published weight header and launcher settings checked; earlier image-profile results documented separately in VISION.md",
-    "historical_profile": true
-  },
-  "draft_depth_override": {
-    "flag": "--draft-tokens",
-    "choices": [
-      0,
-      7,
-      15
-    ],
-    "scope": "C1 at all supported context lengths; C2-8 remains 7",
-    "validation": "CPU-only settings and policy tests; no model loaded"
-  },
-  "runtime_fixes": {
-    "date": "2026-09-13",
-    "notes": "RUNTIME-FIXES.md",
-    "weights_changed": false,
-    "runtime_wheels_changed": false,
-    "native_libraries_changed": false,
-    "changed_files_sha256": {
-      "bundle/plugin-site/ornith_g256/__init__.py": "a15cbb0679a9cc4766a77c57594e941e9a911381f42c815cfdecb189c8b6635c",
-      "bundle/plugin-site/ornith_g256/graph_phase_guard.py": "6cabaf7d4e3be04bbac7a642544d38948de2e07ea7ca38f5927f91bf3ad39163",
-      "bundle/plugin-site/ornith_g256/launch.py": "8c17d8a738ec8b4424a1e451883cd15bd0a480525532de4c05c4b3e2ac8f22bb",
-      "bundle/plugin-site/ornith_g256/parser_contract.py": "2bc5885394786931b2c00693f580ee99da220c6adffc00651fc8227ca9bb1d0d",
-      "bundle/plugin-site/ornith_g256/strict_qwen/__init__.py": "ee8a5055249f1413fc22f46421ab77b5485aa39b199c68eb5b56c75dcc40922d",
-      "bundle/plugin-site/ornith_g256/strict_qwen/qwen3.py": "33c653447a144b72257e508ccd92610db6d30578e28f98c69c99281773dead1c",
-      "bundle/plugin-site/ornith_g256/strict_qwen/qwen3_contract.py": "74a0b03e3532883fa4d70494fbaecae7760e293cbb6b4e588640884f3093052a",
-      "bundle/plugin-site/ciru_ornith_g256-0.0.2a0.dist-info/RECORD": "6605e7606bb5c9f323d2984473ad643ba747616e78c3ebea6938fa872d531f9b",
-      "README.md": "83af97ae9c3ab0ea1185df77aa3436d0792ca730e89268fbaef5aea8b5f6b19a",
-      "RUNTIME-FIXES.md": "084dfd99a09d31a7b4ee30d17087a9cc1d341ab1fced5da644fc9e4c0d6144f9"
-    },
-    "checks": {
-      "shipped_runtime_guard_and_parser_import": "passed without GPU initialization",
-      "shipped_runtime_launcher_argument_inspection": "passed using existing released weight paths",
-      "prior_ornith_gpu_requests": "21/21 Hermes on Strix Halo; includes HE0-9 10/10",
-      "final_parser_saved_response_replays": "6/6 ASGI and 3/3 Pi",
-      "new_public_package_gpu_run": false
-    },
-    "native_parameter_whitespace": "Original Qwen XML wrapping-newline convention retained; exact boundary whitespace not guaranteed."
-  },
-  "runtime_version": "1.0.2",
-  "release_date": "2026-09-15",
-  "accuracy_patch": {
-    "version": "1.0.1",
-    "details": "ACCURACY-PATCH-1.0.1.json",
-    "weights_changed": false,
-    "native_libraries_changed": false,
-    "fixes": [
-      "FP32 products in BF16 causal convolution",
-      "Accepted GDN state and convolution history recovery",
-      "Persistent accepted-state source/count buffers for FULL graphs"
-    ],
-    "single_request_throughput_tradeoff": "Approximately 7% lower in the Ornith screen, accepted by the owner for this correctness release."
-  },
-  "runtime_accuracy_source_commit": "2c50cb2aa31abd0537e9c84b7ec43307e3ea40c7",
-  "runtime_tool_schema_source_commit": "65e32f00c8f13e5ec1f8b941639dc37fdc07424c",
-  "tool_schema_patch": {
-    "version": "1.0.2",
-    "details": "TOOL-SCHEMA-PATCH-1.0.2.json",
-    "default_auto_constraints": true,
-    "prompt_visible_tools_unchanged": true,
-    "weights_changed": false,
-    "native_libraries_changed": false,
-    "runtime_wheels_changed": false
-  },
-  "vision": {
-    "enabled_by_default": true,
-    "upstream_supported": true,
-    "bf16_tensors": 333,
-    "image_limit": 1,
-    "image_max_pixels": 1048576,
-    "encoder_backend": "TRITON_ATTN",
-    "text_only_flag": "--text-only",
-    "live_image_checks_passed": 2,
-    "details": "VISION.md"
+  "payload_sha256": {
+    ".gitattributes": "e8738eab56c30a9c4e6b144e2314425c6cf04e66d46b3a3a44e4d1da65e18ea4",
+    "INSTALL.md": "2878687e04890121e3bbbf280944f33f53b0c373b921d85bc4493b1c35552bca",
+    "README.md": "fa570a0fe074ce2978953a4840852749c77682051f58b2c00181e82c6ff86aef",
+    "RUNTIME-FIXES.md": "9c4fe97a10478429a405e853623d2c0281fe3350227f0505fd4b7ce69d94caf7",
+    "SOURCE-PROVENANCE.json": "dde1551f66d56c18610f541a18b27aabb7932a31c90d27a4390d9d39b7f82776",
+    "assets/ciru-halo-agent.png": "ddbb4819065a652a6feb20ce760c8292b546800cd2e48d01c560809ad5b58765",
+    "bundle/LICENSES/XGRAMMAR-LICENSE": "c71d239df91726fc519c6eb72d318ec65820627232b2f796219e87dcf35d0ab4",
+    "bundle/native/libornith_attention_iu4.so": "3d29cf3616c69f9fe2b89f8d3564d479dda682fdc06572886e7b60faccc9a53a",
+    "bundle/native/libornith_dense_g256.so": "dc9bbf04f02f951fb444feaea2f612671867aad44929a30fdfadef8e5d5513a9",
+    "bundle/native/libornith_dense_g256_n32.so": "db3e265fca6daf51fca412579160db7f98d81cd295fffd54a0c7edc3f5fd45da",
+    "bundle/native/libornith_head_i8_tile.so": "2c03de5ed015a052b407436c7e42b50fa33ef76cb46dccab9b226546ca1d1e09",
+    "bundle/native/libornith_persistent_iu4.so": "c70ba14d5aebf67b761ea54cfc9a38f8bbedf74cfca8b19050048acda9c8a06c",
+    "bundle/native/libornith_routed_direct.so": "5b3ca6826547e0ff1c417e733229e104721cec403ccd8f88f7fcdb1726d6f4ed",
+    "bundle/native/libornith_routed_n32.so": "48415e71f6c85e3aefcecf99d72607bb05900244b4aaa9a7ba2c58378d92f5fb",
+    "bundle/native/libornith_routed_storage_n32.so": "2002bf6f82a637f5842f7c5f09a09f3b7a0db55edc4b3e35181c8f9a12ffb4b1",
+    "bundle/packaging/serve.sh": "78ddf77caf1208aeacbf9fa6ae164a026c3ef1da1bdd43015cdd733a1f1b3b64",
+    "bundle/plugin-site/TOOL-INTEGRATION.json": "bc116f818a397f3f883e9209d4848999a5be51f99689ff2a4063c7f221a464d9",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/INSTALLER": "e6184ce10e266134fdcfa401e8f1a95005bcd4f18d16b62b757323e2833fe9a9",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/METADATA": "c0368af5f84a4158453b069c59491aaafe84386bfa5bdaa3901a8a381d648781",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/RECORD": "591c8a6330b025ab1a91b112d6e6ebc7b7160291d9a1f85a0c3835d398b30449",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/REQUESTED": "e3b0c44298fc1c149afbf4c8996fb92427ae41e4649b934ca495991b7852b855",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/WHEEL": "4a63b1614ee9ccd281a804af409ec38d7dd719417dda5ac684c6f747afe28aa2",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/entry_points.txt": "1127b5c28a5d92219faf90cd5997127a69603ebdeebd1aded4ec87925f3c42a5",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/licenses/LICENSE-APACHE-2.0": "cfc7749b96f63bd31c3c42b5c471bf756814053e847c10f3eb003417bc523d30",
+    "bundle/plugin-site/ciru_ornith_g256-4.0.0.dist-info/top_level.txt": "11d08c2d79e7f2d0cd05447726190cbce8bc844c3a22e85815811c7dd85e9315",
+    "bundle/plugin-site/ornith_g256/__init__.py": "60abbb51a475487e3abcb9e5b34d8836b4d4b36e0bced062d2382dc7b1076cf6",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/__init__.py": "963097106407a87b62b275041516144de15713dee65108accc29e27036ffc059",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/causal_conv1d.py": "230f3e554f5f1d609a37022578c84e42d5c71d46e684101278a7e21af4f7d418",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/chat_completion_serving_r03.py": "5720371f5bcb351e340a3e8d9a0235ec01ff27ed93e4f35542e31b83a055ef45",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/gdn_attn.py": "a3c361d502e8bca2aa46a0baed5ec753a3c62e3d027a298e3ed4692794f9f6ec",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/gpu_model_runner.py": "ef0009acaa67a1a5e066984f34bf619b06e85220ff953a3d9511376f0146dccf",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/manifest.json": "bdef4ca7151a7a5e0410dd3e8bcd2e8a85f125938ec104451cb8777267c3c4d9",
+    "bundle/plugin-site/ornith_g256/_vllm_correctness/qwen_gdn_linear_attn.py": "3e58fdddc7a0b42ae3d9b5eabb375a993a4110bad499c37846bde466901dbeb3",
+    "bundle/plugin-site/ornith_g256/_xgrammar_native/build-receipt.json": "3bdf77657e5157c35938b66dfb46266608ade3b599a62ce74abb3f807d579e84",
+    "bundle/plugin-site/ornith_g256/_xgrammar_native/libxgrammar-candidate.so": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7",
+    "bundle/plugin-site/ornith_g256/_xgrammar_native/selection.json": "7cb93b3fe902de6d604bf96f79fdc3efbd8d612e494da838f6a8f0d731eb89ae",
+    "bundle/plugin-site/ornith_g256/adaptive_c1.py": "8e48fd874fbdf65183c20bcc7c748d2562dbcafa97e004f18a2818b98980a809",
+    "bundle/plugin-site/ornith_g256/attention.py": "2fdd3c0ed84dcb052a9e8f9df4068ddf80a969931bea0a295e022d36b753cda1",
+    "bundle/plugin-site/ornith_g256/attention_compact.py": "50df61a02a3890a2825b205ed18c7d18a2b002333d382c572440b58ca9b2b9fe",
+    "bundle/plugin-site/ornith_g256/attention_fast.py": "27b1a133cadb87ff3a63c027049e1da81e4112fd5d5260a4440fde85491f838a",
+    "bundle/plugin-site/ornith_g256/attention_folded.py": "73dac8203032bce6ceedf213d7eefe5987c921768ae237129e232fefd8d9bba3",
+    "bundle/plugin-site/ornith_g256/attention_iu4.py": "15c5341fc848c43ac956546e563cf8ab7f7649f05a674fb0153b0e2d76438cae",
+    "bundle/plugin-site/ornith_g256/attention_iu4_persistent.py": "cef74ae423974063212637a49cd64a5a10b9077f3a5e2287746722749a1e6297",
+    "bundle/plugin-site/ornith_g256/attention_mixed.py": "cbb530697f2555be6532b7d4d3401f3a33c4c58982659e88892720d21691110a",
+    "bundle/plugin-site/ornith_g256/attention_partition.py": "41bfb44dd0d0c46a2c0a725983155018e051cd89edb8527f1195a092bef7ffdd",
+    "bundle/plugin-site/ornith_g256/attention_storage.py": "3ab813e221835370bdb503b2ab7d48f46294441b1dd9c652c5dc976438bdb6e2",
+    "bundle/plugin-site/ornith_g256/attention_tile.py": "f2bf007132bc712c8e61e036ea8a94b54f18a813d9bdd334eee46436a3ca1fea",
+    "bundle/plugin-site/ornith_g256/attention_verify.py": "3a329599694265f4badecc2f9b709b280d9960b967d7192f3959ad9e0d8c6077",
+    "bundle/plugin-site/ornith_g256/attention_window.py": "78534f9a1b0beae9ecf80cc5b2a613e4fe4827036c8df0acc7e4dfd956aae1c5",
+    "bundle/plugin-site/ornith_g256/cache_full1120.py": "3d4240a749d9671c0bf9b3937200e1cd8d4abfe6699d8c1d6857b188c2ef8a8c",
+    "bundle/plugin-site/ornith_g256/column_backend.py": "a035b839369e733b05f68ce83df0d022cb62455016fccf30c2978127ab3c5403",
+    "bundle/plugin-site/ornith_g256/column_kernel.py": "27fb053d970184732a46212026afd3152ec07a0dfc3dc4c5cc5cc0d539f17c81",
+    "bundle/plugin-site/ornith_g256/config.py": "23210d8b625ef375638e2d6043a92a26181a5b5f1c537d9d4b34c91b7109f36f",
+    "bundle/plugin-site/ornith_g256/dense_n32.py": "119fb7d25e9d28b6e520b5285c10de86c7d8ec2fe29a2399e9c086173d765903",
+    "bundle/plugin-site/ornith_g256/dense_source.py": "cdc7dd6a13c2439cd405596069c301fe28cb95ff32278f1bafe8118fe357d48e",
+    "bundle/plugin-site/ornith_g256/dflash_conv_boundary.py": "a5f46458194cade5ddcbfbb51af30c662d12b9086fc459ecbeb17463339febb1",
+    "bundle/plugin-site/ornith_g256/dflash_spec.py": "c948b99eab22d2af5c5fc0ca9f764a0855a3d86aca0609362bf5474a9beff1b3",
+    "bundle/plugin-site/ornith_g256/dynamic_graphs.py": "6f51f0dd740f6546829e46bbbe840df2780ce8c49978658eb3a05092554ff90c",
+    "bundle/plugin-site/ornith_g256/gdn_compact.py": "8b5c062d10f386d78bc9926ab0576e56359401e5d1621c397a2922e2b426ccab",
+    "bundle/plugin-site/ornith_g256/gdn_compact_kernel.py": "0eaf34a2eccef2dabbc9e4f017e6ce84317a28efdcbabecf92b3b4c4c9429a55",
+    "bundle/plugin-site/ornith_g256/gdn_spec.py": "0931b1b3601d935bb6acdfcb65d307dae4105645cd387bd9b256dd7a525376a3",
+    "bundle/plugin-site/ornith_g256/graph_phase_guard.py": "6cabaf7d4e3be04bbac7a642544d38948de2e07ea7ca38f5927f91bf3ad39163",
+    "bundle/plugin-site/ornith_g256/launch.py": "c2854b86e2f97dd2ea0356b4b333922f62b9b3d1b3e71f7993d803c6d0957479",
+    "bundle/plugin-site/ornith_g256/lifecycle.py": "3285552177cb8ce8a0149a01f8168d2bc44a739a9395ca84e15acccbb20a0b73",
+    "bundle/plugin-site/ornith_g256/loader.py": "a853c6d014214a7adf298548f9175b6e18dd05338917aec16ebcb4fc3c3ecb13",
+    "bundle/plugin-site/ornith_g256/method.py": "fc301af72efd234430dbfb839d56e74328457b5b83e14a044aeb9554d8acc707",
+    "bundle/plugin-site/ornith_g256/moe_base.py": "26fb517415f2497859f2e3ff1aed8f31e7c56a4aacccfb34c6638e65d7df07c4",
+    "bundle/plugin-site/ornith_g256/native.py": "49cce0d67987b9659feb556aafb247423490c2153b7d77b750f055a4fdcda331",
+    "bundle/plugin-site/ornith_g256/native_tool_schema.py": "910060689d72b47ad4a7bb1e51ef301e41ff0be1beb9cc6b3df01c423d2edef5",
+    "bundle/plugin-site/ornith_g256/no_spec.py": "561106b75328fe0e161fa1243aca4cf50c5049d8e785d22c90aacb3e1e1a0104",
+    "bundle/plugin-site/ornith_g256/optimized/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "bundle/plugin-site/ornith_g256/optimized/compile_scope.py": "18971a73b3a5ad80e6212de7f7e419a1dd00eb7380a5e6c8559f5ac6f5ec8d4d",
+    "bundle/plugin-site/ornith_g256/optimized/draft.py": "e801301cf03d69998758d35fb5463754ccc30a18e030ed2b988b26cef43f429c",
+    "bundle/plugin-site/ornith_g256/optimized/draft_kernel.py": "3d6da6ef0c144505f888cbd35ba50b6b95ad8e48c376ab73297746bff228919e",
+    "bundle/plugin-site/ornith_g256/optimized/grammar_parallel.py": "e9c3608438712521980d7fcd60f58862d0b0a2876620183f304c64a03817edea",
+    "bundle/plugin-site/ornith_g256/optimized/greedy.py": "386443ad8729f8fe21ec5fef84ead3acfbd4d6c5dd3b4c1e50a7995ed1d1761f",
+    "bundle/plugin-site/ornith_g256/optimized/greedy_kernel.py": "2ebe2a25d1011a45eea7eda3c0964ca1eb2638df7d60f54db02c91435550dc42",
+    "bundle/plugin-site/ornith_g256/optimized/launch.py": "c1d42602154674916a94b9907b7746fac3ea48cc505fa69babf116f0f127e03d",
+    "bundle/plugin-site/ornith_g256/optimized/mixed_qk_sum.py": "17b18769ae84c2e9707a0257181e5c314e9774fc3b4796f8e6787a8302cb5e1f",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/SOURCE-PROOF.json": "5939bc8a53a831715d05d0fa0290beaa0cb00eda1597fe3bf1f1cfc53699fd0d",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/adapter.py": "6af97a8d5e0d79b7dfc32f2f0fe5f5c53b5c1c50d03542ff99c64cae7f3909e1",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/causal_conv1d.py": "254c540f77b01694696a1ab535abd0e5a23d54e7211e202cf9beead04face16b",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_o.py": "a93f951048958265b4d1c3f9e67341329196e7a57525433934215fde4b16e61b",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_scaled_dot_kkt.py": "71735d2b1c4a19eecbb6058708481817bb610426c48f96a5f646784afcf0c65c",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/solve_tril.py": "bdd04b69c0f91562048d25834bd53d298dc1f1bd068b9494af88776b2bf4b1d1",
+    "bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/wy_fast.py": "4cc0be519bfcf3812d810a4bebb3dd45fb9b13a8d0ab62d0ad463a1a2e36f61f",
+    "bundle/plugin-site/ornith_g256/optimized/postconv.py": "c58ee5d2b8418fbf69eda5d9ee4f342b11cc012ab77ae25b6ce2b2aea51aa50f",
+    "bundle/plugin-site/ornith_g256/optimized/prefix.py": "c695a899ab12b58c3b89b6ecc276e82efa3f0833a319911a8528cc4e237710cd",
+    "bundle/plugin-site/ornith_g256/optimized/qk.py": "d86f1a94086668ac6402eaec7a19c0ad8e55dc8e6377f8d4a9eeab6aa96f9664",
+    "bundle/plugin-site/ornith_g256/optimized/qk_schedule.py": "66038781f0667486fba458ef3118ca89ad9640dc2c2bd05cf8576440f0f178a2",
+    "bundle/plugin-site/ornith_g256/optimized/worker.py": "38579b767235e7fcac3f2e6b0507438f155c134e1e85f3d1c205a43bece95512",
+    "bundle/plugin-site/ornith_g256/parser_contract.py": "2cebab403dc882569f3ce67d12798df14cebd5c162c72f82becd37bcdd7e2a28",
+    "bundle/plugin-site/ornith_g256/phase_dispatch.py": "1a4ebdf2af709e6ca519e0a08f6a537602df178ca575332045617631640f8137",
+    "bundle/plugin-site/ornith_g256/prefill_draft.py": "d79bb29a124639543a64e198908e2eaef4cbcf492f2e5e14debcfc9c867816f5",
+    "bundle/plugin-site/ornith_g256/prefix_cache.py": "65fb972e285e37741f53b221d2f6380ec464d34a39f36bb66a82dcaacab727d0",
+    "bundle/plugin-site/ornith_g256/runtime.py": "db43b1c3766720e0dbee4e8617aee673a533c30d36fcc1fbf43919717d8ab149",
+    "bundle/plugin-site/ornith_g256/runtime_correctness.py": "08f41773e623a4f99cf8ab19e434d7e33bac9cfa5eb05a9dc0f4511aab8ba4bc",
+    "bundle/plugin-site/ornith_g256/strict_qwen/__init__.py": "ee8a5055249f1413fc22f46421ab77b5485aa39b199c68eb5b56c75dcc40922d",
+    "bundle/plugin-site/ornith_g256/strict_qwen/qwen3.py": "33c653447a144b72257e508ccd92610db6d30578e28f98c69c99281773dead1c",
+    "bundle/plugin-site/ornith_g256/strict_qwen/qwen3_contract.py": "29c083659c54c81144fdb6c8329c75aaa62cf69d1c45fa7da8bdeeeca4856eaf",
+    "bundle/plugin-site/ornith_g256/worker.py": "acc27240c21f3478ca452ddf04244df4ee770efaac38192ef64d328ef22701d5",
+    "bundle/plugin-site/ornith_g256/worker_base.py": "a00ef1295172e655c8b3c7d790cfd0316114a9492e483dedfdea036d3b7bf868",
+    "bundle/plugin-site/ornith_g256/xgrammar_native.py": "d19c546e124cb11311e4cc31f5833dc663e2789d45090091c0f8fec8a9626623",
+    "bundle/serve-vision.sh": "21dd61d76c9d82762480f1f8c0364f9692f1b01c8405c79347c1e2c35d931cd6",
+    "bundle/serve.sh": "2950b201e971628425547f13ba68b4a628f230ace612eb1bfe7c9b7b8ead5aea",
+    "evaluation/CIRU-V4-C1-SPEED.md": "19cf7335f60272e13efa4ff94bf7ec2ff79c61ad4d694911bbc6e779e852f790",
+    "evaluation/CIRU-V4-COMPOSITION.json": "ed41fe291530d9c7f9dbda1c9db2b066fc6c1678b81787091597699ae6e62a68",
+    "evaluation/CIRU-V4-FIDELITY32.md": "9858cc15e6b1a767cdaf772587b35abaf1053f5cd39acc18080e425b676cbff9",
+    "evaluation/CIRU-V4-FIDELITY64.md": "149143f2c030d138d4e2c5e4d042401b0b6c47f0f8d7fb8eaefdbd796607c8a7",
+    "evaluation/CIRU-V4-PERFORMANCE.md": "0e953a5167ec8b7d9037c73be018e0c433bd15e23e28472aac3ace17b99cf310",
+    "evaluation/CIRU-V4-SHORT-FIDELITY.md": "84de24ead0949eb28b9697a8f19620312df249456aeddd117cdb3ff434704f1c",
+    "runtime/ciru-v4-ornith-source.tar.gz": "e029f28d4bb89c801437834d2fa0109c7a3c906580ec31e39c053579d6859da5"
   }
 }

RUNTIME-FIXES.md

@@ -1,5 +1,41 @@
 # Ciru runtime patch notes
 
+## Ciru v4 (4.0.0) — September 17, 2026
+
+Continues the Ciru v3 model release line. The earlier 1.0.x identifiers below
+describe runtime patches; R03/R04 component identity 1.0.3 remains provenance,
+while this complete public release is 4.0.0.
+
+- Integrate the retained output32 expert tile and all six qualified GDN
+  traversal changes, including the older Apodex campaign work.
+- Use Q4480 prefill chunks with the required wide-capacity dense/expert
+  libraries and compiler adapters, plus the qualified prefix-cache repair.
+- Include the persistent-attention correctness repair, exact Q/K fusion,
+  fused attention prefill, draft convolution/residual/RMSNorm repair and
+  greedy draft-path fusion. Preserve the default adaptive DFlash policy.
+- Integrate strict-grammar scheduling, the native XGrammar finite-repetition
+  and shared-memory accounting changes, and continuous usage reporting during
+  buffered tool streams. Preserve natural finish reasons and explicit errors
+  for malformed output.
+- Preserve all target, draft and BF16 vision weights, the native chat template,
+  sampler, context/cache limits and previous correctness overlays.
+- Make the public launcher select the complete optimized worker and use
+  relative installation/model defaults when no paths.env override exists.
+
+Update **both** native libraries and the complete plugin/launchers using
+[INSTALL.md](INSTALL.md). The pinned vLLM/AITER wheels are retained.
+
+Final native performance: cold prefill +26.23% at64K and +30.33% near256K;
+cached C1 decode +19.74% at63K and +26.40% near256K. Short C2/C4 aggregate
+throughput regressed, and near256K C8 first-token latency increased. Numerical
+fidelity results and the complete speed comparison are reported on the model card.
+All individual historical gains overlap and must not be added.
+
+The final native performance block completed97 measured requests with90/90
+coding health checks. Native tool integration passed18 exact live calls and
+preserved identical333-token streams across usage modes. See the [Ciru v4 results](README.md#ciru-v4-results).
+
+
 ## Runtime 1.0.2 — September 15, 2026
 
 - Enable native vLLM/XGrammar schema constraints for automatic tool calls when
@@ -43,7 +79,7 @@ all long-generation quality issues are resolved.
 
 This update prevents a reproduced cache-corruption crash and makes malformed tool output fail explicitly. It updates the serving plugin; the released weights, native libraries, runtime wheels, sampler, context pool and adaptive DFlash2 policy are unchanged.
 
-## Update an existing installation
+## Historical September 13 update instructions
 
 Stop your server before replacing its plugin files. Run this from the directory containing `bundle/` and `runtime/`, then restart your usual launcher:
 

SOURCE-PROVENANCE.json

@@ -0,0 +1,159 @@
+{
+  "release": "Ciru v4",
+  "package_version": "4.0.0",
+  "qualified_build_sha256": "c970d175624455ca30dd3efbc6f52225f6ad32595764232050342a23ced99d94",
+  "native": [
+    {
+      "source": "kernels/dense/dense_g256_consumer.cpp",
+      "library": "libornith_dense_g256.so",
+      "source_sha256": "28fe913a6fd5b3e97ae106563a3fe843cf55ba84413a989a7e414ec3357a8150",
+      "measured_library_sha256": "dc9bbf04f02f951fb444feaea2f612671867aad44929a30fdfadef8e5d5513a9"
+    },
+    {
+      "source": "kernels/routed/routed_direct_consumer.cpp",
+      "library": "libornith_routed_direct.so",
+      "source_sha256": "aea3b68893e7e2a8b70ddb1f83a24c444cae6ff8b995ff9794dc545da447a0ff",
+      "measured_library_sha256": "5b3ca6826547e0ff1c417e733229e104721cec403ccd8f88f7fcdb1726d6f4ed"
+    },
+    {
+      "source": "kernels/head/head_i8_tile_consumer.cpp",
+      "library": "libornith_head_i8_tile.so",
+      "source_sha256": "0c14089990d13b63a61fcaa598de2a27d34cae2624ab3c62dd6095e0cb552296",
+      "measured_library_sha256": "2c03de5ed015a052b407436c7e42b50fa33ef76cb46dccab9b226546ca1d1e09"
+    },
+    {
+      "source": "kernels/routed_n32/routed_n32_consumer.cpp",
+      "library": "libornith_routed_n32.so",
+      "source_sha256": "f045389d4cb804fe43d34425ce6c02e817a89bd8cbdb453547f20e36f060be52",
+      "measured_library_sha256": "48415e71f6c85e3aefcecf99d72607bb05900244b4aaa9a7ba2c58378d92f5fb"
+    },
+    {
+      "source": "kernels/routed_storage/routed_storage_n32_consumer.cpp",
+      "library": "libornith_routed_storage_n32.so",
+      "source_sha256": "171dacf56f553dd586dc0bb559b60108c225e93787044c5ee5118c898a5382ea",
+      "measured_library_sha256": "2002bf6f82a637f5842f7c5f09a09f3b7a0db55edc4b3e35181c8f9a12ffb4b1"
+    },
+    {
+      "source": "kernels/attention/prefill.cpp",
+      "library": "libornith_attention_iu4.so",
+      "source_sha256": "7d0a21917fbeff75811d7a7ed9df9438852762c28914eb61ec3b1ba31fc66dae",
+      "measured_library_sha256": "3d29cf3616c69f9fe2b89f8d3564d479dda682fdc06572886e7b60faccc9a53a"
+    },
+    {
+      "source": "kernels/attention/persistent.cpp",
+      "library": "libornith_persistent_iu4.so",
+      "source_sha256": "b430bee407b2c00ee7c2e75dcc39b8dab504767219755c24259f42327975b055",
+      "measured_library_sha256": "c70ba14d5aebf67b761ea54cfc9a38f8bbedf74cfca8b19050048acda9c8a06c"
+    },
+    {
+      "source": "kernels/dense_n32/dense_g256_consumer.cpp",
+      "library": "libornith_dense_g256_n32.so",
+      "source_sha256": "92ff49dc74cdc54a54cbbc2e618ca312f7b17ba502e8c1e83fa2a53570c4567f",
+      "measured_library_sha256": "db3e265fca6daf51fca412579160db7f98d81cd295fffd54a0c7edc3f5fd45da"
+    }
+  ],
+  "native_source_files": {
+    "kernels/attention/persistent.cpp": "b430bee407b2c00ee7c2e75dcc39b8dab504767219755c24259f42327975b055",
+    "kernels/attention/prefill.cpp": "7d0a21917fbeff75811d7a7ed9df9438852762c28914eb61ec3b1ba31fc66dae",
+    "kernels/dense/dense_g256_api.h": "487320b066d2b033982a02cd8e94b3e2075dbec9f7bdb2f7c53b895a060813a2",
+    "kernels/dense/dense_g256_consumer.cpp": "28fe913a6fd5b3e97ae106563a3fe843cf55ba84413a989a7e414ec3357a8150",
+    "kernels/dense_n32/dense_g256_api.h": "487320b066d2b033982a02cd8e94b3e2075dbec9f7bdb2f7c53b895a060813a2",
+    "kernels/dense_n32/dense_g256_consumer.cpp": "92ff49dc74cdc54a54cbbc2e618ca312f7b17ba502e8c1e83fa2a53570c4567f",
+    "kernels/head/head_i8_tile_api.h": "ad11ace0f655b8a883d8888f07f214cb75ba5ff12471ec50a54bd2601f572463",
+    "kernels/head/head_i8_tile_consumer.cpp": "0c14089990d13b63a61fcaa598de2a27d34cae2624ab3c62dd6095e0cb552296",
+    "kernels/routed/portable_exp.h": "d118a773477b611a0f0eb53c1b403479d0c2d35b02e9970ba3c24bb882da8781",
+    "kernels/routed/routed_direct_api.h": "ac2c0076a7db7570f91eed658c8b4646218d98d2dbdb39e84d07c6284ab271d0",
+    "kernels/routed/routed_direct_consumer.cpp": "aea3b68893e7e2a8b70ddb1f83a24c444cae6ff8b995ff9794dc545da447a0ff",
+    "kernels/routed/routed_direct_prefill.h": "0c6a9be68211d3c0fd7c404a7cccaba93a58e996d4326a1e29122c90abe3b5f8",
+    "kernels/routed_n32/portable_exp.h": "d118a773477b611a0f0eb53c1b403479d0c2d35b02e9970ba3c24bb882da8781",
+    "kernels/routed_n32/routed_direct_api.h": "ac2c0076a7db7570f91eed658c8b4646218d98d2dbdb39e84d07c6284ab271d0",
+    "kernels/routed_n32/routed_direct_prefill.h": "0c6a9be68211d3c0fd7c404a7cccaba93a58e996d4326a1e29122c90abe3b5f8",
+    "kernels/routed_n32/routed_n32_consumer.cpp": "f045389d4cb804fe43d34425ce6c02e817a89bd8cbdb453547f20e36f060be52",
+    "kernels/routed_storage/portable_exp.h": "d118a773477b611a0f0eb53c1b403479d0c2d35b02e9970ba3c24bb882da8781",
+    "kernels/routed_storage/routed_direct_api.h": "ac2c0076a7db7570f91eed658c8b4646218d98d2dbdb39e84d07c6284ab271d0",
+    "kernels/routed_storage/routed_storage_n32_consumer.cpp": "171dacf56f553dd586dc0bb559b60108c225e93787044c5ee5118c898a5382ea",
+    "kernels/routed_storage/routed_storage_n32_prefill.h": "15e5f92d7c3f06a05010bb997ee07556007122858c04332ceb794a4c3d1a2bf9"
+  },
+  "plugin_sha256": {
+    "plugin-site/ornith_g256/__init__.py": "60abbb51a475487e3abcb9e5b34d8836b4d4b36e0bced062d2382dc7b1076cf6",
+    "plugin-site/ornith_g256/_vllm_correctness/__init__.py": "963097106407a87b62b275041516144de15713dee65108accc29e27036ffc059",
+    "plugin-site/ornith_g256/_vllm_correctness/causal_conv1d.py": "230f3e554f5f1d609a37022578c84e42d5c71d46e684101278a7e21af4f7d418",
+    "plugin-site/ornith_g256/_vllm_correctness/chat_completion_serving_r03.py": "5720371f5bcb351e340a3e8d9a0235ec01ff27ed93e4f35542e31b83a055ef45",
+    "plugin-site/ornith_g256/_vllm_correctness/gdn_attn.py": "a3c361d502e8bca2aa46a0baed5ec753a3c62e3d027a298e3ed4692794f9f6ec",
+    "plugin-site/ornith_g256/_vllm_correctness/gpu_model_runner.py": "ef0009acaa67a1a5e066984f34bf619b06e85220ff953a3d9511376f0146dccf",
+    "plugin-site/ornith_g256/_vllm_correctness/manifest.json": "bdef4ca7151a7a5e0410dd3e8bcd2e8a85f125938ec104451cb8777267c3c4d9",
+    "plugin-site/ornith_g256/_vllm_correctness/qwen_gdn_linear_attn.py": "3e58fdddc7a0b42ae3d9b5eabb375a993a4110bad499c37846bde466901dbeb3",
+    "plugin-site/ornith_g256/_xgrammar_native/build-receipt.json": "3bdf77657e5157c35938b66dfb46266608ade3b599a62ce74abb3f807d579e84",
+    "plugin-site/ornith_g256/_xgrammar_native/selection.json": "7cb93b3fe902de6d604bf96f79fdc3efbd8d612e494da838f6a8f0d731eb89ae",
+    "plugin-site/ornith_g256/adaptive_c1.py": "8e48fd874fbdf65183c20bcc7c748d2562dbcafa97e004f18a2818b98980a809",
+    "plugin-site/ornith_g256/attention.py": "2fdd3c0ed84dcb052a9e8f9df4068ddf80a969931bea0a295e022d36b753cda1",
+    "plugin-site/ornith_g256/attention_compact.py": "50df61a02a3890a2825b205ed18c7d18a2b002333d382c572440b58ca9b2b9fe",
+    "plugin-site/ornith_g256/attention_fast.py": "27b1a133cadb87ff3a63c027049e1da81e4112fd5d5260a4440fde85491f838a",
+    "plugin-site/ornith_g256/attention_folded.py": "73dac8203032bce6ceedf213d7eefe5987c921768ae237129e232fefd8d9bba3",
+    "plugin-site/ornith_g256/attention_iu4.py": "15c5341fc848c43ac956546e563cf8ab7f7649f05a674fb0153b0e2d76438cae",
+    "plugin-site/ornith_g256/attention_iu4_persistent.py": "cef74ae423974063212637a49cd64a5a10b9077f3a5e2287746722749a1e6297",
+    "plugin-site/ornith_g256/attention_mixed.py": "cbb530697f2555be6532b7d4d3401f3a33c4c58982659e88892720d21691110a",
+    "plugin-site/ornith_g256/attention_partition.py": "41bfb44dd0d0c46a2c0a725983155018e051cd89edb8527f1195a092bef7ffdd",
+    "plugin-site/ornith_g256/attention_storage.py": "3ab813e221835370bdb503b2ab7d48f46294441b1dd9c652c5dc976438bdb6e2",
+    "plugin-site/ornith_g256/attention_tile.py": "f2bf007132bc712c8e61e036ea8a94b54f18a813d9bdd334eee46436a3ca1fea",
+    "plugin-site/ornith_g256/attention_verify.py": "3a329599694265f4badecc2f9b709b280d9960b967d7192f3959ad9e0d8c6077",
+    "plugin-site/ornith_g256/attention_window.py": "78534f9a1b0beae9ecf80cc5b2a613e4fe4827036c8df0acc7e4dfd956aae1c5",
+    "plugin-site/ornith_g256/cache_full1120.py": "3d4240a749d9671c0bf9b3937200e1cd8d4abfe6699d8c1d6857b188c2ef8a8c",
+    "plugin-site/ornith_g256/column_backend.py": "a035b839369e733b05f68ce83df0d022cb62455016fccf30c2978127ab3c5403",
+    "plugin-site/ornith_g256/column_kernel.py": "27fb053d970184732a46212026afd3152ec07a0dfc3dc4c5cc5cc0d539f17c81",
+    "plugin-site/ornith_g256/config.py": "23210d8b625ef375638e2d6043a92a26181a5b5f1c537d9d4b34c91b7109f36f",
+    "plugin-site/ornith_g256/dense_n32.py": "119fb7d25e9d28b6e520b5285c10de86c7d8ec2fe29a2399e9c086173d765903",
+    "plugin-site/ornith_g256/dense_source.py": "cdc7dd6a13c2439cd405596069c301fe28cb95ff32278f1bafe8118fe357d48e",
+    "plugin-site/ornith_g256/dflash_conv_boundary.py": "a5f46458194cade5ddcbfbb51af30c662d12b9086fc459ecbeb17463339febb1",
+    "plugin-site/ornith_g256/dflash_spec.py": "c948b99eab22d2af5c5fc0ca9f764a0855a3d86aca0609362bf5474a9beff1b3",
+    "plugin-site/ornith_g256/dynamic_graphs.py": "6f51f0dd740f6546829e46bbbe840df2780ce8c49978658eb3a05092554ff90c",
+    "plugin-site/ornith_g256/gdn_compact.py": "8b5c062d10f386d78bc9926ab0576e56359401e5d1621c397a2922e2b426ccab",
+    "plugin-site/ornith_g256/gdn_compact_kernel.py": "0eaf34a2eccef2dabbc9e4f017e6ce84317a28efdcbabecf92b3b4c4c9429a55",
+    "plugin-site/ornith_g256/gdn_spec.py": "0931b1b3601d935bb6acdfcb65d307dae4105645cd387bd9b256dd7a525376a3",
+    "plugin-site/ornith_g256/graph_phase_guard.py": "6cabaf7d4e3be04bbac7a642544d38948de2e07ea7ca38f5927f91bf3ad39163",
+    "plugin-site/ornith_g256/launch.py": "c2854b86e2f97dd2ea0356b4b333922f62b9b3d1b3e71f7993d803c6d0957479",
+    "plugin-site/ornith_g256/lifecycle.py": "3285552177cb8ce8a0149a01f8168d2bc44a739a9395ca84e15acccbb20a0b73",
+    "plugin-site/ornith_g256/loader.py": "a853c6d014214a7adf298548f9175b6e18dd05338917aec16ebcb4fc3c3ecb13",
+    "plugin-site/ornith_g256/method.py": "fc301af72efd234430dbfb839d56e74328457b5b83e14a044aeb9554d8acc707",
+    "plugin-site/ornith_g256/moe_base.py": "26fb517415f2497859f2e3ff1aed8f31e7c56a4aacccfb34c6638e65d7df07c4",
+    "plugin-site/ornith_g256/native.py": "49cce0d67987b9659feb556aafb247423490c2153b7d77b750f055a4fdcda331",
+    "plugin-site/ornith_g256/native_tool_schema.py": "910060689d72b47ad4a7bb1e51ef301e41ff0be1beb9cc6b3df01c423d2edef5",
+    "plugin-site/ornith_g256/no_spec.py": "561106b75328fe0e161fa1243aca4cf50c5049d8e785d22c90aacb3e1e1a0104",
+    "plugin-site/ornith_g256/optimized/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "plugin-site/ornith_g256/optimized/compile_scope.py": "18971a73b3a5ad80e6212de7f7e419a1dd00eb7380a5e6c8559f5ac6f5ec8d4d",
+    "plugin-site/ornith_g256/optimized/draft.py": "e801301cf03d69998758d35fb5463754ccc30a18e030ed2b988b26cef43f429c",
+    "plugin-site/ornith_g256/optimized/draft_kernel.py": "3d6da6ef0c144505f888cbd35ba50b6b95ad8e48c376ab73297746bff228919e",
+    "plugin-site/ornith_g256/optimized/grammar_parallel.py": "e9c3608438712521980d7fcd60f58862d0b0a2876620183f304c64a03817edea",
+    "plugin-site/ornith_g256/optimized/greedy.py": "386443ad8729f8fe21ec5fef84ead3acfbd4d6c5dd3b4c1e50a7995ed1d1761f",
+    "plugin-site/ornith_g256/optimized/greedy_kernel.py": "2ebe2a25d1011a45eea7eda3c0964ca1eb2638df7d60f54db02c91435550dc42",
+    "plugin-site/ornith_g256/optimized/launch.py": "c1d42602154674916a94b9907b7746fac3ea48cc505fa69babf116f0f127e03d",
+    "plugin-site/ornith_g256/optimized/mixed_qk_sum.py": "17b18769ae84c2e9707a0257181e5c314e9774fc3b4796f8e6787a8302cb5e1f",
+    "plugin-site/ornith_g256/optimized/pipeline/SOURCE-PROOF.json": "5939bc8a53a831715d05d0fa0290beaa0cb00eda1597fe3bf1f1cfc53699fd0d",
+    "plugin-site/ornith_g256/optimized/pipeline/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "plugin-site/ornith_g256/optimized/pipeline/adapter.py": "6af97a8d5e0d79b7dfc32f2f0fe5f5c53b5c1c50d03542ff99c64cae7f3909e1",
+    "plugin-site/ornith_g256/optimized/pipeline/candidate/causal_conv1d.py": "254c540f77b01694696a1ab535abd0e5a23d54e7211e202cf9beead04face16b",
+    "plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_o.py": "a93f951048958265b4d1c3f9e67341329196e7a57525433934215fde4b16e61b",
+    "plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_scaled_dot_kkt.py": "71735d2b1c4a19eecbb6058708481817bb610426c48f96a5f646784afcf0c65c",
+    "plugin-site/ornith_g256/optimized/pipeline/candidate/solve_tril.py": "bdd04b69c0f91562048d25834bd53d298dc1f1bd068b9494af88776b2bf4b1d1",
+    "plugin-site/ornith_g256/optimized/pipeline/candidate/wy_fast.py": "4cc0be519bfcf3812d810a4bebb3dd45fb9b13a8d0ab62d0ad463a1a2e36f61f",
+    "plugin-site/ornith_g256/optimized/postconv.py": "c58ee5d2b8418fbf69eda5d9ee4f342b11cc012ab77ae25b6ce2b2aea51aa50f",
+    "plugin-site/ornith_g256/optimized/prefix.py": "c695a899ab12b58c3b89b6ecc276e82efa3f0833a319911a8528cc4e237710cd",
+    "plugin-site/ornith_g256/optimized/qk.py": "d86f1a94086668ac6402eaec7a19c0ad8e55dc8e6377f8d4a9eeab6aa96f9664",
+    "plugin-site/ornith_g256/optimized/qk_schedule.py": "66038781f0667486fba458ef3118ca89ad9640dc2c2bd05cf8576440f0f178a2",
+    "plugin-site/ornith_g256/optimized/worker.py": "38579b767235e7fcac3f2e6b0507438f155c134e1e85f3d1c205a43bece95512",
+    "plugin-site/ornith_g256/parser_contract.py": "2cebab403dc882569f3ce67d12798df14cebd5c162c72f82becd37bcdd7e2a28",
+    "plugin-site/ornith_g256/phase_dispatch.py": "1a4ebdf2af709e6ca519e0a08f6a537602df178ca575332045617631640f8137",
+    "plugin-site/ornith_g256/prefill_draft.py": "d79bb29a124639543a64e198908e2eaef4cbcf492f2e5e14debcfc9c867816f5",
+    "plugin-site/ornith_g256/prefix_cache.py": "65fb972e285e37741f53b221d2f6380ec464d34a39f36bb66a82dcaacab727d0",
+    "plugin-site/ornith_g256/runtime.py": "db43b1c3766720e0dbee4e8617aee673a533c30d36fcc1fbf43919717d8ab149",
+    "plugin-site/ornith_g256/runtime_correctness.py": "08f41773e623a4f99cf8ab19e434d7e33bac9cfa5eb05a9dc0f4511aab8ba4bc",
+    "plugin-site/ornith_g256/strict_qwen/__init__.py": "ee8a5055249f1413fc22f46421ab77b5485aa39b199c68eb5b56c75dcc40922d",
+    "plugin-site/ornith_g256/strict_qwen/qwen3.py": "33c653447a144b72257e508ccd92610db6d30578e28f98c69c99281773dead1c",
+    "plugin-site/ornith_g256/strict_qwen/qwen3_contract.py": "29c083659c54c81144fdb6c8329c75aaa62cf69d1c45fa7da8bdeeeca4856eaf",
+    "plugin-site/ornith_g256/worker.py": "acc27240c21f3478ca452ddf04244df4ee770efaac38192ef64d328ef22701d5",
+    "plugin-site/ornith_g256/worker_base.py": "a00ef1295172e655c8b3c7d790cfd0316114a9492e483dedfdea036d3b7bf868",
+    "plugin-site/ornith_g256/xgrammar_native.py": "d19c546e124cb11311e4cc31f5833dc663e2789d45090091c0f8fec8a9626623"
+  },
+  "runtime_native_grammar": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7",
+  "note": "Same measured model libraries and plugin code; R04 source pinned by patches/xgrammar/SOURCE.json. Portable recipe includes output32 compile definition. Rebuilding with another toolchain need not produce identical binaries."
+}

bundle/LICENSES/XGRAMMAR-LICENSE

@@ -0,0 +1,201 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

bundle/native/libornith_attention_iu4.so

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:2905806824bece62ce9d9140608859df869f3e3d3663c67bf79ad4ea495cd61f
-size 186640
+oid sha256:3d29cf3616c69f9fe2b89f8d3564d479dda682fdc06572886e7b60faccc9a53a
+size 252976

bundle/native/libornith_persistent_iu4.so

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:40e5cf35157c3e716de8067f7eef1312a3b9e787dd52a6f2e232ccd2d07fa68c
-size 103792
+oid sha256:c70ba14d5aebf67b761ea54cfc9a38f8bbedf74cfca8b19050048acda9c8a06c
+size 104904

bundle/native/libornith_routed_storage_n32.so

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:3ea288cd0e37a4c373038744c3cc4cd4eac09675ac137e953d5ae7a6cd43dd3d
-size 158536
+oid sha256:2002bf6f82a637f5842f7c5f09a09f3b7a0db55edc4b3e35181c8f9a12ffb4b1
+size 166536

bundle/packaging/serve.sh

@@ -28,4 +28,4 @@ export XDG_CACHE_HOME="$cache_directory" AITER_JIT_DIR="$cache_directory/aiter"
 source "$runtime_root/runtime-env.sh"
 unset VLLM_SOURCE VLLM_VENV
 export PYTHONPATH="$plugin_site${PYTHONPATH:+:$PYTHONPATH}"
-exec "$runtime_root/venv/bin/python" -m ornith_g256.launch "${launch_args[@]}" --cache-directory "$cache_directory"
+exec "$runtime_root/venv/bin/python" -m ornith_g256.optimized.launch "${launch_args[@]}" --cache-directory "$cache_directory"

bundle/plugin-site/TOOL-INTEGRATION.json

@@ -0,0 +1,205 @@
+{
+  "status": "STAGED_PENDING_QUALIFICATION",
+  "parent": "/srv/ssd/p3700ba/scratch/crown/optimization-campaign-20260916/artifacts/ornith-combined-v3/bundle/plugin-site/ornith_g256",
+  "parent_files": {
+    "attention_window.py": "78534f9a1b0beae9ecf80cc5b2a613e4fe4827036c8df0acc7e4dfd956aae1c5",
+    "dflash_spec.py": "c948b99eab22d2af5c5fc0ca9f764a0855a3d86aca0609362bf5474a9beff1b3",
+    "gdn_compact.py": "8b5c062d10f386d78bc9926ab0576e56359401e5d1621c397a2922e2b426ccab",
+    "cache_full1120.py": "3d4240a749d9671c0bf9b3937200e1cd8d4abfe6699d8c1d6857b188c2ef8a8c",
+    "attention_storage.py": "3ab813e221835370bdb503b2ab7d48f46294441b1dd9c652c5dc976438bdb6e2",
+    "prefill_draft.py": "d79bb29a124639543a64e198908e2eaef4cbcf492f2e5e14debcfc9c867816f5",
+    "column_kernel.py": "27fb053d970184732a46212026afd3152ec07a0dfc3dc4c5cc5cc0d539f17c81",
+    "moe_base.py": "26fb517415f2497859f2e3ff1aed8f31e7c56a4aacccfb34c6638e65d7df07c4",
+    "attention_partition.py": "41bfb44dd0d0c46a2c0a725983155018e051cd89edb8527f1195a092bef7ffdd",
+    "lifecycle.py": "3285552177cb8ce8a0149a01f8168d2bc44a739a9395ca84e15acccbb20a0b73",
+    "attention_verify.py": "3a329599694265f4badecc2f9b709b280d9960b967d7192f3959ad9e0d8c6077",
+    "attention_iu4_persistent.py": "cef74ae423974063212637a49cd64a5a10b9077f3a5e2287746722749a1e6297",
+    "loader.py": "a853c6d014214a7adf298548f9175b6e18dd05338917aec16ebcb4fc3c3ecb13",
+    "attention_mixed.py": "cbb530697f2555be6532b7d4d3401f3a33c4c58982659e88892720d21691110a",
+    "graph_phase_guard.py": "6cabaf7d4e3be04bbac7a642544d38948de2e07ea7ca38f5927f91bf3ad39163",
+    "phase_dispatch.py": "1a4ebdf2af709e6ca519e0a08f6a537602df178ca575332045617631640f8137",
+    "method.py": "fc301af72efd234430dbfb839d56e74328457b5b83e14a044aeb9554d8acc707",
+    "gdn_compact_kernel.py": "0eaf34a2eccef2dabbc9e4f017e6ce84317a28efdcbabecf92b3b4c4c9429a55",
+    "attention_tile.py": "f2bf007132bc712c8e61e036ea8a94b54f18a813d9bdd334eee46436a3ca1fea",
+    "no_spec.py": "561106b75328fe0e161fa1243aca4cf50c5049d8e785d22c90aacb3e1e1a0104",
+    "__init__.py": "1f526dde80a76eb467b1fd30e8b2babdf0b677cb2f246217ac3ab2509d8cf288",
+    "gdn_spec.py": "0931b1b3601d935bb6acdfcb65d307dae4105645cd387bd9b256dd7a525376a3",
+    "column_backend.py": "a035b839369e733b05f68ce83df0d022cb62455016fccf30c2978127ab3c5403",
+    "launch.py": "c2854b86e2f97dd2ea0356b4b333922f62b9b3d1b3e71f7993d803c6d0957479",
+    "dense_n32.py": "119fb7d25e9d28b6e520b5285c10de86c7d8ec2fe29a2399e9c086173d765903",
+    "prefix_cache.py": "65fb972e285e37741f53b221d2f6380ec464d34a39f36bb66a82dcaacab727d0",
+    "runtime.py": "db43b1c3766720e0dbee4e8617aee673a533c30d36fcc1fbf43919717d8ab149",
+    "parser_contract.py": "2cebab403dc882569f3ce67d12798df14cebd5c162c72f82becd37bcdd7e2a28",
+    "attention_fast.py": "27b1a133cadb87ff3a63c027049e1da81e4112fd5d5260a4440fde85491f838a",
+    "attention.py": "2fdd3c0ed84dcb052a9e8f9df4068ddf80a969931bea0a295e022d36b753cda1",
+    "native_tool_schema.py": "910060689d72b47ad4a7bb1e51ef301e41ff0be1beb9cc6b3df01c423d2edef5",
+    "attention_folded.py": "73dac8203032bce6ceedf213d7eefe5987c921768ae237129e232fefd8d9bba3",
+    "dynamic_graphs.py": "6f51f0dd740f6546829e46bbbe840df2780ce8c49978658eb3a05092554ff90c",
+    "worker.py": "acc27240c21f3478ca452ddf04244df4ee770efaac38192ef64d328ef22701d5",
+    "runtime_correctness.py": "08f41773e623a4f99cf8ab19e434d7e33bac9cfa5eb05a9dc0f4511aab8ba4bc",
+    "config.py": "23210d8b625ef375638e2d6043a92a26181a5b5f1c537d9d4b34c91b7109f36f",
+    "dflash_conv_boundary.py": "a5f46458194cade5ddcbfbb51af30c662d12b9086fc459ecbeb17463339febb1",
+    "native.py": "49cce0d67987b9659feb556aafb247423490c2153b7d77b750f055a4fdcda331",
+    "attention_compact.py": "50df61a02a3890a2825b205ed18c7d18a2b002333d382c572440b58ca9b2b9fe",
+    "worker_base.py": "a00ef1295172e655c8b3c7d790cfd0316114a9492e483dedfdea036d3b7bf868",
+    "attention_iu4.py": "15c5341fc848c43ac956546e563cf8ab7f7649f05a674fb0153b0e2d76438cae",
+    "dense_source.py": "cdc7dd6a13c2439cd405596069c301fe28cb95ff32278f1bafe8118fe357d48e",
+    "adaptive_c1.py": "8e48fd874fbdf65183c20bcc7c748d2562dbcafa97e004f18a2818b98980a809",
+    "_vllm_correctness/gpu_model_runner.py": "ef0009acaa67a1a5e066984f34bf619b06e85220ff953a3d9511376f0146dccf",
+    "_vllm_correctness/__init__.py": "963097106407a87b62b275041516144de15713dee65108accc29e27036ffc059",
+    "_vllm_correctness/manifest.json": "0ddfa489f4f9ead135006e29aa675e7522afa5cee10596c2115e07f13186478c",
+    "_vllm_correctness/causal_conv1d.py": "230f3e554f5f1d609a37022578c84e42d5c71d46e684101278a7e21af4f7d418",
+    "_vllm_correctness/gdn_attn.py": "a3c361d502e8bca2aa46a0baed5ec753a3c62e3d027a298e3ed4692794f9f6ec",
+    "_vllm_correctness/qwen_gdn_linear_attn.py": "3e58fdddc7a0b42ae3d9b5eabb375a993a4110bad499c37846bde466901dbeb3",
+    "strict_qwen/__init__.py": "ee8a5055249f1413fc22f46421ab77b5485aa39b199c68eb5b56c75dcc40922d",
+    "strict_qwen/qwen3_contract.py": "29c083659c54c81144fdb6c8329c75aaa62cf69d1c45fa7da8bdeeeca4856eaf",
+    "strict_qwen/qwen3.py": "33c653447a144b72257e508ccd92610db6d30578e28f98c69c99281773dead1c",
+    "optimized/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "optimized/compile_scope.py": "18971a73b3a5ad80e6212de7f7e419a1dd00eb7380a5e6c8559f5ac6f5ec8d4d",
+    "optimized/draft.py": "e801301cf03d69998758d35fb5463754ccc30a18e030ed2b988b26cef43f429c",
+    "optimized/draft_kernel.py": "3d6da6ef0c144505f888cbd35ba50b6b95ad8e48c376ab73297746bff228919e",
+    "optimized/greedy.py": "386443ad8729f8fe21ec5fef84ead3acfbd4d6c5dd3b4c1e50a7995ed1d1761f",
+    "optimized/greedy_kernel.py": "2ebe2a25d1011a45eea7eda3c0964ca1eb2638df7d60f54db02c91435550dc42",
+    "optimized/launch.py": "c1d42602154674916a94b9907b7746fac3ea48cc505fa69babf116f0f127e03d",
+    "optimized/mixed_qk_sum.py": "17b18769ae84c2e9707a0257181e5c314e9774fc3b4796f8e6787a8302cb5e1f",
+    "optimized/postconv.py": "c58ee5d2b8418fbf69eda5d9ee4f342b11cc012ab77ae25b6ce2b2aea51aa50f",
+    "optimized/prefix.py": "c695a899ab12b58c3b89b6ecc276e82efa3f0833a319911a8528cc4e237710cd",
+    "optimized/qk.py": "d86f1a94086668ac6402eaec7a19c0ad8e55dc8e6377f8d4a9eeab6aa96f9664",
+    "optimized/qk_schedule.py": "66038781f0667486fba458ef3118ca89ad9640dc2c2bd05cf8576440f0f178a2",
+    "optimized/worker.py": "38579b767235e7fcac3f2e6b0507438f155c134e1e85f3d1c205a43bece95512",
+    "optimized/grammar_parallel.py": "e9c3608438712521980d7fcd60f58862d0b0a2876620183f304c64a03817edea",
+    "optimized/pipeline/SOURCE-PROOF.json": "5939bc8a53a831715d05d0fa0290beaa0cb00eda1597fe3bf1f1cfc53699fd0d",
+    "optimized/pipeline/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "optimized/pipeline/adapter.py": "6af97a8d5e0d79b7dfc32f2f0fe5f5c53b5c1c50d03542ff99c64cae7f3909e1",
+    "optimized/pipeline/candidate/causal_conv1d.py": "254c540f77b01694696a1ab535abd0e5a23d54e7211e202cf9beead04face16b",
+    "optimized/pipeline/candidate/chunk_o.py": "a93f951048958265b4d1c3f9e67341329196e7a57525433934215fde4b16e61b",
+    "optimized/pipeline/candidate/chunk_scaled_dot_kkt.py": "71735d2b1c4a19eecbb6058708481817bb610426c48f96a5f646784afcf0c65c",
+    "optimized/pipeline/candidate/solve_tril.py": "bdd04b69c0f91562048d25834bd53d298dc1f1bd068b9494af88776b2bf4b1d1",
+    "optimized/pipeline/candidate/wy_fast.py": "4cc0be519bfcf3812d810a4bebb3dd45fb9b13a8d0ab62d0ad463a1a2e36f61f"
+  },
+  "changed_existing_files": [
+    "__init__.py",
+    "_vllm_correctness/manifest.json"
+  ],
+  "files": {
+    "ornith_g256/__init__.py": "60abbb51a475487e3abcb9e5b34d8836b4d4b36e0bced062d2382dc7b1076cf6",
+    "ornith_g256/_vllm_correctness/__init__.py": "963097106407a87b62b275041516144de15713dee65108accc29e27036ffc059",
+    "ornith_g256/_vllm_correctness/causal_conv1d.py": "230f3e554f5f1d609a37022578c84e42d5c71d46e684101278a7e21af4f7d418",
+    "ornith_g256/_vllm_correctness/chat_completion_serving_r03.py": "5720371f5bcb351e340a3e8d9a0235ec01ff27ed93e4f35542e31b83a055ef45",
+    "ornith_g256/_vllm_correctness/gdn_attn.py": "a3c361d502e8bca2aa46a0baed5ec753a3c62e3d027a298e3ed4692794f9f6ec",
+    "ornith_g256/_vllm_correctness/gpu_model_runner.py": "ef0009acaa67a1a5e066984f34bf619b06e85220ff953a3d9511376f0146dccf",
+    "ornith_g256/_vllm_correctness/manifest.json": "bdef4ca7151a7a5e0410dd3e8bcd2e8a85f125938ec104451cb8777267c3c4d9",
+    "ornith_g256/_vllm_correctness/qwen_gdn_linear_attn.py": "3e58fdddc7a0b42ae3d9b5eabb375a993a4110bad499c37846bde466901dbeb3",
+    "ornith_g256/_xgrammar_native/build-receipt.json": "3bdf77657e5157c35938b66dfb46266608ade3b599a62ce74abb3f807d579e84",
+    "ornith_g256/_xgrammar_native/libxgrammar-candidate.so": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7",
+    "ornith_g256/_xgrammar_native/selection.json": "7cb93b3fe902de6d604bf96f79fdc3efbd8d612e494da838f6a8f0d731eb89ae",
+    "ornith_g256/adaptive_c1.py": "8e48fd874fbdf65183c20bcc7c748d2562dbcafa97e004f18a2818b98980a809",
+    "ornith_g256/attention.py": "2fdd3c0ed84dcb052a9e8f9df4068ddf80a969931bea0a295e022d36b753cda1",
+    "ornith_g256/attention_compact.py": "50df61a02a3890a2825b205ed18c7d18a2b002333d382c572440b58ca9b2b9fe",
+    "ornith_g256/attention_fast.py": "27b1a133cadb87ff3a63c027049e1da81e4112fd5d5260a4440fde85491f838a",
+    "ornith_g256/attention_folded.py": "73dac8203032bce6ceedf213d7eefe5987c921768ae237129e232fefd8d9bba3",
+    "ornith_g256/attention_iu4.py": "15c5341fc848c43ac956546e563cf8ab7f7649f05a674fb0153b0e2d76438cae",
+    "ornith_g256/attention_iu4_persistent.py": "cef74ae423974063212637a49cd64a5a10b9077f3a5e2287746722749a1e6297",
+    "ornith_g256/attention_mixed.py": "cbb530697f2555be6532b7d4d3401f3a33c4c58982659e88892720d21691110a",
+    "ornith_g256/attention_partition.py": "41bfb44dd0d0c46a2c0a725983155018e051cd89edb8527f1195a092bef7ffdd",
+    "ornith_g256/attention_storage.py": "3ab813e221835370bdb503b2ab7d48f46294441b1dd9c652c5dc976438bdb6e2",
+    "ornith_g256/attention_tile.py": "f2bf007132bc712c8e61e036ea8a94b54f18a813d9bdd334eee46436a3ca1fea",
+    "ornith_g256/attention_verify.py": "3a329599694265f4badecc2f9b709b280d9960b967d7192f3959ad9e0d8c6077",
+    "ornith_g256/attention_window.py": "78534f9a1b0beae9ecf80cc5b2a613e4fe4827036c8df0acc7e4dfd956aae1c5",
+    "ornith_g256/cache_full1120.py": "3d4240a749d9671c0bf9b3937200e1cd8d4abfe6699d8c1d6857b188c2ef8a8c",
+    "ornith_g256/column_backend.py": "a035b839369e733b05f68ce83df0d022cb62455016fccf30c2978127ab3c5403",
+    "ornith_g256/column_kernel.py": "27fb053d970184732a46212026afd3152ec07a0dfc3dc4c5cc5cc0d539f17c81",
+    "ornith_g256/config.py": "23210d8b625ef375638e2d6043a92a26181a5b5f1c537d9d4b34c91b7109f36f",
+    "ornith_g256/dense_n32.py": "119fb7d25e9d28b6e520b5285c10de86c7d8ec2fe29a2399e9c086173d765903",
+    "ornith_g256/dense_source.py": "cdc7dd6a13c2439cd405596069c301fe28cb95ff32278f1bafe8118fe357d48e",
+    "ornith_g256/dflash_conv_boundary.py": "a5f46458194cade5ddcbfbb51af30c662d12b9086fc459ecbeb17463339febb1",
+    "ornith_g256/dflash_spec.py": "c948b99eab22d2af5c5fc0ca9f764a0855a3d86aca0609362bf5474a9beff1b3",
+    "ornith_g256/dynamic_graphs.py": "6f51f0dd740f6546829e46bbbe840df2780ce8c49978658eb3a05092554ff90c",
+    "ornith_g256/gdn_compact.py": "8b5c062d10f386d78bc9926ab0576e56359401e5d1621c397a2922e2b426ccab",
+    "ornith_g256/gdn_compact_kernel.py": "0eaf34a2eccef2dabbc9e4f017e6ce84317a28efdcbabecf92b3b4c4c9429a55",
+    "ornith_g256/gdn_spec.py": "0931b1b3601d935bb6acdfcb65d307dae4105645cd387bd9b256dd7a525376a3",
+    "ornith_g256/graph_phase_guard.py": "6cabaf7d4e3be04bbac7a642544d38948de2e07ea7ca38f5927f91bf3ad39163",
+    "ornith_g256/launch.py": "c2854b86e2f97dd2ea0356b4b333922f62b9b3d1b3e71f7993d803c6d0957479",
+    "ornith_g256/lifecycle.py": "3285552177cb8ce8a0149a01f8168d2bc44a739a9395ca84e15acccbb20a0b73",
+    "ornith_g256/loader.py": "a853c6d014214a7adf298548f9175b6e18dd05338917aec16ebcb4fc3c3ecb13",
+    "ornith_g256/method.py": "fc301af72efd234430dbfb839d56e74328457b5b83e14a044aeb9554d8acc707",
+    "ornith_g256/moe_base.py": "26fb517415f2497859f2e3ff1aed8f31e7c56a4aacccfb34c6638e65d7df07c4",
+    "ornith_g256/native.py": "49cce0d67987b9659feb556aafb247423490c2153b7d77b750f055a4fdcda331",
+    "ornith_g256/native_tool_schema.py": "910060689d72b47ad4a7bb1e51ef301e41ff0be1beb9cc6b3df01c423d2edef5",
+    "ornith_g256/no_spec.py": "561106b75328fe0e161fa1243aca4cf50c5049d8e785d22c90aacb3e1e1a0104",
+    "ornith_g256/optimized/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "ornith_g256/optimized/compile_scope.py": "18971a73b3a5ad80e6212de7f7e419a1dd00eb7380a5e6c8559f5ac6f5ec8d4d",
+    "ornith_g256/optimized/draft.py": "e801301cf03d69998758d35fb5463754ccc30a18e030ed2b988b26cef43f429c",
+    "ornith_g256/optimized/draft_kernel.py": "3d6da6ef0c144505f888cbd35ba50b6b95ad8e48c376ab73297746bff228919e",
+    "ornith_g256/optimized/grammar_parallel.py": "e9c3608438712521980d7fcd60f58862d0b0a2876620183f304c64a03817edea",
+    "ornith_g256/optimized/greedy.py": "386443ad8729f8fe21ec5fef84ead3acfbd4d6c5dd3b4c1e50a7995ed1d1761f",
+    "ornith_g256/optimized/greedy_kernel.py": "2ebe2a25d1011a45eea7eda3c0964ca1eb2638df7d60f54db02c91435550dc42",
+    "ornith_g256/optimized/launch.py": "c1d42602154674916a94b9907b7746fac3ea48cc505fa69babf116f0f127e03d",
+    "ornith_g256/optimized/mixed_qk_sum.py": "17b18769ae84c2e9707a0257181e5c314e9774fc3b4796f8e6787a8302cb5e1f",
+    "ornith_g256/optimized/pipeline/SOURCE-PROOF.json": "5939bc8a53a831715d05d0fa0290beaa0cb00eda1597fe3bf1f1cfc53699fd0d",
+    "ornith_g256/optimized/pipeline/__init__.py": "0297f5d476afdddb0556f4f236f95e514db6b9b96788df672a4a0bc0074fd835",
+    "ornith_g256/optimized/pipeline/adapter.py": "6af97a8d5e0d79b7dfc32f2f0fe5f5c53b5c1c50d03542ff99c64cae7f3909e1",
+    "ornith_g256/optimized/pipeline/candidate/causal_conv1d.py": "254c540f77b01694696a1ab535abd0e5a23d54e7211e202cf9beead04face16b",
+    "ornith_g256/optimized/pipeline/candidate/chunk_o.py": "a93f951048958265b4d1c3f9e67341329196e7a57525433934215fde4b16e61b",
+    "ornith_g256/optimized/pipeline/candidate/chunk_scaled_dot_kkt.py": "71735d2b1c4a19eecbb6058708481817bb610426c48f96a5f646784afcf0c65c",
+    "ornith_g256/optimized/pipeline/candidate/solve_tril.py": "bdd04b69c0f91562048d25834bd53d298dc1f1bd068b9494af88776b2bf4b1d1",
+    "ornith_g256/optimized/pipeline/candidate/wy_fast.py": "4cc0be519bfcf3812d810a4bebb3dd45fb9b13a8d0ab62d0ad463a1a2e36f61f",
+    "ornith_g256/optimized/postconv.py": "c58ee5d2b8418fbf69eda5d9ee4f342b11cc012ab77ae25b6ce2b2aea51aa50f",
+    "ornith_g256/optimized/prefix.py": "c695a899ab12b58c3b89b6ecc276e82efa3f0833a319911a8528cc4e237710cd",
+    "ornith_g256/optimized/qk.py": "d86f1a94086668ac6402eaec7a19c0ad8e55dc8e6377f8d4a9eeab6aa96f9664",
+    "ornith_g256/optimized/qk_schedule.py": "66038781f0667486fba458ef3118ca89ad9640dc2c2bd05cf8576440f0f178a2",
+    "ornith_g256/optimized/worker.py": "38579b767235e7fcac3f2e6b0507438f155c134e1e85f3d1c205a43bece95512",
+    "ornith_g256/parser_contract.py": "2cebab403dc882569f3ce67d12798df14cebd5c162c72f82becd37bcdd7e2a28",
+    "ornith_g256/phase_dispatch.py": "1a4ebdf2af709e6ca519e0a08f6a537602df178ca575332045617631640f8137",
+    "ornith_g256/prefill_draft.py": "d79bb29a124639543a64e198908e2eaef4cbcf492f2e5e14debcfc9c867816f5",
+    "ornith_g256/prefix_cache.py": "65fb972e285e37741f53b221d2f6380ec464d34a39f36bb66a82dcaacab727d0",
+    "ornith_g256/runtime.py": "db43b1c3766720e0dbee4e8617aee673a533c30d36fcc1fbf43919717d8ab149",
+    "ornith_g256/runtime_correctness.py": "08f41773e623a4f99cf8ab19e434d7e33bac9cfa5eb05a9dc0f4511aab8ba4bc",
+    "ornith_g256/strict_qwen/__init__.py": "ee8a5055249f1413fc22f46421ab77b5485aa39b199c68eb5b56c75dcc40922d",
+    "ornith_g256/strict_qwen/qwen3.py": "33c653447a144b72257e508ccd92610db6d30578e28f98c69c99281773dead1c",
+    "ornith_g256/strict_qwen/qwen3_contract.py": "29c083659c54c81144fdb6c8329c75aaa62cf69d1c45fa7da8bdeeeca4856eaf",
+    "ornith_g256/worker.py": "acc27240c21f3478ca452ddf04244df4ee770efaac38192ef64d328ef22701d5",
+    "ornith_g256/worker_base.py": "a00ef1295172e655c8b3c7d790cfd0316114a9492e483dedfdea036d3b7bf868",
+    "ornith_g256/xgrammar_native.py": "d19c546e124cb11311e4cc31f5833dc663e2789d45090091c0f8fec8a9626623"
+  },
+  "r01_sha256": "e9c3608438712521980d7fcd60f58862d0b0a2876620183f304c64a03817edea",
+  "r01_existing_preserved": true,
+  "r04": {
+    "mode": "candidate",
+    "library": "libxgrammar-candidate.so",
+    "sha256": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7",
+    "source_commit": "557becfb64c503ae9c04344b0047661f43f44320",
+    "patch_sha256": "2c42719e948809242f230ec534f967eab1babad9be1ba6dda9c59fd706f20d45",
+    "sources": {
+      "cpp/grammar.cc": "01d1552b37ecaa72da601a3b6d43975a2c3a0de02cdf797e82c243b30eef8c9d",
+      "cpp/grammar_functor.cc": "9410c20b665093b9fb67dc30db148121b819d3aee5ff94eb82558bf79b1f8279"
+    },
+    "configure": [
+      "cmake",
+      "-S",
+      "/srv/ssd/p3700ba/scratch/crown/runtime-tool-integration-20260917/native-build-r2/source-candidate",
+      "-B",
+      "/srv/ssd/p3700ba/scratch/crown/runtime-tool-integration-20260917/native-build-r2/build-candidate",
+      "-G",
+      "Ninja",
+      "-DCMAKE_BUILD_TYPE=Release",
+      "-DXGRAMMAR_BUILD_PYTHON_BINDINGS=ON",
+      "-DCMAKE_CXX_FLAGS_RELEASE=-O3 -DNDEBUG -fno-lto",
+      "-DCMAKE_CXX_FLAGS_RELWITHDEBINFO=-O3 -g -DNDEBUG -fno-lto",
+      "-DPython_EXECUTABLE=/opt/ciru/glm53-iu4/venv/bin/python",
+      "-DPython3_EXECUTABLE=/opt/ciru/glm53-iu4/venv/bin/python"
+    ],
+    "build_seconds": 67.22386657004245,
+    "compiler": "g++ (GCC) 15.2.0\nCopyright (C) 2025 Free Software Foundation, Inc.\nThis is free software; see the source for copying conditions.  There is NO\nwarranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.\n\n",
+    "submodules": " 6689d14c203eed390ae7bb64f56a983cfd7dff9c 3rdparty/cpptrace (v0.1-644-g6689d14)\n bbd2f4d32427e548797929af08cfe2a9cbb3cf12 3rdparty/dlpack (v1.0)\n df1544bcee0c7ce35cd5ea0b3eb8cc81855a4140 3rdparty/googletest (release-1.8.0-3498-gdf1544bc)\n"
+  },
+  "r03": {
+    "file": "chat_completion_serving_r03.py",
+    "native_path": "entrypoints/openai/chat_completion/serving.py",
+    "native_sha256": "9982953285e9df469032a82fffa4095d0e9d86278bede6e2b91d03d02373d182",
+    "sha256": "5720371f5bcb351e340a3e8d9a0235ec01ff27ed93e4f35542e31b83a055ef45"
+  },
+  "model_settings_changed": false,
+  "weights_changed": false,
+  "strict_policy_changed": false
+}

bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/direct_url.json

@@ -1 +0,0 @@
-{"url":"file:///srv/ssd/sn850x/scratch/crown/ornith-prefix-64k-v1/src/runtime/ornith_g256","dir_info":{}}

bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/uv_build.json

@@ -1 +0,0 @@
-{}

bundle/plugin-site/ciru_ornith_g256-1.0.2.dist-info/uv_cache.json

@@ -1 +0,0 @@
-{"timestamp":{"secs_since_epoch":1788727205,"nanos_since_epoch":456565118},"commit":null,"tags":null,"env":{},"directories":{"src":null}}

bundle/plugin-site/{ciru_ornith_g256-1.0.2.dist-info → ciru_ornith_g256-4.0.0.dist-info}/METADATA

@@ -1,6 +1,6 @@
 Metadata-Version: 2.4
 Name: ciru-ornith-g256
-Version: 1.0.2
+Version: 4.0.0
 Summary: Self-contained project adapter for Ornith G256 and DFlash2 on Ciru vLLM
 Author-email: Ciru <ciru@ciru.ai>
 Requires-Python: >=3.10

bundle/plugin-site/{ciru_ornith_g256-1.0.2.dist-info → ciru_ornith_g256-4.0.0.dist-info}/RECORD

@@ -1,27 +1,29 @@
-ciru_ornith_g256-1.0.2.dist-info/INSTALLER,sha256=5hhM4Q4mYTT9z6QB6PGpUAW81PGNFrYrdXMj4oM_6ak,2
-ciru_ornith_g256-1.0.2.dist-info/METADATA,sha256=hoIAew6fYBXVq6WqaJFx3YgykAk-iDq3d2OU9LaQhRc,254
-ciru_ornith_g256-1.0.2.dist-info/REQUESTED,sha256=47DEQpj8HBSa-_TImW-5JCeuQeRkm5NMpJWZG3hSuFU,0
-ciru_ornith_g256-1.0.2.dist-info/WHEEL,sha256=SmOxYU7pzNKBqASvQJ7DjX3XGUF92lrGhMb3R6_iiqI,91
-ciru_ornith_g256-1.0.2.dist-info/direct_url.json,sha256=MwFMTtCGcu6fn-TruM1lo2XsI7q0gT0JZbfW7sb3OEA,105
-ciru_ornith_g256-1.0.2.dist-info/entry_points.txt,sha256=ESe1wopdkiGfr5DNWZcSemlgPr3uvRre1OyHkl88QqU,121
-ciru_ornith_g256-1.0.2.dist-info/licenses/LICENSE-APACHE-2.0,sha256=z8d0m5b2O9McPEK1xHG_dWgUBT6EfBDz6wA0F7xSPTA,11358
-ciru_ornith_g256-1.0.2.dist-info/top_level.txt,sha256=EdCMLXnn8tDNBUR3JhkMvOi8hEw6IuhYFYEcfdhekxU,12
-ciru_ornith_g256-1.0.2.dist-info/uv_build.json,sha256=RBNvo1WzZ4oRRq0W9-hknpT7T8If536DEMBg9hyq_4o,2
-ciru_ornith_g256-1.0.2.dist-info/uv_cache.json,sha256=L73WKGonFia8yJ8vK90qJmbQxtUhlvCNsq7D4umUJX4,137
-ornith_g256/__init__.py,sha256=tqhBbqjsIA7JBuJINSvgLpqqNdviaoSPeG-jINdHqPM,1368
+TOOL-INTEGRATION.json,sha256=vBFvgYo5fz-IPpIJ1ISJmaW-UfmWif8qQGPH8iGkZNk,18419
+ciru_ornith_g256-4.0.0.dist-info/INSTALLER,sha256=5hhM4Q4mYTT9z6QB6PGpUAW81PGNFrYrdXMj4oM_6ak,2
+ciru_ornith_g256-4.0.0.dist-info/METADATA,sha256=wDaK9fhKQVhFOwacWUkaqv6EOGv6W9qjkBqKOB1kh4E,254
+ciru_ornith_g256-4.0.0.dist-info/REQUESTED,sha256=47DEQpj8HBSa-_TImW-5JCeuQeRkm5NMpJWZG3hSuFU,0
+ciru_ornith_g256-4.0.0.dist-info/WHEEL,sha256=SmOxYU7pzNKBqASvQJ7DjX3XGUF92lrGhMb3R6_iiqI,91
+ciru_ornith_g256-4.0.0.dist-info/entry_points.txt,sha256=ESe1wopdkiGfr5DNWZcSemlgPr3uvRre1OyHkl88QqU,121
+ciru_ornith_g256-4.0.0.dist-info/licenses/LICENSE-APACHE-2.0,sha256=z8d0m5b2O9McPEK1xHG_dWgUBT6EfBDz6wA0F7xSPTA,11358
+ciru_ornith_g256-4.0.0.dist-info/top_level.txt,sha256=EdCMLXnn8tDNBUR3JhkMvOi8hEw6IuhYFYEcfdhekxU,12
+ornith_g256/__init__.py,sha256=YKu7UaR1SH46vLnls02INrTUs24LztBi0jgtx7EHbPY,1571
 ornith_g256/_vllm_correctness/__init__.py,sha256=ljCXEGQHqHtisnUEFRYUTeFXE97mUQiszCnicDb_wFk,79
 ornith_g256/_vllm_correctness/causal_conv1d.py,sha256=Iw8-VU9fHWCaNwIleMhOQtXHHUbmhBASeKfiGvT31Bg,51774
+ornith_g256/_vllm_correctness/chat_completion_serving_r03.py,sha256=VyA3H1vLNR40Cj6NmgI17AH_J-2T5PNVQuMbg6BV70U,57061
 ornith_g256/_vllm_correctness/gdn_attn.py,sha256=o8Nh1QLovKKqRqC67V7HU6PGLj0CeimOPtRpJ5T59uw,26224
 ornith_g256/_vllm_correctness/gpu_model_runner.py,sha256=7wAJrKpnoaXgZphPNL9hmwboUiD_lTo9lRE3bwFG3M8,345523
-ornith_g256/_vllm_correctness/manifest.json,sha256=Dd-kifT56tE1AG4pqmdedSKvpc7hBZbCEV4H8TGGR4w,1607
+ornith_g256/_vllm_correctness/manifest.json,sha256=ve9MpxUael4EEN0-i80uioXxJZOOwQRFHLh3cmfDxNk,1963
 ornith_g256/_vllm_correctness/qwen_gdn_linear_attn.py,sha256=Plj93cegtCrj2bXquzdamTpBELrUmcN4Rr3kZpAdvrM,78334
+ornith_g256/_xgrammar_native/build-receipt.json,sha256=O993ZX5RV8NZOLZt-0YmZgit47WZpiznSrs_gH1XnoQ,1654
+ornith_g256/_xgrammar_native/libxgrammar-candidate.so,sha256=kqs2v64gxsTfP1MhyjdSBHmayDMlKSHZ2PaiBkkdNfc,60211248
+ornith_g256/_xgrammar_native/selection.json,sha256=fLk7P-kC3m1gS_lvef3D772NYS5JTag49qjw1zHria4,106
 ornith_g256/adaptive_c1.py,sha256=jkj9h0-99lGDwgvMfHSNJWLbyvqX4ATxiigYuYmAqAk,14985
 ornith_g256/attention.py,sha256=L908DthNywUqno-d9AaN34CpaZMb6gopXgItNrdTzaE,772
-ornith_g256/attention_compact.py,sha256=ZrrGSCCgtUdW2amQLV-AZmdaorCz6w7HP_G0O0NAeWM,10412
+ornith_g256/attention_compact.py,sha256=UN9hoCo4kKKCWyBe0Yx9GKKwAjM9OCxXJEC1jKmyuf4,10459
 ornith_g256/attention_fast.py,sha256=J7GhM8rbh_86Y8AnBJ4dqB5BEv1dUmCkRA_ehUkfg4o,12274
 ornith_g256/attention_folded.py,sha256=c9rIIDAyvObO7fIT1-7-WYfJIXaK4jcSniMv79jZu6M,9570
-ornith_g256/attention_iu4.py,sha256=NCNbEQTq6WxVmGs1mFvimK7F7GbE-oFCQJRewD7ax_U,2575
-ornith_g256/attention_iu4_persistent.py,sha256=yivkrUYx6ie2L88HXng732woKS9ZD0nuh0jpkwmmzQ4,12418
+ornith_g256/attention_iu4.py,sha256=FcU0H8hIxDrJVlRuVjz4q392SfBaZ0-wFTsOLXZDjK4,3337
+ornith_g256/attention_iu4_persistent.py,sha256=zvdK5COXQGMhJjeknNZKWhC5B386XiKHdGcidJoeYpc,12418
 ornith_g256/attention_mixed.py,sha256=y7UwaX8lVb5lMrfU00AfOjPExYmCZZ6IiScg0haREQo,4138
 ornith_g256/attention_partition.py,sha256=Qb-0TdDQxGosCnJZgxVQGOBRzYntuFJ_EZWgkr73_90,8642
 ornith_g256/attention_storage.py,sha256=OrgT4iGDU3C9tQOyq31I9GKURBsd2cZSxdyXZDi9tuI,4117
@@ -49,6 +51,28 @@ ornith_g256/moe_base.py,sha256=JvtRdBXySXhZ8uP_Gu2PMefFakqszPs0xmOOZdffB8Q,3683
 ornith_g256/native.py,sha256=Sczg1nmHuWWf61Vqr7JHQjSQwhU7fXe3UPBVpP3NozE,12196
 ornith_g256/native_tool_schema.py,sha256=kQBgaJ1ytHrUp7seUe8wHkH_C-G-ucxrPfAcQj0u3vU,1897
 ornith_g256/no_spec.py,sha256=VhEGt1Mo_g4WH6EkOspM9QxQSdjnhdIskKrLPh4aAQQ,2618
+ornith_g256/optimized/__init__.py,sha256=Apf11Hav3dsFVvTyNvleUU22ublniN9nKkoLwAdP2DU,45
+ornith_g256/optimized/compile_scope.py,sha256=GJcac7OlrYDmIS3n9-QZod0A63OApebIVZ9axvXsjU0,1334
+ornith_g256/optimized/draft.py,sha256=6AEwHPA9aZmHWNNftUY3VMzDChjgMO0rmIsmzvQ_Qpw,4327
+ornith_g256/optimized/draft_kernel.py,sha256=PW2m7wwURQX4iMvTW6ULa5WtjkjDdqtzKXdGv_IokZ4,2569
+ornith_g256/optimized/grammar_parallel.py,sha256=6cNghDhxJSGYDX_NYPWIYtCwoodmIBg_MExkoDgX7eo,5431
+ornith_g256/optimized/greedy.py,sha256=OGRDrYcp-P4h7F_vhOrTrPvU1sXdO0weUKeZXtHRdh8,2212
+ornith_g256/optimized/greedy_kernel.py,sha256=Lr4qJdEBGkXup-2jwJZMoesmON99YPVNsCyRQ1VQ3EI,1930
+ornith_g256/optimized/launch.py,sha256=wdQmAhVGdJFqlLmQe3dG-sPqSMxQX6abq_EW8PEn4D0,435
+ornith_g256/optimized/mixed_qk_sum.py,sha256=F7GHaa6EwulwegJXGB5cMU6XdPw7R5b45nh6gwLLXh8,2761
+ornith_g256/optimized/pipeline/SOURCE-PROOF.json,sha256=WTm8ilOoMXFdBdD6ApC-qgywDtoVl_478fHPxTaZ_Q0,1709
+ornith_g256/optimized/pipeline/__init__.py,sha256=Apf11Hav3dsFVvTyNvleUU22ublniN9nKkoLwAdP2DU,45
+ornith_g256/optimized/pipeline/adapter.py,sha256=avl6jV4Nebffwy8vD-X1xTtcHFDQNUL_mcZMrn85CeE,9092
+ornith_g256/optimized/pipeline/candidate/causal_conv1d.py,sha256=JUxUD3ewFpRpahq1NavQ5aI9VOchHiAs-b7q0E-s4Ws,51774
+ornith_g256/optimized/pipeline/candidate/chunk_o.py,sha256=qT-VEEiVgmW00cP55nNBMpGW56V1JUM5NCFf3ksW5hs,5392
+ornith_g256/optimized/pipeline/candidate/chunk_scaled_dot_kkt.py,sha256=cXNdKxxKGe7LtgWHCEgYF7thBCbEj5al9kZ4Svzwxlw,5720
+ornith_g256/optimized/pipeline/candidate/solve_tril.py,sha256=vdBLacD5FWIEjSWDS9U9KY3B8b0Gi5SUr4h3ayv0sdE,19567
+ornith_g256/optimized/pipeline/candidate/wy_fast.py,sha256=TMC-UZv884EtgQpL67PdRfubE6jQq2LQrUY6Gi429h8,4424
+ornith_g256/optimized/postconv.py,sha256=xY7l0rhBj79p7aXZ7k80KxHMASq3euJbbOKyrqUapQ8,7935
+ornith_g256/optimized/prefix.py,sha256=xpWomasStYw7ibbswnboLvo_CDOjGZEahSjMTiN3EM0,2353
+ornith_g256/optimized/qk.py,sha256=2G8alAhmaKxkAursehnArY5V3I5jd_jUqe6raqlvlmQ,4957
+ornith_g256/optimized/qk_schedule.py,sha256=ZgOHgfBmdIb7pFjvMRjKia2WQNwsK9Bc-FdkQPDxeKI,5793
+ornith_g256/optimized/worker.py,sha256=OFebdnI15_ysPy5rBQdDjxVcE04ehfPRwgWkO-zpVRI,2988
 ornith_g256/parser_contract.py,sha256=LOurQD3IglafPOZ9EnmN8Uzr1cFixy-Cvs03vN1-Kig,4203
 ornith_g256/phase_dispatch.py,sha256=Gk698q9wnmylGeCgj2pTdgLfF4yldTMgRWF2MWQPgTc,2184
 ornith_g256/prefill_draft.py,sha256=15uymhJGOVQ6ZOGYkI4urvTLz0kvLl4U3rz8nIZ4FvU,5630
@@ -58,6 +82,7 @@ ornith_g256/runtime_correctness.py,sha256=CPQXc-YjpPmc-KsZ5DTX4zusnPpesFqdwPRRGq
 ornith_g256/strict_qwen/__init__.py,sha256=7opQVSSfFBP8IvRkIat3tUhao5sZnGjrW1bHXcxAki0,81
 ornith_g256/strict_qwen/qwen3.py,sha256=M8ZTRHoUS3IlflCMzZJhDbbTBXjij5jGnJkoF3PerRw,11658
 ornith_g256/strict_qwen/qwen3_contract.py,sha256=KcCDZZxUyBFE_bbIMpx1qqYs9p0cRfp9qL3u7KSFbq8,8719
-ornith_g256/worker.py,sha256=En2VXcVk520ONHYD_aQmSaisvXJHfIMfg5Z5g--n63o,13579
+ornith_g256/worker.py,sha256=rMJyQMIfNHjKRS3fBCRN9O53DvqsOBku9k0yjvInAdU,13579
 ornith_g256/worker_base.py,sha256=oA7xKVFy5lXIs8fXkM_QMWEUqUkuSD3t_eoDbTt7-Gg,4601
-ciru_ornith_g256-1.0.2.dist-info/RECORD,,
+ornith_g256/xgrammar_native.py,sha256=0ZxUbhJMsRMR5Mwx9YM9xmPieJ1FCQCRwPj-yKliZiM,2059
+ciru_ornith_g256-4.0.0.dist-info/RECORD,,

bundle/plugin-site/ornith_g256/__init__.py

@@ -1,6 +1,9 @@
 """Ciru G256 prototype; no installed vLLM files are modified."""
 
 
+from .xgrammar_native import install as _install_native_grammar
+_install_native_grammar()
+
 from .runtime_correctness import install as _install_correctness
 _install_correctness()
 
@@ -30,3 +33,6 @@ def register():
 
     from .adaptive_c1 import install as install_adaptive_c1
     install_adaptive_c1()
+
+    from .optimized.grammar_parallel import install as install_parallel_grammar
+    install_parallel_grammar()

bundle/plugin-site/ornith_g256/_vllm_correctness/chat_completion_serving_r03.py

@@ -0,0 +1,1319 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import asyncio
+import time
+from collections.abc import AsyncGenerator, AsyncIterator
+from collections.abc import Sequence as GenericSequence
+from http import HTTPStatus
+from typing import Any, Final, cast
+
+from fastapi import Request
+
+from vllm.engine.protocol import EngineClient
+from vllm.entrypoints.chat_utils import (
+    ChatTemplateContentFormatOption,
+    ConversationMessage,
+    make_tool_call_id,
+)
+from vllm.entrypoints.generate.base.serving import (
+    GenerateBaseServing,
+    GenerationError,
+    build_per_request_timing_metrics,
+    build_spec_decoding_metrics,
+    clamp_prompt_logprobs,
+    format_token_id_placeholder,
+)
+from vllm.entrypoints.openai.chat_completion.protocol import (
+    ChatCompletionLogProb,
+    ChatCompletionLogProbs,
+    ChatCompletionLogProbsContent,
+    ChatCompletionNamedToolChoiceParam,
+    ChatCompletionRequest,
+    ChatCompletionResponse,
+    ChatCompletionResponseChoice,
+    ChatCompletionResponseStreamChoice,
+    ChatCompletionStreamResponse,
+    ChatMessage,
+)
+from vllm.entrypoints.openai.engine.protocol import (
+    CompletionTokenUsageInfo,
+    DeltaMessage,
+    ErrorResponse,
+    FunctionCall,
+    PerRequestMetrics,
+    PromptTokenUsageInfo,
+    RequestResponseMetadata,
+    ToolCall,
+    UsageInfo,
+)
+from vllm.entrypoints.openai.models.serving import OpenAIServingModels
+from vllm.entrypoints.serve.utils.api_utils import get_max_tokens, should_include_usage
+from vllm.entrypoints.serve.utils.request_logger import RequestLogger
+from vllm.entrypoints.serve.utils.tool_calls_utils import (
+    maybe_filter_parallel_tool_calls,
+)
+from vllm.inputs import EngineInput, MultiModalPlaceholders
+from vllm.logger import init_logger
+from vllm.logprobs import Logprob
+from vllm.outputs import RequestOutput
+from vllm.parser import ParserManager
+from vllm.parser.abstract_parser import Parser
+from vllm.renderers.online_renderer import OnlineRenderer
+from vllm.sampling_params import BeamSearchParams, SamplingParams
+from vllm.tokenizers import TokenizerLike
+from vllm.utils.collection_utils import as_list
+from vllm.utils.serial_utils import numpy2base64
+
+logger = init_logger(__name__)
+
+
+def _get_mm_token_counts(engine_input: EngineInput) -> dict[str, int]:
+    """Sum per-modality placeholder tokens from ``mm_placeholders``.
+
+    Keyed by modality name; ``PlaceholderRange.length`` is the placeholder's
+    prompt token span, so each sum matches the placeholder tokens already
+    counted in ``usage.prompt_tokens``.
+    """
+    mm_placeholders = cast(
+        MultiModalPlaceholders | None, engine_input.get("mm_placeholders")
+    )
+    return {
+        modality: sum(p.length for p in ranges)
+        for modality, ranges in (mm_placeholders or {}).items()
+        if ranges
+    }
+
+
+def _make_prompt_tokens_details(
+    enable_prompt_tokens_details: bool,
+    num_cached_tokens: int | None,
+    num_cache_creation_tokens: int | None,
+    mm_token_counts: dict[str, int] | None,
+) -> PromptTokenUsageInfo | None:
+    """Build ``prompt_tokens_details`` from cached + multimodal token counts."""
+    if not enable_prompt_tokens_details:
+        return None
+    if (
+        num_cached_tokens is None
+        and num_cache_creation_tokens is None
+        and not mm_token_counts
+    ):
+        return None
+    return PromptTokenUsageInfo(
+        cached_tokens=num_cached_tokens,
+        created_cache_tokens=num_cache_creation_tokens,
+        multimodal_tokens=mm_token_counts or None,
+    )
+
+
+def _make_completion_tokens_details(
+    reasoning_tokens: int,
+) -> CompletionTokenUsageInfo:
+    return CompletionTokenUsageInfo(reasoning_tokens=reasoning_tokens)
+
+
+class OpenAIServingChat(GenerateBaseServing):
+    def __init__(
+        self,
+        engine_client: EngineClient,
+        models: OpenAIServingModels,
+        response_role: str,
+        *,
+        online_renderer: "OnlineRenderer",
+        request_logger: RequestLogger | None,
+        chat_template: str | None,
+        chat_template_content_format: ChatTemplateContentFormatOption,
+        trust_request_chat_template: bool = False,
+        return_tokens_as_token_ids: bool = False,
+        reasoning_parser: str = "",
+        enable_auto_tools: bool = False,
+        exclude_tools_when_tool_choice_none: bool = False,
+        tool_parser: str | None = None,
+        enable_prompt_tokens_details: bool = False,
+        enable_force_include_usage: bool = False,
+        enable_log_outputs: bool = False,
+        enable_log_deltas: bool = True,
+        default_chat_template_kwargs: dict[str, Any] | None = None,
+        enable_per_request_metrics: bool = False,
+    ) -> None:
+        super().__init__(
+            engine_client=engine_client,
+            models=models,
+            request_logger=request_logger,
+            return_tokens_as_token_ids=return_tokens_as_token_ids,
+        )
+
+        self.online_renderer = online_renderer
+        self.response_role = response_role
+        self.chat_template = chat_template
+        self.chat_template_content_format: Final = chat_template_content_format
+        self.trust_request_chat_template = trust_request_chat_template
+        self.default_chat_template_kwargs = default_chat_template_kwargs or {}
+        self.enable_log_outputs = enable_log_outputs
+        self.enable_log_deltas = enable_log_deltas
+
+        self.enable_auto_tools: bool = enable_auto_tools
+        self._include_reasoning_tokens_details = bool(reasoning_parser)
+        self.parser_cls = ParserManager.get_parser(
+            tool_parser_name=tool_parser,
+            reasoning_parser_name=reasoning_parser,
+            enable_auto_tools=enable_auto_tools,
+            model_name=self.model_config.model,
+            is_harmony=self.model_config.hf_config.model_type == "gpt_oss",
+        )
+        self.exclude_tools_when_tool_choice_none = exclude_tools_when_tool_choice_none
+
+        self.enable_prompt_tokens_details = enable_prompt_tokens_details
+        self.enable_force_include_usage = enable_force_include_usage
+        self.enable_per_request_metrics = enable_per_request_metrics
+        self.default_sampling_params = self.model_config.get_diff_sampling_param()
+        mc = self.model_config
+        self.override_max_tokens = (
+            self.default_sampling_params.get("max_tokens")
+            if mc.generation_config not in ("auto", "vllm")
+            else getattr(mc, "override_generation_config", {}).get("max_new_tokens")
+        )
+        # NOTE(woosuk): While OpenAI's chat completion API supports browsing
+        # for some models, currently vLLM doesn't support it. Please use the
+        # Responses API instead.
+        self.supports_browsing = False
+        self.browser_tool = None
+        # NOTE(woosuk): Chat completion API does not support code interpreter.
+        # Please use the Responses API instead.
+        self.supports_code_interpreter = False
+        self.python_tool = None
+
+    def _effective_chat_template_kwargs(
+        self, request: ChatCompletionRequest
+    ) -> dict[str, Any]:
+        return (
+            request.build_chat_params(
+                self.chat_template,
+                self.chat_template_content_format,
+            )
+            .with_defaults(self.default_chat_template_kwargs)
+            .chat_template_kwargs
+        )
+
+    def _engine_chat_template_kwargs(
+        self, chat_template_kwargs: dict[str, Any]
+    ) -> dict[str, Any]:
+        """Subclass hook to narrow ``chat_template_kwargs`` for the engine.
+
+        The same dict is used twice: to build the API-server-side parser
+        instance, and to populate
+        ``EngineCoreRequest.reasoning_parser_kwargs`` for the engine-core
+        side. The latter crosses ZMQ as msgpack, so a handler that stashes
+        request-scoped state only the API-server-side parser needs (values
+        msgpack can't encode, or payloads not worth shipping) can drop
+        those entries here without affecting the in-process parser.
+
+        Must not mutate the argument -- the caller still needs the full
+        dict. The default forwards it unchanged.
+        """
+        return chat_template_kwargs
+
+    async def render_chat_request(
+        self,
+        request: ChatCompletionRequest,
+    ) -> tuple[list[ConversationMessage], list[EngineInput]] | ErrorResponse:
+        """
+        Validate the model and preprocess a chat completion request.
+
+        Delegates preprocessing logic to OnlineRenderer, adding the
+        engine-aware checks (LoRA model validation, engine health).
+
+        Returns:
+            A tuple of (conversation, engine_inputs) on success,
+            or an ErrorResponse on failure.
+        """
+        error_check_ret = await self._check_model(request)
+        if error_check_ret is not None:
+            logger.error("Error with model %s", error_check_ret)
+            return error_check_ret
+
+        # If the engine is dead, raise the engine's DEAD_ERROR.
+        # This is required for the streaming case, where we return a
+        # success status before we actually start generating text :).
+        if self.engine_client.errored:
+            raise self.engine_client.dead_error
+
+        return await self.online_renderer.render_chat(request)
+
+    async def create_chat_completion(
+        self,
+        request: ChatCompletionRequest,
+        raw_request: Request | None = None,
+    ) -> AsyncGenerator[str, None] | ChatCompletionResponse | ErrorResponse:
+        """
+        Chat Completion API similar to OpenAI's API.
+
+        See https://platform.openai.com/docs/api-reference/chat/create
+        for the API specification. This API mimics the OpenAI
+        Chat Completion API.
+        """
+        return await self._with_kv_transfer_rejection_cleanup(
+            self._create_chat_completion(request, raw_request), request, raw_request
+        )
+
+    async def _create_chat_completion(
+        self,
+        request: ChatCompletionRequest,
+        raw_request: Request | None = None,
+    ) -> AsyncGenerator[str, None] | ChatCompletionResponse | ErrorResponse:
+        # Streaming response
+        tokenizer = self.renderer.tokenizer
+        assert tokenizer is not None
+        chat_template_kwargs = self._effective_chat_template_kwargs(request)
+        parser: Parser | None = None
+        if self.parser_cls is not None:
+            parser = self.parser_cls(
+                tokenizer,
+                request.tools,
+                chat_template_kwargs=chat_template_kwargs,
+                model_config=self.model_config,
+            )
+        result = await self.render_chat_request(request)
+        if isinstance(result, ErrorResponse):
+            return result
+
+        conversation, engine_inputs = result
+
+        request_id = (
+            f"chatcmpl-{self._base_request_id(raw_request, request.request_id)}"
+        )
+
+        request_metadata = RequestResponseMetadata(request_id=request_id)
+        if raw_request:
+            raw_request.state.request_metadata = request_metadata
+
+        lora_request = self._maybe_get_adapters(request, supports_default_mm_loras=True)
+
+        model_name = self.models.model_name(lora_request)
+
+        # Extract data_parallel_rank from header (router can inject it)
+        data_parallel_rank = self._get_data_parallel_rank(raw_request)
+
+        # Schedule the request and get the result generator.
+        max_model_len = self.model_config.max_model_len
+        generators: list[AsyncGenerator[RequestOutput, None]] = []
+        mm_token_counts: dict[str, int] | None = None
+        for i, engine_input in enumerate(engine_inputs):
+            prompt_token_ids = self._extract_prompt_components(engine_input).token_ids
+            mm_token_counts = _get_mm_token_counts(engine_input)
+
+            # If we are creating sub requests for multiple prompts, ensure that they
+            # have unique request ids.
+            sub_request_id = (
+                request_id if len(engine_inputs) == 1 else f"{request_id}_{i}"
+            )
+
+            max_tokens = get_max_tokens(
+                max_model_len,
+                request.max_completion_tokens
+                if request.max_completion_tokens is not None
+                else request.max_tokens,
+                self._extract_prompt_len(engine_input),
+                self.default_sampling_params,
+                self.override_max_tokens,
+                truncate_prompt_tokens=request.truncate_prompt_tokens,
+            )
+
+            sampling_params: SamplingParams | BeamSearchParams
+            if request.use_beam_search:
+                sampling_params = request.to_beam_search_params(
+                    max_tokens, self.default_sampling_params
+                )
+            else:
+                sampling_params = request.to_sampling_params(
+                    max_tokens,
+                    self.default_sampling_params,
+                )
+
+            self._log_inputs(
+                sub_request_id,
+                engine_input,
+                params=sampling_params,
+                lora_request=lora_request,
+            )
+
+            trace_headers = (
+                None
+                if raw_request is None
+                else await self._get_trace_headers(raw_request.headers)
+            )
+            session_id = self._get_session_id(request, raw_request)
+
+            if isinstance(sampling_params, BeamSearchParams):
+                generator = self.beam_search(
+                    prompt=engine_input,
+                    request_id=sub_request_id,
+                    params=sampling_params,
+                    lora_request=lora_request,
+                    trace_headers=trace_headers,
+                    session_id=session_id,
+                )
+            else:
+                if not request.include_reasoning:
+                    reasoning_ended = True
+                elif request._grammar_from_parser:
+                    # The Mistral grammar already includes an optional
+                    # `think?` rule that handles both reasoning and
+                    # non-reasoning outputs.
+                    reasoning_ended = True
+                elif parser is not None and parser.reasoning_parser is not None:
+                    reasoning_ended = parser.is_reasoning_end(prompt_token_ids or [])
+                else:
+                    reasoning_ended = None
+
+                generator = self.engine_client.generate(
+                    engine_input,
+                    sampling_params,
+                    sub_request_id,
+                    lora_request=lora_request,
+                    trace_headers=trace_headers,
+                    priority=self._get_priority(request, raw_request),
+                    data_parallel_rank=data_parallel_rank,
+                    session_id=session_id,
+                    reasoning_ended=reasoning_ended,
+                    reasoning_parser_kwargs={
+                        "chat_template_kwargs": self._engine_chat_template_kwargs(
+                            chat_template_kwargs
+                        ),
+                    }
+                    if parser is not None and parser.reasoning_parser is not None
+                    else None,
+                )
+
+            generators.append(generator)
+
+        assert len(generators) == 1
+        (result_generator,) = generators
+
+        if request.stream:
+            return self.chat_completion_stream_generator(
+                request,
+                result_generator,
+                request_id,
+                model_name,
+                conversation,
+                tokenizer,
+                request_metadata,
+                chat_template_kwargs=chat_template_kwargs,
+                mm_token_counts=mm_token_counts,
+            )
+
+        return await self.chat_completion_full_generator(
+            request,
+            result_generator,
+            request_id,
+            model_name,
+            conversation,
+            tokenizer,
+            request_metadata,
+            parser=parser,
+            mm_token_counts=mm_token_counts,
+        )
+
+    def get_chat_request_role(self, request: ChatCompletionRequest) -> str:
+        if request.add_generation_prompt:
+            return self.response_role
+        return request.messages[-1]["role"]
+
+    def _create_chat_message(self, *args: Any, **kwargs: Any) -> ChatMessage:
+        """Construct the response :class:`ChatMessage` for the non-streaming path.
+
+        The full-generator calls this at every construction site so
+        subclasses can swap in a specialized :class:`ChatMessage`
+        subclass (e.g. :class:`CohereServingChatV2` returning
+        :class:`CohereChatMessage`) without duplicating the branchy
+        tool-choice / auto-tools logic that decides which fields are
+        populated. The default returns a plain :class:`ChatMessage`.
+        """
+        return ChatMessage(*args, **kwargs)
+
+    def _finalize_response_message(
+        self,
+        message: ChatMessage,
+        *,
+        parser: Parser | None,
+    ) -> ChatMessage:
+        """Subclass hook to enrich a fully-constructed :class:`ChatMessage`.
+
+        Default is a no-op. Subclasses that need to surface parser-side
+        extras (e.g. :class:`CohereServingChatV2` reading grounding
+        citations off the reasoning parser and populating
+        :class:`CohereChatMessage.citations`) override this to inspect
+        ``parser`` and mutate/replace ``message``.
+        """
+        return message
+
+    async def chat_completion_stream_generator(
+        self,
+        request: ChatCompletionRequest,
+        result_generator: AsyncIterator[RequestOutput],
+        request_id: str,
+        model_name: str,
+        conversation: list[ConversationMessage],
+        tokenizer: TokenizerLike,
+        request_metadata: RequestResponseMetadata,
+        chat_template_kwargs: dict[str, Any] | None = None,
+        mm_token_counts: dict[str, int] | None = None,
+    ) -> AsyncGenerator[str, None]:
+        created_time = int(time.time())
+        chunk_object_type: Final = "chat.completion.chunk"
+        first_iteration = True
+
+        # Send response for each token for each request.n (index)
+        num_choices = 1 if request.n is None else request.n
+        previous_num_tokens = [0] * num_choices
+        # TODO: Remove once all reasoning parsers use the Parser Engine.
+        generated_token_ids: list[list[int]] = [[] for _ in range(num_choices)]
+        previous_reasoning_tokens = [0] * num_choices
+        finish_reason_sent = [False] * num_choices
+        num_prompt_tokens = 0
+        num_cached_tokens = None
+        num_cache_creation_tokens = None
+        tools_streamed = [False] * num_choices
+
+        if isinstance(request.tool_choice, ChatCompletionNamedToolChoiceParam):
+            tool_choice_function_name = request.tool_choice.function.name
+        else:
+            tool_choice_function_name = None
+
+        previous_texts = [""] * num_choices
+
+        try:
+            if self.parser_cls is not None:
+                if tokenizer is None:
+                    raise ValueError(
+                        "Tokenizer not available when `skip_tokenizer_init=True`"
+                    )
+                parsers: list[Parser | None] = [
+                    self.parser_cls(
+                        tokenizer,
+                        request.tools,
+                        chat_template_kwargs=chat_template_kwargs,
+                        model_config=self.model_config,
+                    )
+                    for _ in range(num_choices)
+                ]
+            else:
+                parsers = [None] * num_choices
+        except Exception as e:
+            logger.exception("Error in parser creation.")
+            data = self.create_streaming_error_response(e)
+            yield f"data: {data}\n\n"
+            yield "data: [DONE]\n\n"
+            return
+
+        stream_options = request.stream_options
+        include_usage, include_continuous_usage = should_include_usage(
+            stream_options, self.enable_force_include_usage
+        )
+
+        last_res: RequestOutput | None = None
+        try:
+            async for res in result_generator:
+                last_res = res
+                if res.prompt_token_ids is not None:
+                    num_prompt_tokens = len(res.prompt_token_ids)
+                    if res.encoder_prompt_token_ids is not None:
+                        num_prompt_tokens += len(res.encoder_prompt_token_ids)
+
+                # We need to do it here, because if there are exceptions in
+                # the result_generator, it needs to be sent as the FIRST
+                # response (by the try...catch).
+                if first_iteration:
+                    num_cached_tokens = res.num_cached_tokens
+                    num_cache_creation_tokens = res.num_cache_creation_tokens
+                    # Send first response for each request.n (index) with
+                    # the role
+                    role = self.get_chat_request_role(request)
+
+                    # ``res.prompt`` is the rendered chat-templated prompt
+                    prompt_text = res.prompt if request.return_prompt_text else None
+
+                    # NOTE num_choices defaults to 1 so this usually executes
+                    # once per request
+                    for i in range(num_choices):
+                        choice_data = ChatCompletionResponseStreamChoice(
+                            index=i,
+                            delta=DeltaMessage(
+                                role=role,
+                                content="",
+                            ),
+                            logprobs=None,
+                            finish_reason=None,
+                        )
+
+                        # return prompt_token_ids at the first chunk ever
+                        chunk = ChatCompletionStreamResponse(
+                            id=request_id,
+                            object=chunk_object_type,
+                            created=created_time,
+                            choices=[choice_data],
+                            model=model_name,
+                            prompt_token_ids=(
+                                res.prompt_token_ids
+                                if request.return_token_ids
+                                else None
+                            ),
+                            prompt_text=prompt_text,
+                        )
+
+                        # if continuous usage stats are requested, add it
+                        if include_continuous_usage:
+                            chunk.usage = UsageInfo(
+                                prompt_tokens=num_prompt_tokens,
+                                completion_tokens=0,
+                                total_tokens=num_prompt_tokens,
+                                completion_tokens_details=(
+                                    _make_completion_tokens_details(0)
+                                    if self._include_reasoning_tokens_details
+                                    else None
+                                ),
+                            )
+
+                        data = chunk.model_dump_json(exclude_unset=True)
+                        yield f"data: {data}\n\n"
+
+                    # Send response to echo the input portion of the
+                    # last message
+                    if request.echo:
+                        last_msg_content: str | list[dict[str, str]] = ""
+                        if (
+                            conversation
+                            and "content" in conversation[-1]
+                            and conversation[-1].get("role") == role
+                        ):
+                            last_msg_content = conversation[-1]["content"] or ""
+
+                        if last_msg_content:
+                            for i in range(num_choices):
+                                choice_data = ChatCompletionResponseStreamChoice(
+                                    index=i,
+                                    delta=DeltaMessage(content=last_msg_content),
+                                    logprobs=None,
+                                    finish_reason=None,
+                                )
+                                chunk = ChatCompletionStreamResponse(
+                                    id=request_id,
+                                    object=chunk_object_type,
+                                    created=created_time,
+                                    choices=[choice_data],
+                                    model=model_name,
+                                )
+                                if include_continuous_usage:
+                                    chunk.usage = UsageInfo(
+                                        prompt_tokens=num_prompt_tokens,
+                                        completion_tokens=0,
+                                        total_tokens=num_prompt_tokens,
+                                        completion_tokens_details=(
+                                            _make_completion_tokens_details(0)
+                                            if self._include_reasoning_tokens_details
+                                            else None
+                                        ),
+                                    )
+
+                                data = chunk.model_dump_json(exclude_unset=True)
+                                yield f"data: {data}\n\n"
+                    first_iteration = False
+
+                for output in res.outputs:
+                    i = output.index
+                    parser = parsers[i]
+                    if finish_reason_sent[i]:
+                        continue
+
+                    if request.logprobs and (
+                        request.top_logprobs is not None or request.logprob_token_ids
+                    ):
+                        assert output.logprobs is not None, "Did not output logprobs"
+                        logprobs = self._create_chat_logprobs(
+                            token_ids=output.token_ids,
+                            top_logprobs=output.logprobs,
+                            tokenizer=tokenizer,
+                            num_output_top_logprobs=request.top_logprobs,
+                            logprob_token_ids=request.logprob_token_ids,
+                            return_as_token_id=request.return_tokens_as_token_ids,
+                        )
+                    else:
+                        logprobs = None
+
+                    delta_text = output.text
+
+                    if (
+                        not delta_text
+                        and not output.token_ids
+                        and not previous_num_tokens[i]
+                    ):
+                        # Chunked prefill case, don't return empty chunks
+                        continue
+
+                    delta_message: DeltaMessage | None
+
+                    if parser is not None:
+                        delta_message = parser.parse_delta(
+                            delta_text=delta_text,
+                            delta_token_ids=as_list(output.token_ids),
+                            request=request,
+                            prompt_token_ids=res.prompt_token_ids,
+                            finished=output.finish_reason is not None,
+                        )
+                        if delta_message is not None and delta_message.tool_calls:
+                            tools_streamed[i] = True
+
+                    # handle streaming just a content delta (no parsers)
+                    else:
+                        delta_message = DeltaMessage(content=delta_text)
+
+                    previous_texts[i] += delta_text
+
+                    # set the previous values for the next iteration
+                    previous_num_tokens[i] += len(output.token_ids)
+                    if parser is not None:
+                        generated_token_ids[i].extend(output.token_ids)
+                        previous_reasoning_tokens[i] = parser.count_reasoning_tokens(
+                            tuple(generated_token_ids[i])
+                        )
+
+                    # if the message delta is None (e.g. because it was a
+                    # "control token" for tool calls or the parser otherwise
+                    # wasn't ready to send a token, then
+                    #   get the next token without streaming a chunk
+                    # When reasoning is hidden, suppress per-token
+                    # metadata (logprobs, token_ids) on every chunk to
+                    # prevent leaking reasoning tokens through decoded
+                    # token text in logprob entries or raw token IDs.
+                    hide_stream_metadata = (
+                        not request.include_reasoning and parser is not None
+                    )
+                    if hide_stream_metadata:
+                        logprobs = None
+
+                    if delta_message is None:
+                        # NOTE: If return_token_ids is enabled, we still need to
+                        # send a chunk with token_ids even if delta_message is None
+                        # to ensure all tokens are included in the response
+                        # A parser may buffer a whole JSON array/object.
+                        # Real token progress is independent of argument deltas.
+                        # Preserve requested usage without exposing partial calls
+                        # or hidden reasoning metadata. Empty prefill stays quiet.
+                        if (
+                            output.finish_reason is None
+                            and not (include_continuous_usage and output.token_ids)
+                            and (not request.return_token_ids or hide_stream_metadata)
+                        ):
+                            continue
+                        delta_message = DeltaMessage()
+
+                    # Log streaming delta if output logging is enabled
+                    if self.enable_log_outputs and self.request_logger:
+                        delta_content_parts = []
+                        if delta_message.content:
+                            delta_content_parts.append(delta_message.content)
+                        if delta_message.reasoning:
+                            reasoning = delta_message.reasoning
+                            delta_content_parts.append(f"[reasoning: {reasoning}]")
+                        if delta_message.tool_calls:
+                            tool_args = "".join(
+                                tc.function.arguments
+                                for tc in delta_message.tool_calls
+                                if tc.function and tc.function.arguments
+                            )
+                            if tool_args:
+                                delta_content_parts.append(f"[tool_calls: {tool_args}]")
+
+                        if delta_content_parts and self.enable_log_deltas:
+                            delta_content = " ".join(delta_content_parts)
+                            self.request_logger.log_outputs(
+                                request_id=request_id,
+                                outputs=delta_content,
+                                output_token_ids=as_list(output.token_ids),
+                                finish_reason=output.finish_reason,
+                                is_streaming=True,
+                                delta=True,
+                            )
+
+                    include_token_ids = (
+                        request.return_token_ids and not hide_stream_metadata
+                    )
+
+                    if output.finish_reason is None:
+                        # Send token-by-token response for each request.n
+                        choice_data = ChatCompletionResponseStreamChoice(
+                            index=i,
+                            delta=delta_message,
+                            logprobs=logprobs,
+                            finish_reason=None,
+                            token_ids=(
+                                as_list(output.token_ids) if include_token_ids else None
+                            ),
+                        )
+
+                    # if the model is finished generating
+                    else:
+                        # check for error finish reason and abort streaming
+                        # finish_reason='error' indicates a retryable error
+                        self._raise_if_error(output.finish_reason, request_id)
+
+                        # Send the finish response for each request.n only once
+                        # In OpenAI's API, when a tool is called, the
+                        # finish_reason is:
+                        # "tool_calls" for "auto" or "required" tool calls,
+                        # and "stop" for named tool calls.
+                        if (
+                            tools_streamed[i]
+                            and not tool_choice_function_name
+                            and output.finish_reason == "stop"
+                        ):
+                            finish_reason_ = "tool_calls"
+                        else:
+                            finish_reason_ = (
+                                output.finish_reason if output.finish_reason else "stop"
+                            )
+                        choice_data = ChatCompletionResponseStreamChoice(
+                            index=i,
+                            delta=delta_message,
+                            logprobs=logprobs,
+                            finish_reason=finish_reason_,
+                            stop_reason=output.stop_reason,
+                            token_ids=(
+                                as_list(output.token_ids) if include_token_ids else None
+                            ),
+                        )
+
+                        finish_reason_sent[i] = True
+
+                    choice_data = maybe_filter_parallel_tool_calls(choice_data, request)
+                    chunk = ChatCompletionStreamResponse(
+                        id=request_id,
+                        object=chunk_object_type,
+                        created=created_time,
+                        choices=[choice_data],
+                        model=model_name,
+                    )
+                    # Stamp the fingerprint on terminal chunks only (those with
+                    # finish_reason set). When ``include_usage`` is on, the
+                    # trailing usage chunk below overrides this as the true
+                    # final message.
+                    if (
+                        not include_usage
+                        and self.system_fingerprint is not None
+                        and choice_data.finish_reason is not None
+                    ):
+                        chunk.system_fingerprint = self.system_fingerprint
+
+                    # handle usage stats if requested & if continuous
+                    if include_continuous_usage:
+                        completion_tokens = previous_num_tokens[i]
+                        chunk.usage = UsageInfo(
+                            prompt_tokens=num_prompt_tokens,
+                            completion_tokens=completion_tokens,
+                            total_tokens=num_prompt_tokens + completion_tokens,
+                            completion_tokens_details=(
+                                _make_completion_tokens_details(
+                                    previous_reasoning_tokens[i]
+                                )
+                                if self._include_reasoning_tokens_details
+                                else None
+                            ),
+                        )
+
+                    data = chunk.model_dump_json(exclude_unset=True)
+                    yield f"data: {data}\n\n"
+
+            # once the final token is handled, if stream_options.include_usage
+            # is sent, send the usage
+            if include_usage:
+                completion_tokens = sum(previous_num_tokens)
+                final_usage = UsageInfo(
+                    prompt_tokens=num_prompt_tokens,
+                    completion_tokens=completion_tokens,
+                    total_tokens=num_prompt_tokens + completion_tokens,
+                    completion_tokens_details=_make_completion_tokens_details(
+                        sum(previous_reasoning_tokens)
+                    )
+                    if self._include_reasoning_tokens_details
+                    else None,
+                )
+                final_usage.prompt_tokens_details = _make_prompt_tokens_details(
+                    self.enable_prompt_tokens_details,
+                    num_cached_tokens,
+                    num_cache_creation_tokens,
+                    mm_token_counts,
+                )
+
+                # In streaming, metrics ride on this final usage chunk, which is
+                # only emitted when usage reporting is enabled (i.e.
+                # ``stream_options.include_usage=true`` or
+                # ``--enable-force-include-usage``).
+                stream_per_request_metrics: PerRequestMetrics | None = None
+                # See note in chat_completion_full_generator: suppress for n>1.
+                if (request.n or 1) == 1:
+                    if self.enable_per_request_metrics:
+                        last_metrics = (
+                            last_res.metrics if last_res is not None else None
+                        )
+                        stream_per_request_metrics = build_per_request_timing_metrics(
+                            last_metrics, completion_tokens
+                        )
+                    spec_stats = build_spec_decoding_metrics(last_res)
+                    if spec_stats is not None:
+                        if stream_per_request_metrics is None:
+                            stream_per_request_metrics = PerRequestMetrics()
+                        stream_per_request_metrics.speculative_decoding = spec_stats
+
+                final_usage_chunk = ChatCompletionStreamResponse(
+                    id=request_id,
+                    object=chunk_object_type,
+                    created=created_time,
+                    choices=[],
+                    model=model_name,
+                    usage=final_usage,
+                    system_fingerprint=self.system_fingerprint,
+                    metrics=stream_per_request_metrics,
+                )
+                final_usage_data = final_usage_chunk.model_dump_json(
+                    exclude_unset=True, exclude_none=True
+                )
+                yield f"data: {final_usage_data}\n\n"
+
+            # report to FastAPI middleware aggregate usage across all choices
+            num_completion_tokens = sum(previous_num_tokens)
+            request_metadata.final_usage_info = UsageInfo(
+                prompt_tokens=num_prompt_tokens,
+                completion_tokens=num_completion_tokens,
+                total_tokens=num_prompt_tokens + num_completion_tokens,
+                completion_tokens_details=_make_completion_tokens_details(
+                    sum(previous_reasoning_tokens)
+                )
+                if self._include_reasoning_tokens_details
+                else None,
+            )
+
+            # Log complete streaming response if output logging is enabled
+            if self.enable_log_outputs and self.request_logger:
+                # Log the complete response for each choice
+                for i in range(num_choices):
+                    full_text = (
+                        previous_texts[i]
+                        if previous_texts and i < len(previous_texts)
+                        else f"<streaming_complete: {previous_num_tokens[i]} tokens>"
+                    )
+                    self.request_logger.log_outputs(
+                        request_id=request_id,
+                        outputs=full_text,
+                        output_token_ids=None,  # Consider also logging all token IDs
+                        finish_reason="streaming_complete",
+                        is_streaming=True,
+                        delta=False,
+                    )
+
+        except GenerationError as e:
+            yield f"data: {self._convert_generation_error_to_streaming_response(e)}\n\n"
+        except Exception as e:
+            logger.exception("Error in chat completion stream generator.")
+            data = self.create_streaming_error_response(e)
+            yield f"data: {data}\n\n"
+        # Send the final done message after all response.n are finished
+        yield "data: [DONE]\n\n"
+
+    async def chat_completion_full_generator(
+        self,
+        request: ChatCompletionRequest,
+        result_generator: AsyncIterator[RequestOutput],
+        request_id: str,
+        model_name: str,
+        conversation: list[ConversationMessage],
+        tokenizer: TokenizerLike,
+        request_metadata: RequestResponseMetadata,
+        parser: Parser | None = None,
+        mm_token_counts: dict[str, int] | None = None,
+    ) -> ErrorResponse | ChatCompletionResponse:
+        created_time = int(time.time())
+        final_res: RequestOutput | None = None
+
+        try:
+            async for res in result_generator:
+                final_res = res
+        except asyncio.CancelledError:
+            return self.create_error_response("Client disconnected")
+
+        if final_res is None:
+            return self.create_error_response(
+                "No output received from the engine.",
+                err_type="InternalServerError",
+                status_code=HTTPStatus.INTERNAL_SERVER_ERROR,
+            )
+
+        choices: list[ChatCompletionResponseChoice] = []
+        total_reasoning_tokens = 0
+
+        role = self.get_chat_request_role(request)
+        tool_parser_cls = (
+            self.parser_cls.tool_parser_cls if self.parser_cls is not None else None
+        )
+        for output in final_res.outputs:
+            # check for error finish reason and raise GenerationError
+            # finish_reason='error' indicates a retryable request-level internal error
+            self._raise_if_error(output.finish_reason, request_id)
+            token_ids = output.token_ids
+            out_logprobs = output.logprobs
+
+            if request.logprobs and (
+                request.top_logprobs is not None or request.logprob_token_ids
+            ):
+                assert out_logprobs is not None, "Did not output logprobs"
+                logprobs = self._create_chat_logprobs(
+                    token_ids=token_ids,
+                    top_logprobs=out_logprobs,
+                    num_output_top_logprobs=request.top_logprobs,
+                    logprob_token_ids=request.logprob_token_ids,
+                    tokenizer=tokenizer,
+                    return_as_token_id=request.return_tokens_as_token_ids,
+                )
+            else:
+                logprobs = None
+
+            if parser is not None:
+                reasoning, content, tool_calls = parser.parse(
+                    output.text,
+                    request,
+                    enable_auto_tools=self.enable_auto_tools,
+                    model_output_token_ids=token_ids,
+                )
+                suppress_metadata = not request.include_reasoning and parser is not None
+                if not request.include_reasoning:
+                    reasoning = None
+                total_reasoning_tokens += parser.count_reasoning_tokens(token_ids)
+                if suppress_metadata:
+                    logprobs = None
+            else:
+                reasoning = None
+                content = output.text
+                tool_calls = []
+                suppress_metadata = False
+
+            auto_tools_called = False
+            is_named_tool_choice = (
+                request.tool_choice is not None
+                and type(request.tool_choice) is ChatCompletionNamedToolChoiceParam
+            )
+            is_required_tool_choice = request.tool_choice == "required"
+
+            # All six construction sites route through ``self._create_chat_message``
+            # so subclasses can swap in a specialized :class:`ChatMessage`
+            # (e.g. the Cohere v2 handler's ``CohereChatMessage``) without
+            # having to duplicate this branch logic.
+            if (not self.enable_auto_tools or not tool_parser_cls) and (
+                not is_named_tool_choice and not is_required_tool_choice
+            ):
+                message = self._create_chat_message(
+                    role=role, reasoning=reasoning, content=content
+                )
+
+            elif is_named_tool_choice or is_required_tool_choice:
+                message = self._create_chat_message(
+                    role=role,
+                    reasoning=reasoning,
+                    content=content or "",
+                    tool_calls=[
+                        ToolCall(id=tc.id or make_tool_call_id(), function=tc)
+                        for tc in (tool_calls or [])
+                    ],
+                )
+
+            # if the request doesn't use tool choice
+            # OR specifies to not use a tool
+            elif not request.tool_choice or request.tool_choice == "none":
+                message = self._create_chat_message(
+                    role=role, reasoning=reasoning, content=content
+                )
+
+            # handle when there are tools and tool choice is auto
+            elif (
+                request.tools
+                and (request.tool_choice == "auto" or request.tool_choice is None)
+                and self.enable_auto_tools
+                and tool_parser_cls
+            ):
+                auto_tools_called = tool_calls is not None and len(tool_calls) > 0
+                if tool_calls:
+                    message = self._create_chat_message(
+                        role=role,
+                        reasoning=reasoning,
+                        content=content,
+                        tool_calls=[
+                            ToolCall(id=tc.id or make_tool_call_id(), function=tc)
+                            for tc in tool_calls
+                        ],
+                    )
+
+                else:
+                    message = self._create_chat_message(
+                        role=role,
+                        reasoning=reasoning,
+                        content=content,
+                    )
+
+            # undetermined case that is still important to handle
+            else:
+                logger.error(
+                    "Error in chat_completion_full_generator - cannot determine"
+                    " if tools should be extracted. Returning a standard chat "
+                    "completion."
+                )
+                message = self._create_chat_message(
+                    role=role, reasoning=reasoning, content=content
+                )
+
+            # Subclass hook: enrich the constructed message with any
+            # parser-side extras that don't fit through the plain
+            # ``(reasoning, content, tool_calls)`` tuple. Base is a no-op;
+            # citation-aware handlers use this to surface grounding
+            # metadata cached on the reasoning parser.
+            message = self._finalize_response_message(message, parser=parser)
+
+            # In OpenAI's API, when a tool is called, the finish_reason is:
+            # "tool_calls" for "auto" or "required" tool calls,
+            # and "stop" for named tool calls.
+            is_finish_reason_tool_calls = (
+                auto_tools_called and output.finish_reason == "stop"
+            ) or (
+                request.tool_choice
+                and request.tool_choice == "required"
+                and output.finish_reason == "stop"
+            )
+
+            routed_experts_b64 = (
+                numpy2base64(output.routed_experts)
+                if output.routed_experts is not None
+                else None
+            )
+
+            choice_data = ChatCompletionResponseChoice(
+                index=output.index,
+                message=message,
+                logprobs=logprobs,
+                finish_reason="tool_calls"
+                if is_finish_reason_tool_calls
+                else output.finish_reason
+                if output.finish_reason
+                else "stop",
+                stop_reason=output.stop_reason,
+                token_ids=(
+                    as_list(output.token_ids)
+                    if request.return_token_ids and not suppress_metadata
+                    else None
+                ),
+                routed_experts=routed_experts_b64,
+            )
+            choice_data = maybe_filter_parallel_tool_calls(choice_data, request)
+
+            choices.append(choice_data)
+
+        if request.echo:
+            last_msg_content: str | list[dict[str, str]] = ""
+            if (
+                conversation
+                and "content" in conversation[-1]
+                and conversation[-1].get("role") == role
+            ):
+                last_msg_content = conversation[-1]["content"] or ""
+            if isinstance(last_msg_content, list):
+                last_msg_content = "\n".join(msg["text"] for msg in last_msg_content)
+
+            for choice in choices:
+                full_message = last_msg_content + (choice.message.content or "")
+                choice.message.content = full_message
+
+        assert final_res.prompt_token_ids is not None
+        num_prompt_tokens = len(final_res.prompt_token_ids)
+        if final_res.encoder_prompt_token_ids is not None:
+            num_prompt_tokens += len(final_res.encoder_prompt_token_ids)
+        num_generated_tokens = sum(
+            len(output.token_ids) for output in final_res.outputs
+        )
+        usage = UsageInfo(
+            prompt_tokens=num_prompt_tokens,
+            completion_tokens=num_generated_tokens,
+            total_tokens=num_prompt_tokens + num_generated_tokens,
+            completion_tokens_details=_make_completion_tokens_details(
+                total_reasoning_tokens
+            )
+            if self._include_reasoning_tokens_details
+            else None,
+        )
+        usage.prompt_tokens_details = _make_prompt_tokens_details(
+            self.enable_prompt_tokens_details,
+            final_res.num_cached_tokens,
+            final_res.num_cache_creation_tokens,
+            mm_token_counts,
+        )
+
+        request_metadata.final_usage_info = usage
+
+        per_request_metrics: PerRequestMetrics | None = None
+        # Per-request metrics (timing + spec-decode acceptance) describe a single
+        # generation stream. For n>1 the stats belong to only one of the n
+        # sequences, so they cannot be attributed to the request; suppress.
+        if (request.n or 1) == 1:
+            if self.enable_per_request_metrics:
+                per_request_metrics = build_per_request_timing_metrics(
+                    final_res.metrics, num_generated_tokens
+                )
+            spec_stats = build_spec_decoding_metrics(final_res)
+            if spec_stats is not None:
+                if per_request_metrics is None:
+                    per_request_metrics = PerRequestMetrics()
+                per_request_metrics.speculative_decoding = spec_stats
+
+        # ``final_res.prompt`` is the rendered chat-templated prompt text
+        prompt_text = final_res.prompt if request.return_prompt_text else None
+
+        response = ChatCompletionResponse(
+            id=request_id,
+            created=created_time,
+            model=model_name,
+            choices=choices,
+            usage=usage,
+            system_fingerprint=self.system_fingerprint,
+            prompt_logprobs=clamp_prompt_logprobs(final_res.prompt_logprobs),
+            prompt_token_ids=(
+                final_res.prompt_token_ids if request.return_token_ids else None
+            ),
+            prompt_text=prompt_text,
+            kv_transfer_params=final_res.kv_transfer_params,
+            ec_transfer_params=final_res.ec_transfer_params,
+            metrics=per_request_metrics,
+        )
+
+        # Log complete response if output logging is enabled
+        if self.enable_log_outputs and self.request_logger:
+            for choice in choices:
+                output_text = ""
+                if choice.message.content:
+                    output_text = choice.message.content
+                elif choice.message.tool_calls:
+                    # For tool calls, log the function name and arguments
+                    tool_call_descriptions = []
+                    for tc in choice.message.tool_calls:  # type: ignore
+                        function_call: FunctionCall = tc.function  # type: ignore
+                        tool_call_descriptions.append(
+                            f"{function_call.name}({function_call.arguments})"
+                        )
+                    tool_calls_str = ", ".join(tool_call_descriptions)
+                    output_text = f"[tool_calls: {tool_calls_str}]"
+
+                if output_text:
+                    # Get the corresponding output token IDs
+                    output_token_ids = None
+                    if choice.index < len(final_res.outputs):
+                        output_token_ids = final_res.outputs[choice.index].token_ids
+
+                    self.request_logger.log_outputs(
+                        request_id=request_id,
+                        outputs=output_text,
+                        output_token_ids=output_token_ids,
+                        finish_reason=choice.finish_reason,
+                        is_streaming=False,
+                        delta=False,
+                    )
+
+        return response
+
+    def _get_top_logprobs(
+        self,
+        logprobs: dict[int, Logprob],
+        top_logprobs: int | None,
+        tokenizer: TokenizerLike | None,
+        should_return_as_token_id: bool,
+        return_all: bool = False,
+    ) -> list[ChatCompletionLogProb]:
+        return [
+            ChatCompletionLogProb(
+                token=(
+                    token := self._get_decoded_token(
+                        p[1],
+                        p[0],
+                        tokenizer,
+                        return_as_token_id=should_return_as_token_id,
+                    )
+                ),
+                logprob=max(p[1].logprob, -9999.0),
+                bytes=list(token.encode("utf-8", errors="replace")),
+            )
+            for i, p in enumerate(logprobs.items())
+            if return_all
+            or top_logprobs == -1
+            or (top_logprobs is not None and i < top_logprobs)
+        ]
+
+    def _create_chat_logprobs(
+        self,
+        token_ids: GenericSequence[int],
+        top_logprobs: GenericSequence[dict[int, Logprob] | None],
+        tokenizer: TokenizerLike | None,
+        num_output_top_logprobs: int | None = None,
+        logprob_token_ids: list[int] | None = None,
+        return_as_token_id: bool | None = None,
+    ) -> ChatCompletionLogProbs:
+        """Create OpenAI-style logprobs."""
+        logprobs_content: list[ChatCompletionLogProbsContent] = []
+
+        should_return_as_token_id = (
+            return_as_token_id
+            if return_as_token_id is not None
+            else self.return_tokens_as_token_ids
+        )
+        for i, token_id in enumerate(token_ids):
+            step_top_logprobs = top_logprobs[i]
+            if step_top_logprobs is None or step_top_logprobs.get(token_id) is None:
+                if should_return_as_token_id:
+                    token = format_token_id_placeholder(token_id)
+                else:
+                    if tokenizer is None:
+                        raise ValueError(
+                            "Unable to get tokenizer because `skip_tokenizer_init=True`"
+                        )
+
+                    token = tokenizer.decode(token_id)
+
+                logprobs_content.append(
+                    ChatCompletionLogProbsContent(
+                        token=token,
+                        bytes=list(token.encode("utf-8", errors="replace")),
+                    )
+                )
+            else:
+                step_token = step_top_logprobs[token_id]
+                step_decoded = step_token.decoded_token
+
+                logprobs_content.append(
+                    ChatCompletionLogProbsContent(
+                        token=self._get_decoded_token(
+                            step_token,
+                            token_id,
+                            tokenizer,
+                            should_return_as_token_id,
+                        ),
+                        logprob=max(step_token.logprob, -9999.0),
+                        bytes=(
+                            None
+                            if step_decoded is None
+                            else list(step_decoded.encode("utf-8", errors="replace"))
+                        ),
+                        top_logprobs=self._get_top_logprobs(
+                            step_top_logprobs,
+                            num_output_top_logprobs,
+                            tokenizer,
+                            should_return_as_token_id,
+                            return_all=bool(logprob_token_ids),
+                        ),
+                    )
+                )
+
+        return ChatCompletionLogProbs(content=logprobs_content)

bundle/plugin-site/ornith_g256/_vllm_correctness/manifest.json

@@ -1,5 +1,5 @@
 {
-  "runtime_version": "1.0.1",
+  "runtime_version": "1.0.3",
   "modules": {
     "vllm.model_executor.layers.mamba.gdn.qwen_gdn_linear_attn": {
       "file": "qwen_gdn_linear_attn.py",
@@ -24,6 +24,12 @@
       "native_path": "v1/worker/gpu_model_runner.py",
       "native_sha256": "4706fcf4b85158173d88e3b19ffb0f6a280fcb60554cfdac861a01a0cab4d4df",
       "sha256": "ef0009acaa67a1a5e066984f34bf619b06e85220ff953a3d9511376f0146dccf"
+    },
+    "vllm.entrypoints.openai.chat_completion.serving": {
+      "file": "chat_completion_serving_r03.py",
+      "native_path": "entrypoints/openai/chat_completion/serving.py",
+      "native_sha256": "9982953285e9df469032a82fffa4095d0e9d86278bede6e2b91d03d02373d182",
+      "sha256": "5720371f5bcb351e340a3e8d9a0235ec01ff27ed93e4f35542e31b83a055ef45"
     }
   },
   "upstream": [

bundle/plugin-site/ornith_g256/_xgrammar_native/build-receipt.json

@@ -0,0 +1,29 @@
+{
+  "mode": "candidate",
+  "library": "libxgrammar-candidate.so",
+  "sha256": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7",
+  "source_commit": "557becfb64c503ae9c04344b0047661f43f44320",
+  "patch_sha256": "2c42719e948809242f230ec534f967eab1babad9be1ba6dda9c59fd706f20d45",
+  "sources": {
+    "cpp/grammar.cc": "01d1552b37ecaa72da601a3b6d43975a2c3a0de02cdf797e82c243b30eef8c9d",
+    "cpp/grammar_functor.cc": "9410c20b665093b9fb67dc30db148121b819d3aee5ff94eb82558bf79b1f8279"
+  },
+  "configure": [
+    "cmake",
+    "-S",
+    "/srv/ssd/p3700ba/scratch/crown/runtime-tool-integration-20260917/native-build-r2/source-candidate",
+    "-B",
+    "/srv/ssd/p3700ba/scratch/crown/runtime-tool-integration-20260917/native-build-r2/build-candidate",
+    "-G",
+    "Ninja",
+    "-DCMAKE_BUILD_TYPE=Release",
+    "-DXGRAMMAR_BUILD_PYTHON_BINDINGS=ON",
+    "-DCMAKE_CXX_FLAGS_RELEASE=-O3 -DNDEBUG -fno-lto",
+    "-DCMAKE_CXX_FLAGS_RELWITHDEBINFO=-O3 -g -DNDEBUG -fno-lto",
+    "-DPython_EXECUTABLE=/opt/ciru/glm53-iu4/venv/bin/python",
+    "-DPython3_EXECUTABLE=/opt/ciru/glm53-iu4/venv/bin/python"
+  ],
+  "build_seconds": 67.22386657004245,
+  "compiler": "g++ (GCC) 15.2.0\nCopyright (C) 2025 Free Software Foundation, Inc.\nThis is free software; see the source for copying conditions.  There is NO\nwarranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.\n\n",
+  "submodules": " 6689d14c203eed390ae7bb64f56a983cfd7dff9c 3rdparty/cpptrace (v0.1-644-g6689d14)\n bbd2f4d32427e548797929af08cfe2a9cbb3cf12 3rdparty/dlpack (v1.0)\n df1544bcee0c7ce35cd5ea0b3eb8cc81855a4140 3rdparty/googletest (release-1.8.0-3498-gdf1544bc)\n"
+}

bundle/plugin-site/ornith_g256/_xgrammar_native/libxgrammar-candidate.so

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

bundle/plugin-site/ornith_g256/_xgrammar_native/selection.json

@@ -0,0 +1,4 @@
+{
+  "mode": "candidate",
+  "sha256": "92ab36bfae20c6c4df3f5321ca375204799ac833252921d9d8f6a206491d35f7"
+}

bundle/plugin-site/ornith_g256/attention_compact.py

@@ -107,9 +107,9 @@ def configure(*, max_num_seqs, max_model_len, max_num_batched_tokens, device,
             raise RuntimeError('Compact attention arena was configured differently')
         return _arena['bytes']
     if (not 1 <= max_num_seqs <= 8 or not 1 <= max_model_len <= 262144
-            or not 1 <= max_num_batched_tokens <= 2048
+            or not 1 <= max_num_batched_tokens <= 8192
             or torch.cuda.is_current_stream_capturing()):
-        raise ValueError('Configure target compact prefill before capture, within C8/256K/2048 tokens')
+        raise ValueError('Configure target compact prefill before capture, within C8/256K/8192 tokens')
     _prefill = _graph_safe_prefill()
     key = torch.empty((max_num_seqs * max_model_len, 2, 256),
                       dtype=torch.bfloat16, device=device)
@@ -177,14 +177,14 @@ def try_forward(query, key_cache, value_cache, output, block_table,
         num_warps=8, num_stages=1)
     # Graph-padded rows are not real requests and receive defined zero output.
     output.zero_()
-    if (_arena['iu4'] is not None and requests == 1 and max_query_len == 1120
+    if (_arena['iu4'] is not None and requests == 1 and max_query_len in (1120, 4480)
             and max_seq_len >= 4096 and max_seq_len % 32 == 0
             and max_query_len <= rows and sm_scale == 256**-.5):
         _arena['iu4'].forward(query, key, value, output, query_start_loc,
-                              seq_lens, max_seq_len)
+                              seq_lens, max_seq_len, query_rows=max_query_len)
         logger.info_once('Ornith C1 IU4 prefill active: BF16 cache gather + '
                          'normalized Q/K H256 and P/V H32, signed IU4 QK/PV; '
-                         'Q1120 and aligned K>=4096, existing BF16 fallback elsewhere')
+                         'Q%d and aligned K>=4096, existing BF16 fallback elsewhere', max_query_len)
         return True
     _prefill(q=query, k=key, v=value, o=output, softmax_lse=lse,
         sd_mask=None, sm_scale=sm_scale, alibi_slopes=None, causal=True,

bundle/plugin-site/ornith_g256/attention_iu4.py

@@ -11,31 +11,42 @@ class NativeAttention:
             raise ValueError("IU4 prefill capacity must be1120..262144")
         self.max_keys = 65536 if max_model_len <= 65536 else 262144
         groups = self.max_keys // 32
+        self.query_capacity = 4480
         self.library = ctypes.CDLL(str(library_path))
         self.library.iu4_prepare.argtypes = [ctypes.c_void_p] * 9 + [ctypes.c_int] * 3 + [ctypes.c_void_p]
         self.library.iu4_prepare.restype = ctypes.c_int
         self.library.iu4_attention.argtypes = [ctypes.c_void_p] * 7 + [ctypes.c_int] + [ctypes.c_void_p]
         self.library.iu4_attention.restype = ctypes.c_int
-        specs = [((16, 1120, 32), torch.int32), ((16, 1120), torch.float16),
+        self.library.iu4_prepare_q4480.argtypes = self.library.iu4_prepare.argtypes
+        self.library.iu4_prepare_q4480.restype = ctypes.c_int
+        self.library.iu4_attention_q4480.argtypes = self.library.iu4_attention.argtypes
+        self.library.iu4_attention_q4480.restype = ctypes.c_int
+        specs = [((16, self.query_capacity, 32), torch.int32), ((16, self.query_capacity), torch.float16),
                  ((2, groups, 32, 32), torch.int32), ((2, self.max_keys), torch.float16),
                  ((2, groups, 4, 256), torch.int32), ((2, groups, 256), torch.float16)]
         self.packed = [torch.empty(shape, dtype=dtype, device=device) for shape, dtype in specs]
         self.pointers = [tensor.data_ptr() for tensor in self.packed]
-        self.output = torch.empty((1120, 16, 256), dtype=torch.bfloat16, device=device)
+        self.output = torch.empty((self.query_capacity, 16, 256), dtype=torch.bfloat16, device=device)
         self.bytes = self.output.numel() * self.output.element_size() + sum(tensor.numel() * tensor.element_size() for tensor in self.packed)
 
-    def forward(self, query, key, value, output, starts, lengths, max_keys):
-        # C1 max_query_len1120 and exact CPU metadata max_seq_len are required.
-        # The unchanged consumer uses its original contiguous fixed1120 output.
+    def forward(self, query, key, value, output, starts, lengths, max_keys, *, query_rows=1120):
+        # CPU query metadata selects one of two fixed-Q consumers. The shared
+        # arena is capacity-sized; each native call packs with its own Q stride.
         # No allocation, tensor copy to CPU, synchronization or cache mutation here.
-        if not 1120 <= max_keys <= self.max_keys or max_keys % 32:
+        if (query_rows not in (1120, 4480) or query.shape[0] < query_rows
+                or output.shape[0] < query_rows
+                or not query_rows <= max_keys <= self.max_keys or max_keys % 32):
             raise ValueError("IU4 prefill requires aligned keys within reserved capacity")
+        prepare = (self.library.iu4_prepare_q4480 if query_rows == 4480
+                   else self.library.iu4_prepare)
+        attend = (self.library.iu4_attention_q4480 if query_rows == 4480
+                  else self.library.iu4_attention)
         stream = torch.cuda.current_stream().cuda_stream
-        rc = self.library.iu4_prepare(query.data_ptr(), key.data_ptr(), value.data_ptr(),
+        rc = prepare(query.data_ptr(), key.data_ptr(), value.data_ptr(),
             *self.pointers, max_keys, query.stride(0), query.stride(1), stream)
         if rc:
             raise RuntimeError(f'IU4 attention preparation launch failed: HIP {rc}')
-        rc = self.library.iu4_attention(*self.pointers, self.output.data_ptr(), max_keys, stream)
+        rc = attend(*self.pointers, self.output.data_ptr(), max_keys, stream)
         if rc:
             raise RuntimeError(f'IU4 attention launch failed: HIP {rc}')
-        output[:1120].copy_(self.output)
+        output[:query_rows].copy_(self.output[:query_rows])

bundle/plugin-site/ornith_g256/attention_iu4_persistent.py

@@ -99,7 +99,7 @@ class State:
             raise RuntimeError(f'Expected ten target layers aliasing five banks, got {targets}/{len(self.banks)}')
         self.device = next(iter(self.banks.values())).cache.device
         device = self.device
-        self.groups = torch.empty(2048, dtype=torch.int32, device=device)
+        self.groups = torch.empty(8192, dtype=torch.int32, device=device)
         self.contexts = torch.empty(8, dtype=torch.int32, device=device)
         self.counts = torch.empty(8, dtype=torch.int32, device=device)
         self.owners = torch.empty(64, dtype=torch.int32, device=device)
@@ -122,7 +122,7 @@ class State:
         bank = self.banks.get(cache.data_ptr())
         if bank is None:
             return
-        if slots.numel() > 2048 or slots.dtype != torch.int64:
+        if slots.numel() > 8192 or slots.dtype != torch.int64:
             raise RuntimeError('Unexpected persistent IU4 slot capacity/type')
         self.call('update_cache', [cache, slots, self.groups, bank.busy,
                                    bank.valid, *bank.packed], slots.numel())

bundle/plugin-site/ornith_g256/optimized/__init__.py

@@ -0,0 +1 @@
+"""Retained Ornith runtime optimizations."""

bundle/plugin-site/ornith_g256/optimized/compile_scope.py

@@ -0,0 +1,30 @@
+"""Keep model-only fusion hooks outside compiler autotuning microbenchmarks."""
+import json
+import threading
+from pathlib import Path
+
+def install(original_run, output):
+    from torch._inductor.runtime.triton_heuristics import CachingAutotuner
+    from torch._inductor.codegen.wrapper import PythonWrapperCodegen
+    selected_run=CachingAutotuner.run
+    original_tuning=PythonWrapperCodegen.generate_and_run_autotune_block
+    local=threading.local()
+    state=dict(autotune_blocks=0,original_autotune_calls=0,active=False)
+    out=Path(output)/'compile-scope.json'
+    def persist():out.write_text(json.dumps(state,indent=2)+'\n')
+    def tuning(wrapper,*args,**kwargs):
+        previous=getattr(local,'autotuning',False)
+        local.autotuning=True;state['active']=True
+        state['autotune_blocks']+=1
+        try:return original_tuning(wrapper,*args,**kwargs)
+        finally:
+            local.autotuning=previous;state['active']=previous;persist()
+    def run(kernel,*args,**kwargs):
+        if getattr(local,'autotuning',False):
+            state['original_autotune_calls']+=1
+            return original_run(kernel,*args,**kwargs)
+        return selected_run(kernel,*args,**kwargs)
+    PythonWrapperCodegen.generate_and_run_autotune_block=tuning
+    CachingAutotuner.run=run
+    persist()
+    return lambda:dict(state)

bundle/plugin-site/ornith_g256/optimized/draft.py

@@ -0,0 +1,90 @@
+"""Isolated draft residual/RMSNorm repair; public target and sampler preserved."""
+import ast
+import hashlib
+import json
+import os
+from pathlib import Path
+import torch
+from ornith_g256.worker import OrnithG256Worker as Parent
+from .draft_kernel import launch
+
+R = Path(__file__).resolve().parents[1]
+NAME = 'triton_red_fused_add_arange_bitwise_and_constant_pad_nd_fused_add_rms_norm_ge_mul_select_slice_unsqueeze_view_3'
+KINDS = {
+    '3945f633918257a81244518749ca8c2613735dbedb656d49d811eba9e07a5f9a': 'middle',
+    '456e4167b1199b272395eae07c665ad6b501b39cb0018097a4abf87c72861bd7': 'final',
+}
+
+def install(output, cache_root):
+    out = Path(output)
+    state = dict(candidate_installed=True, kernels=[], calls={}, captured_rows={},
+                 graph_replay_poison_check=False, sampler_changed=False)
+    def persist():
+        temp = out / 'kernel-provenance.tmp'
+        temp.write_text(json.dumps(state, indent=2))
+        temp.replace(out / 'kernel-provenance.json')
+    # Compile all configured graph shapes before graph capture, without RNG use.
+    for rows in [1, 8, 16, 32, 48, 64]:
+        residual = torch.zeros((rows, 2048), device='cuda', dtype=torch.float32)
+        base = torch.zeros((2, 2, 2048), device='cuda', dtype=torch.bfloat16)
+        coef = torch.zeros((rows, 512), device='cuda', dtype=torch.bfloat16)
+        hidden = torch.zeros_like(residual, dtype=torch.bfloat16)
+        mask = torch.full((1,), 15, device='cuda', dtype=torch.int32)
+        weight = torch.ones((2048,), device='cuda', dtype=torch.bfloat16)
+        out_residual = torch.empty_like(hidden)
+        out_norm = torch.empty_like(hidden)
+        for kind in ['middle', 'final']:
+            inputs = [residual, base, coef, hidden, mask, weight]
+            if kind == 'middle':
+                inputs.append(out_residual)
+            inputs.extend([out_norm, rows, 2048])
+            launch(kind, inputs, torch.cuda.current_stream().cuda_stream)
+            torch.cuda.synchronize()
+            graph = torch.cuda.CUDAGraph()
+            with torch.cuda.graph(graph):
+                launch(kind, inputs, torch.cuda.current_stream().cuda_stream)
+            out_norm.fill_(float('nan'))
+            out_residual.fill_(float('nan'))
+            graph.replay()
+            torch.cuda.synchronize()
+            assert torch.count_nonzero(out_norm) == 0
+            if kind == 'middle':
+                assert torch.count_nonzero(out_residual) == 0
+            del graph
+    state['graph_replay_poison_check'] = True
+    from torch._inductor.runtime.triton_heuristics import CachingAutotuner
+    original = CachingAutotuner.run
+    classified, seen = {}, set()
+    draft_root = str(Path(cache_root) / 'vllm/torch_compile_cache/torch_aot_compile')
+
+    def run(kernel, *inputs, stream, benchmark_run=False, **kwargs):
+        if kernel not in classified:
+            kind = None
+            if kernel.fn.__name__ == NAME:
+                relative = Path(kernel.filename).relative_to(draft_root)
+                assert len(relative.parts[0]) == 64 and 'inductor_cache' in relative.parts, kernel.filename
+                fn = next(n for n in ast.parse(kernel.fn.src).body if isinstance(n, ast.FunctionDef))
+                fn.decorator_list = []
+                digest = hashlib.sha256(ast.dump(fn, include_attributes=False).encode()).hexdigest()
+                assert digest in KINDS, digest
+                kind = KINDS[digest]
+                state['kernels'].append(dict(kind=kind, function=NAME, filename=kernel.filename, ast_sha256=digest))
+            classified[kernel] = kind
+        kind = classified[kernel]
+        if kind is None:
+            return original(kernel, *inputs, stream=stream, benchmark_run=benchmark_run, **kwargs)
+        assert not benchmark_run and not kwargs
+        launch(kind, inputs, stream)
+        rows = int(inputs[-2])
+        captured = torch.cuda.is_current_stream_capturing()
+        key = kind + ':' + str(rows)
+        state['calls'][key] = state['calls'].get(key, 0) + 1
+        if captured and rows not in state['captured_rows'].setdefault(kind, []):
+            state['captured_rows'][kind].append(rows)
+        if (kind, rows, captured) not in seen:
+            seen.add((kind, rows, captured))
+            persist()
+        return None
+
+    CachingAutotuner.run = run
+    persist()

bundle/plugin-site/ornith_g256/optimized/draft_kernel.py

@@ -0,0 +1,50 @@
+"""Single-pass convolution + residual RMSNorm; no scratch reread."""
+import triton
+import triton.language as tl
+
+@triton.jit
+def fused_norm_register(residual, base, coefficients, hidden, query_mask, weight,
+                        out_residual, out_norm, rows: tl.constexpr,
+                        RETURN_RESIDUAL: tl.constexpr, XBLOCK: tl.constexpr = 2):
+    row = tl.program_id(0) * XBLOCK + tl.arange(0, XBLOCK)[:, None]
+    col = tl.arange(0, 2048)[None, :]
+    valid = row < rows
+    mask = tl.load(query_mask)
+    b0 = tl.load(base + 4096 + col).to(tl.float32)
+    b1 = tl.load(base + 6144 + col).to(tl.float32)
+    c0 = tl.load(coefficients + 256 + 512 * row + col // 16, valid, 0).to(tl.float32)
+    c1 = tl.load(coefficients + 384 + 512 * row + col // 16, valid, 0).to(tl.float32)
+    x = tl.load(hidden + 2048 * row + col, valid, 0).to(tl.float32)
+    previous = tl.load(hidden + 2048 * (row - 1) + col, valid & (row >= 1), 0).to(tl.float32)
+    old = tl.load(residual + 2048 * row + col, valid, 0)
+    w = tl.load(weight + col).to(tl.float32)
+    coefficient0 = b0 + c0
+    coefficient1 = b1 + c1
+    term0 = coefficient0 * x
+    term1 = coefficient1 * previous
+    term1 = term1 * ((row & mask) >= 1).to(tl.float32)
+    convolved = term0 + term1
+    updated = convolved + old
+    square_sum = tl.sum(updated * updated, axis=1)[:, None]
+    inv = tl.rsqrt(square_sum / 2048.0 + 1e-6)
+    normalized = (updated * inv) * w
+    tl.store(residual + 2048 * row + col, updated, valid)
+    if RETURN_RESIDUAL:
+        tl.store(out_residual + 2048 * row + col, updated, valid)
+    tl.store(out_norm + 2048 * row + col, normalized, valid)
+
+def launch(kind, inputs, stream):
+    import torch
+    assert stream == torch.cuda.current_stream().cuda_stream
+    if kind == 'middle':
+        residual, base, coef, hidden, mask, weight, out_residual, out_norm, rows, width = inputs
+    else:
+        residual, base, coef, hidden, mask, weight, out_norm, rows, width = inputs
+        out_residual = out_norm  # ignored by the final-layer specialization
+    assert width == 2048
+    assert residual.dtype == torch.float32 and tuple(residual.shape) == (rows, 2048)
+    assert all(t.is_contiguous() for t in [residual, base, coef, hidden, mask, weight, out_residual, out_norm])
+    assert tuple(coef.shape) == (rows, 512) and tuple(base.shape) == (2, 2, 2048)
+    fused_norm_register[(triton.cdiv(rows, 2),)](residual, base, coef, hidden, mask, weight,
+        out_residual, out_norm, rows, kind == 'middle', XBLOCK=2, num_warps=8,
+        num_stages=1, enable_fp_fusion=False)

bundle/plugin-site/ornith_g256/optimized/grammar_parallel.py

@@ -0,0 +1,122 @@
+"""Parallelize independent vLLM grammar chains without changing their semantics.
+
+Copyright 2026 Ciru. Licensed under Apache-2.0.
+The original runtime implements every mask, reasoning transition and rollback.
+Each worker receives a disjoint view of the original output buffer and exactly
+one request, so the original speculative sequence remains ordered per matcher.
+"""
+from concurrent.futures import ThreadPoolExecutor
+from copy import copy
+from dataclasses import dataclass
+import functools
+import hashlib
+import inspect
+import os
+from pathlib import Path
+import time
+
+BASELINE_SHA256 = '057f1ac57ec7d5f28b47a0ad85e03ae568951d408ceff345dbcac4c19b5d6ee9'
+MIN_SERIAL_COST_NS = 750_000
+
+
+@dataclass
+class MaskState:
+    cost_per_request_ns: float = 0.0
+    pool: ThreadPoolExecutor | None = None
+    serial_batches: int = 0
+    parallel_batches: int = 0
+
+
+def _fill_one(original, manager, requests, req_id, scheduled, start, stop):
+    # A shallow facade shares immutable configuration/backend and owns its
+    # output-buffer view. All request-local state belongs to this one request.
+    facade = copy(manager)
+    facade._grammar_bitmask = manager._grammar_bitmask[start:stop]
+    begin = time.perf_counter_ns()
+    original(facade, requests, [req_id], scheduled)
+    return time.perf_counter_ns() - begin
+
+
+def _available_workers(count):
+    cpus = len(os.sched_getaffinity(0)) if hasattr(os, 'sched_getaffinity') else (os.cpu_count() or 1)
+    return max(1, min(count, cpus // 2, 8))
+
+
+def make_wrapper(original, grammar_type=None):
+    if grammar_type is None:
+        from vllm.v1.structured_output.backend_xgrammar import XgrammarGrammar
+        grammar_type = XgrammarGrammar
+    @functools.wraps(original)
+    def grammar_bitmask(self, requests, structured_output_request_ids, scheduled_spec_decode_tokens):
+        manager = self
+        request_ids = structured_output_request_ids
+        scheduled = scheduled_spec_decode_tokens
+        count = len(request_ids)
+        # Keep C0/C1, diffusion and the existing large-batch optimization on
+        # their original paths. This patch fills the small-batch serving gap.
+        if count < 2 or count > 128 or manager.vllm_config.model_config.is_diffusion:
+            return original(manager, requests, request_ids, scheduled)
+
+        state = getattr(manager, '_ciru_mask_state', None)
+        if state is None:
+            state = manager._ciru_mask_state = MaskState()
+        grammars = [requests[r].structured_output_request.grammar for r in request_ids]
+        if not all(isinstance(g, grammar_type) for g in grammars):
+            return original(manager, requests, request_ids, scheduled)
+        unique = len({id(g) for g in grammars}) == count
+        thinking = not manager.enable_in_reasoning and all(
+            requests[r].structured_output_request.reasoning_ended is False for r in request_ids
+        )
+        parallel = (
+            manager._grammar_bitmask is not None
+            and unique
+            and not thinking
+            and state.cost_per_request_ns * count >= MIN_SERIAL_COST_NS
+            and _available_workers(count) > 1
+        )
+        if not parallel:
+            begin = time.perf_counter_ns()
+            result = original(manager, requests, request_ids, scheduled)
+            state.cost_per_request_ns = (time.perf_counter_ns() - begin) / count
+            state.serial_batches += 1
+            return result
+
+        if state.pool is None:
+            state.pool = ThreadPoolExecutor(max_workers=_available_workers(manager.vllm_config.scheduler_config.max_num_seqs), thread_name_prefix='ciru-grammar')
+        futures = []
+        offset = 0
+        for req_id in request_ids:
+            rows = len(scheduled.get(req_id, ())) + 1
+            futures.append(state.pool.submit(_fill_one, original, manager, requests, req_id, scheduled, offset, offset + rows))
+            offset += rows
+        elapsed = []
+        first_error = None
+        # Drain every future before returning or raising: no worker can keep
+        # writing into a buffer the scheduler might reuse for its next step.
+        for future in futures:
+            try:
+                elapsed.append(future.result())
+            except BaseException as error:
+                if first_error is None:
+                    first_error = error
+        state.parallel_batches += 1
+        if first_error is not None:
+            raise first_error
+        state.cost_per_request_ns = sum(elapsed) / count
+        return manager._grammar_bitmask[:offset].numpy()
+    return grammar_bitmask
+
+
+def install():
+    """Install once in the engine process, requiring the tested source identity."""
+    from vllm.v1.structured_output import StructuredOutputManager
+    if getattr(StructuredOutputManager, '_ciru_parallel_masks_installed', False):
+        return
+    source = Path(inspect.getsourcefile(StructuredOutputManager))
+    actual = hashlib.sha256(source.read_bytes()).hexdigest()
+    if actual != BASELINE_SHA256:
+        raise RuntimeError(f'Ciru parallel grammar requires its tested vLLM source; found {actual}')
+    StructuredOutputManager.grammar_bitmask = make_wrapper(StructuredOutputManager.grammar_bitmask)
+    StructuredOutputManager._ciru_parallel_masks_installed = True
+    from vllm.logger import init_logger
+    init_logger(__name__).info('Ciru parallel grammar masks R01 enabled; strict constraints and original per-request transitions retained')

bundle/plugin-site/ornith_g256/optimized/greedy.py

@@ -0,0 +1,40 @@
+"""Install the exact greedy walk beneath the unchanged public prefill wrapper."""
+import ast, copy, hashlib, inspect, json, os
+from pathlib import Path
+import torch
+from .greedy_kernel import greedy_path
+from ornith_g256.worker import OrnithG256Worker as Parent
+
+EXPECTED_SOURCE='56b817daecb3740cad899a3bc1ebc481def465457d5e1a7a516e0535522e9e7f'
+
+def install(out):
+    from ornith_g256 import prefill_draft
+    from vllm.v1.spec_decode.dflash import DFlashProposer
+    prefill_draft.install()
+    assert DFlashProposer._sample_draft_tokens is prefill_draft._sample_draft_tokens
+    original=prefill_draft._upstream_sample
+    source=Path(inspect.getsourcefile(original))
+    assert hashlib.sha256(source.read_bytes()).hexdigest()==EXPECTED_SOURCE
+    assert original.__module__=='vllm.v1.spec_decode.dflash'
+    tree=ast.parse(source.read_text())
+    cls=next(n for n in tree.body if isinstance(n,ast.ClassDef) and n.name=='DFlashProposer')
+    fn=copy.deepcopy(next(n for n in cls.body if isinstance(n,ast.FunctionDef) and n.name=='_sample_draft_tokens'))
+    fn.decorator_list=[]
+    index=next(i for i,n in enumerate(fn.body) if isinstance(n,ast.Assign)
+               and any(isinstance(t,ast.Name) and t.id=='rows' for t in n.targets))
+    fn.body.insert(index,ast.parse('if sampling_metadata.all_greedy:\n return qualified_path(scores, candidate_ids), None').body[0])
+    fn.name='_s13_sample'
+    out=Path(out)
+    globals_=dict(original.__globals__,qualified_path=greedy_path)
+    module=ast.fix_missing_locations(ast.Module(body=ast.parse('from __future__ import annotations').body+[fn],type_ignores=[]))
+    exec(compile(module,str(source)+':s13', 'exec'),globals_)
+    candidate=globals_['_s13_sample']
+    def dispatch(self,hidden,metadata):
+        if self.is_dflash2 and metadata.all_greedy:
+            return candidate(self,hidden,metadata)
+        return original(self,hidden,metadata)
+    prefill_draft._upstream_sample=dispatch
+    with (out/'sampler-installed.json').open('x') as f:
+        json.dump(dict(source_sha256=EXPECTED_SOURCE,prefill_wrapper_preserved=True,
+                       non_greedy_original=True,delta='Greedy early return after unchanged score calculation'),f,indent=2)
+

bundle/plugin-site/ornith_g256/optimized/greedy_kernel.py

@@ -0,0 +1,40 @@
+"""Exact greedy DFlash2 path walk; candidate logits and edge scores unchanged."""
+import torch
+import triton
+import triton.language as tl
+
+
+@triton.jit
+def _walk(S, I, O, S0: tl.constexpr, S1: tl.constexpr, S2: tl.constexpr,
+          S3: tl.constexpr, I0: tl.constexpr, I1: tl.constexpr, I2: tl.constexpr,
+          L: tl.constexpr, K: tl.constexpr):
+    row = tl.program_id(0)
+    col = tl.arange(0, K)
+    predecessor = tl.full((), 0, tl.int32)
+    for step in tl.static_range(L):
+        values = tl.load(S + row*S0 + step*S1 + predecessor*S2 + col*S3).to(tl.float32)
+        # Torch argmax selects the first NaN, or the first maximum otherwise.
+        nan_index = tl.min(tl.where(values != values, col, K), 0)
+        maximum = tl.max(tl.where(values != values, -float('inf'), values), 0)
+        max_index = tl.min(tl.where(values == maximum, col, K), 0)
+        selected = tl.where(nan_index < K, nan_index, max_index)
+        token = tl.load(I + row*I0 + step*I1 + selected*I2)
+        tl.store(O + row*L + step, token)
+        predecessor = selected
+
+
+def greedy_path(scores, candidate_ids):
+    b, length, k = candidate_ids.shape
+    if scores.shape != (b, length, k, k) or k != 16 or not 1 <= length <= 15:
+        raise ValueError('Unqualified DFlash2 score/ID shape')
+    if scores.dtype not in (torch.float32, torch.bfloat16, torch.float16):
+        raise ValueError('Unqualified score dtype')
+    if candidate_ids.dtype not in (torch.int32, torch.int64):
+        raise ValueError('Unqualified candidate-ID dtype')
+    if not scores.is_cuda or candidate_ids.device != scores.device:
+        raise ValueError('DFlash path requires matching GPU tensors')
+    output = torch.empty((b * length,), device=candidate_ids.device, dtype=candidate_ids.dtype)
+    if b:
+        _walk[(b,)](scores, candidate_ids, output, *scores.stride(),
+                    *candidate_ids.stride(), length, k, num_warps=1)
+    return output

bundle/plugin-site/ornith_g256/optimized/launch.py

@@ -0,0 +1,17 @@
+"""Published serving settings with the qualified wide prefill policy."""
+from ornith_g256 import launch
+
+original = launch.make_settings
+
+
+def settings(args):
+    result = original(args)
+    result.update(worker_cls='ornith_g256.optimized.worker.Worker',
+                  max_num_batched_tokens=8192, long_prefill_token_threshold=4480)
+    return result
+
+
+launch.make_settings = settings
+
+if __name__ == '__main__':
+    launch.main()

bundle/plugin-site/ornith_g256/optimized/mixed_qk_sum.py

@@ -0,0 +1,66 @@
+import triton
+import triton.language as tl
+
+@triton.jit
+def mixed_qk_sum(in_ptr0, out_ptr0, out_ptr1, xnumel_0, xnumel_1, XBLOCK: tl.constexpr, R0_BLOCK: tl.constexpr, K_XBLOCK: tl.constexpr, K_RBLOCK: tl.constexpr):
+    pid = tl.program_id(0)
+    num_xblocks_0 = tl.cdiv(xnumel_0, XBLOCK)
+    num_xblocks_1 = num_xblocks_0 + tl.cdiv(xnumel_1, K_XBLOCK)
+    if pid < num_xblocks_0:
+        pid_offset = pid
+        r0_numel = 256
+        rnumel = r0_numel
+        RBLOCK: tl.constexpr = R0_BLOCK
+        xoffset = pid_offset * XBLOCK
+        xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
+        xmask = xindex < xnumel_0
+        r0_base = tl.arange(0, R0_BLOCK)[None, :]
+        rbase = r0_base
+        x0 = xindex % 16
+        x1 = xindex // 16
+        _tmp4 = tl.full([XBLOCK, R0_BLOCK], 0, tl.float32)
+        x3 = xindex
+        for r0_offset in tl.range(0, r0_numel, R0_BLOCK, num_stages=2):
+            r0_index = r0_offset + r0_base
+            r0_mask = r0_index < r0_numel
+            roffset = r0_offset
+            rindex = r0_index
+            r0_2 = r0_index
+            tmp0 = tl.load(in_ptr0 + (r0_2 + 512 * x0 + 9216 * x1), r0_mask & xmask, eviction_policy='evict_first', other=0.0).to(tl.float32)
+            tmp1 = tmp0.to(tl.float32)
+            tmp2 = tmp1 * tmp1
+            tmp3 = tl.broadcast_to(tmp2, [XBLOCK, R0_BLOCK])
+            tmp5 = _tmp4 + tmp3
+            _tmp4 = tl.where(r0_mask & xmask, tmp5, _tmp4)
+        tmp4 = tl.sum(_tmp4, 1)[:, None]
+        tl.store(out_ptr0 + x3, tmp4, xmask)
+    elif pid < num_xblocks_1:
+        pid_offset = pid - num_xblocks_0
+        r0_numel = 256
+        rnumel = r0_numel
+        RBLOCK: tl.constexpr = K_RBLOCK
+        xoffset = pid_offset * K_XBLOCK
+        xindex = xoffset + tl.arange(0, K_XBLOCK)[:, None]
+        xmask = xindex < xnumel_1
+        r0_base = tl.arange(0, K_RBLOCK)[None, :]
+        rbase = r0_base
+        x4 = xindex % 2
+        x5 = xindex // 2
+        _tmp10 = tl.full([K_XBLOCK, K_RBLOCK], 0, tl.float32)
+        x7 = xindex
+        for r0_offset in tl.range(0, r0_numel, K_RBLOCK, num_stages=2):
+            r0_index = r0_offset + r0_base
+            r0_mask = r0_index < r0_numel
+            roffset = r0_offset
+            rindex = r0_index
+            r0_6 = r0_index
+            tmp6 = tl.load(in_ptr0 + (8192 + r0_6 + 256 * x4 + 9216 * x5), r0_mask & xmask, eviction_policy='evict_first', other=0.0).to(tl.float32)
+            tmp7 = tmp6.to(tl.float32)
+            tmp8 = tmp7 * tmp7
+            tmp9 = tl.broadcast_to(tmp8, [K_XBLOCK, K_RBLOCK])
+            tmp11 = _tmp10 + tmp9
+            _tmp10 = tl.where(r0_mask & xmask, tmp11, _tmp10)
+        tmp10 = tl.sum(_tmp10, 1)[:, None]
+        tl.store(out_ptr1 + x7, tmp10, xmask)
+    else:
+        pass

bundle/plugin-site/ornith_g256/optimized/pipeline/SOURCE-PROOF.json

@@ -0,0 +1,34 @@
+{
+  "chunk_scaled_dot_kkt.py": {
+    "parent_sha256": "7ae41df3755acc7e925a134d848f78bb5fc00b90069039d86403c5853323ea42",
+    "candidate_sha256": "71735d2b1c4a19eecbb6058708481817bb610426c48f96a5f646784afcf0c65c",
+    "reverse_exact_grid_program_id_edits_restores_identical_bytes": true,
+    "ast_equal_after_reverse": true
+  },
+  "solve_tril.py": {
+    "parent_sha256": "d1a4ff27623a938825052500afa7e2aa24fadd2fa4c9be6376c9128824f86f82",
+    "candidate_sha256": "bdd04b69c0f91562048d25834bd53d298dc1f1bd068b9494af88776b2bf4b1d1",
+    "reverse_exact_grid_program_id_edits_restores_identical_bytes": true,
+    "ast_equal_after_reverse": true
+  },
+  "wy_fast.py": {
+    "parent_sha256": "c9fc8259db0a8746dc04f1d799cb73333457118cf89a0bdce8dde8a80a28cd96",
+    "candidate_sha256": "4cc0be519bfcf3812d810a4bebb3dd45fb9b13a8d0ab62d0ad463a1a2e36f61f",
+    "reverse_exact_grid_program_id_edits_restores_identical_bytes": true,
+    "ast_equal_after_reverse": true
+  },
+  "chunk_o.py": {
+    "parent_sha256": "3b63b836a8cc561723c23db5a528a0b474d9cd78508d610f0b348c585a889a19",
+    "candidate_sha256": "a93f951048958265b4d1c3f9e67341329196e7a57525433934215fde4b16e61b",
+    "reverse_exact_grid_program_id_edits_restores_identical_bytes": true,
+    "ast_equal_after_reverse": true
+  },
+  "causal_conv1d.py": {
+    "parent_sha256": "230f3e554f5f1d609a37022578c84e42d5c71d46e684101278a7e21af4f7d418",
+    "candidate_sha256": "254c540f77b01694696a1ab535abd0e5a23d54e7211e202cf9beead04face16b",
+    "reverse_exact_grid_program_id_edits_restores_identical_bytes": true,
+    "ast_equal_after_reverse": true,
+    "preserves_HF_correctness_overlay": true,
+    "scope": "Existing guarded single-sequence non-APC scope only"
+  }
+}

bundle/plugin-site/ornith_g256/optimized/pipeline/__init__.py

@@ -0,0 +1 @@
+"""Retained Ornith runtime optimizations."""

bundle/plugin-site/ornith_g256/optimized/pipeline/adapter.py

@@ -0,0 +1,172 @@
+"""Isolated GDN spatial traversal adapter; leaves postconv/h/cumsum untouched.
+
+Install with enabled=False during normal model warmup. Candidate dispatch reuses
+only configurations already selected by the original Autotuner. A missing key
+fails before compilation/benchmarking of alternative configurations.
+"""
+from __future__ import annotations
+import copy
+import hashlib
+import importlib
+import importlib.util
+import inspect
+from pathlib import Path
+import sys
+import types
+import json
+
+class ParentConfigurationRequired(Exception):
+    pass
+
+ROOT = Path(__file__).resolve().parent
+OPS = 'vllm.third_party.flash_linear_attention.ops'
+SPECS = {
+    'chunk_scaled_dot_kkt_fwd': (OPS+'.chunk_scaled_dot_kkt','chunk_scaled_dot_kkt_fwd_kernel'),
+    'solve_tril': (OPS+'.solve_tril','merge_16x16_to_64x64_inverse_kernel'),
+    'recompute_w_u_fwd': (OPS+'.wy_fast','recompute_w_u_fwd_kernel'),
+    'chunk_fwd_o': (OPS+'.chunk_o','chunk_fwd_kernel_o'),
+    'causal_conv1d_fn': ('vllm.model_executor.layers.mamba.ops.causal_conv1d','_causal_conv1d_fwd_kernel'),
+}
+
+def exact_tuner_key(tuner, args, kwargs):
+    # Exact installed Autotuner.run key construction, including dtype order.
+    nargs = dict(zip(tuner.arg_names,args))
+    all_args = {**nargs,**kwargs}
+    selected = {k:v for k,v in all_args.items() if k in tuner.arg_names}
+    key = [selected[k] for k in tuner.keys if k in selected]
+    for value in selected.values():
+        if hasattr(value,'dtype'): key.append(str(value.dtype))
+    return tuple(key)
+
+def safe_conv_scope(arguments):
+    # Only one live sequence, no APC intermediate cache writes and no x/state alias.
+    if any(arguments.get(n) is not None for n in ('block_idx_first_scheduled_token','block_idx_last_scheduled_token','initial_state_idx','num_computed_tokens')):
+        return False
+    x = arguments['x']; starts=arguments['query_start_loc']; state=arguments['conv_states']
+    if tuple(starts.shape)!=(2,) or x.ndim!=2 or x.shape[1]<=8:
+        return False
+    if arguments.get('has_initial_state') is not None and tuple(arguments['has_initial_state'].shape)!=(1,):
+        return False
+    if x.untyped_storage().data_ptr()==state.untyped_storage().data_ptr():
+        return False
+    return True
+
+class SpatialPipeline:
+    def __init__(self, enabled):
+        from triton.runtime.autotuner import Autotuner
+        from triton.runtime.jit import JITFunction
+        self.enabled=enabled
+        self.patches=[]
+        self.original={}
+        self.candidate={}
+        self.tuners={}
+        self.calls={name:{'parent':0,'candidate':0,'scope_fallback':0} for name in SPECS}
+        self.activations=[]
+        self.configs=[]
+        self.seen=set()
+        proof=json.loads((ROOT/'SOURCE-PROOF.json').read_text())
+        for stage,(module_name,kernel_name) in SPECS.items():
+            parent=importlib.import_module(module_name)
+            filename=module_name.rsplit('.',1)[1]+'.py'
+            assert hashlib.sha256(Path(parent.__file__).read_bytes()).hexdigest()==proof[filename]['parent_sha256'],('Installed source changed',module_name)
+            path=ROOT/'candidate'/filename
+            assert hashlib.sha256(path.read_bytes()).hexdigest()==proof[filename]['candidate_sha256']
+            # Loading inside the real package preserves every unchanged relative import.
+            name=module_name.rsplit('.',1)[0]+'._gdn_spatial_'+filename[:-3]
+            spec=importlib.util.spec_from_file_location(name,path)
+            candidate=importlib.util.module_from_spec(spec)
+            sys.modules[name]=candidate
+            spec.loader.exec_module(candidate)
+            new_leaf=getattr(candidate,kernel_name)
+            while not isinstance(new_leaf,JITFunction): new_leaf=new_leaf.fn
+            old_kernel=getattr(parent,kernel_name)
+            def clone(node, *, stage=stage, new_leaf=new_leaf):
+                if isinstance(node,JITFunction):
+                    assert node.arg_names==new_leaf.arg_names
+                    return new_leaf
+                duplicate=copy.copy(node)
+                duplicate.fn=clone(node.fn)
+                if isinstance(node,Autotuner):
+                    assert duplicate.cache is node.cache and duplicate.configs is node.configs
+                    duplicate.base_fn=new_leaf.fn
+                    self.tuners[stage]=(node,duplicate)
+                    run_impl=type(node).run
+                    def guarded_run(this,*args,**kwargs):
+                        key=exact_tuner_key(this,args,kwargs)
+                        if len(this.configs)>1 and key not in this.cache:
+                            raise ParentConfigurationRequired(stage, key)
+                        config=this.cache[key] if len(this.configs)>1 else this.configs[0]
+                        record_key=(stage,repr(key),repr(config.all_kwargs()))
+                        if record_key not in self.seen:
+                            self.seen.add(record_key)
+                            self.configs.append(dict(stage=stage,key=repr(key),parent_selected_config=config.all_kwargs(),same_config_object=True))
+                        return run_impl(this,*args,**kwargs)
+                    duplicate.run=types.MethodType(guarded_run,duplicate)
+                return duplicate
+            candidate.__dict__[kernel_name]=clone(old_kernel)
+            self.original[stage]=getattr(parent,stage)
+            self.candidate[stage]=getattr(candidate,stage)
+        from vllm.third_party.flash_linear_attention.ops import chunk
+        from vllm.model_executor.layers.mamba.gdn import qwen_gdn_linear_attn as qwen
+        chunk_globals=chunk.chunk_gated_delta_rule_fwd.__globals__
+        core_globals=qwen.QwenGatedDeltaNetAttention._forward_core.__globals__
+        # Validate the whole hook set before the first global mutation.
+        for stage in SPECS:
+            namespace=core_globals if stage=='causal_conv1d_fn' else chunk_globals
+            if namespace.get(stage) is not self.original[stage]:
+                raise RuntimeError(f'Unexpected actual installed caller alias for {stage}')
+        for stage in SPECS:
+            namespace=core_globals if stage=='causal_conv1d_fn' else chunk_globals
+            original=self.original[stage]
+            signature=inspect.signature(original)
+            def dispatch(*args,_stage=stage,_signature=signature,**kwargs):
+                active=bool(self.enabled())
+                if active:
+                    bound=_signature.bind(*args,**kwargs);bound.apply_defaults()
+                    values=bound.arguments
+                    if _stage=='causal_conv1d_fn':
+                        safe=safe_conv_scope(values)
+                        rows=int(values['x'].shape[1])
+                    elif _stage=='solve_tril':
+                        rows=int(values['A'].shape[1])
+                        safe=values['A'].shape[-1]==64 and rows>8
+                    else:
+                        input_tensor=values['q'] if _stage=='chunk_fwd_o' else values['k']
+                        rows=int(input_tensor.shape[1]);safe=rows>8
+                    if not safe:
+                        self.calls[_stage]['scope_fallback']+=1
+                        active=False
+                    else:
+                        key=('activation',_stage,rows)
+                        if key not in self.seen:
+                            self.seen.add(key)
+                            self.activations.append(dict(stage=_stage,rows=rows,scope='single-sequence non-APC' if _stage=='causal_conv1d_fn' else 'independent chunk/head tiles'))
+                arm='candidate' if active else 'parent'
+                self.calls[_stage][arm]+=1
+                if not active:
+                    return self.original[_stage](*args, **kwargs)
+                try:
+                    return self.candidate[_stage](*args, **kwargs)
+                except ParentConfigurationRequired:
+                    self.calls[_stage]['candidate'] -= 1
+                    self.calls[_stage]['parent'] += 1
+                    self.calls[_stage]['tuner_cold'] = self.calls[_stage].get('tuner_cold', 0) + 1
+                    return self.original[_stage](*args, **kwargs)
+            dispatch.__name__=stage
+            dispatch.__signature__=signature
+            self.patches.append((namespace,stage,original))
+            namespace[stage]=dispatch
+
+    def snapshot(self):
+        return dict(calls=copy.deepcopy(self.calls),activations=copy.deepcopy(self.activations),
+                    parent_selected_configs=copy.deepcopy(self.configs),
+                    shared_autotuner_cache={name:parent.cache is candidate.cache for name,(parent,candidate) in self.tuners.items()},
+                    unchanged=['postconv','chunk_local_cumsum','chunk_gated_delta_rule_fwd_h','convolution update/decode','solve BT16/BT32'])
+
+    def restore(self):
+        for namespace,name,original in reversed(self.patches): namespace[name]=original
+        self.patches.clear()
+
+def install(enabled):
+    """Return a handle exposing snapshot()/restore(); enabled is a no-argument callable."""
+    return SpatialPipeline(enabled)

bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/causal_conv1d.py

@@ -0,0 +1,1307 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+# Copyright (c) 2024, Tri Dao.
+# Adapted from https://github.com/Dao-AILab/causal-conv1d/blob/main/causal_conv1d/causal_conv1d_interface.py
+
+
+import numpy as np
+import torch
+
+from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
+from vllm.v1.attention.backends.utils import NULL_BLOCK_ID, PAD_SLOT_ID
+
+
+@triton.jit(do_not_specialize_on_alignment=["num_cache_lines"])
+def _causal_conv1d_fwd_kernel(  # continuous batching
+    # Pointers to matrices
+    x_ptr,  # (dim, cu_seqlen) holding `batch` of actual sequences + padded sequences
+    w_ptr,  # (dim, width)
+    bias_ptr,
+    initial_states_ptr,  # conv_states_ptr
+    cache_indices_ptr,  # (batch, n_blocks + padding) The second dimension contains
+    # the block indices relevant for each sequence
+    # plus potential 0-padding at the beginning and at the end
+    has_initial_states_ptr,
+    query_start_loc_ptr,
+    batch_ptr,
+    token_chunk_offset_ptr,
+    block_idx_first_scheduled_token,  # (batch,)
+    block_idx_last_scheduled_token,  # (batch,)
+    initial_state_idx,  # (batch,)
+    num_computed_tokens,  # (batch,)
+    num_accepted_tokens_ptr,  # (batch,) or None
+    o_ptr,  # (dim, seqlen) - actually pointing to x_ptr
+    # Matrix dimensions
+    dim: tl.constexpr,
+    num_cache_lines,  # added to support vLLM larger cache lines
+    # Strides
+    stride_x_dim: tl.constexpr,  # stride to get to next feature-value,
+    stride_x_token: tl.int64,  # stride to get to next token (same feature-index, same sequence-index)
+    stride_w_dim: tl.constexpr,  # stride to get to next dim-axis value
+    stride_w_width: tl.constexpr,  # stride to get to next width-axis value
+    stride_istate_seq: tl.constexpr,
+    stride_istate_dim: tl.constexpr,
+    stride_istate_token: tl.constexpr,
+    stride_cache_indices: tl.constexpr,
+    stride_o_dim: tl.constexpr,
+    stride_o_token: tl.int64,
+    stride_block_m: tl.constexpr,  # Stride block to align divided by BLOCK_M
+    # others
+    pad_slot_id: tl.constexpr,
+    null_block_id: tl.constexpr,
+    # Meta-parameters
+    HAS_BIAS: tl.constexpr,
+    KERNEL_WIDTH: tl.constexpr,
+    SILU_ACTIVATION: tl.constexpr,
+    IS_APC_ENABLED: tl.constexpr,
+    IS_SPEC_DECODING: tl.constexpr,
+    HAS_NULL_BLOCK: tl.constexpr,
+    NP2_STATELEN: tl.constexpr,
+    BLOCK_M: tl.constexpr,
+    BLOCK_N: tl.constexpr,
+    launch_pdl: tl.constexpr,
+):
+    conv_states_ptr = initial_states_ptr
+    conv_state_indices_ptr = cache_indices_ptr
+    stride_conv_state_seq = stride_istate_seq
+    stride_conv_state_dim = stride_istate_dim
+    stride_conv_state_tok = stride_istate_token
+    state_len = (
+        KERNEL_WIDTH - 1
+    )  # can be passed via argument if it's not the same as this value
+
+    if launch_pdl:
+        tl.extra.cuda.gdc_wait()
+
+    # one program handles one chunk in a single sequence
+    # rather than mixing sequences - to make updating initial_states across sequences efficiently
+
+    # single-sequence id
+    idx_seq = tl.load(batch_ptr + tl.program_id(1)).to(tl.int64)
+
+    if IS_SPEC_DECODING:
+        conv_state_token_offset = (
+            tl.load(num_accepted_tokens_ptr + idx_seq).to(tl.int64) - 1
+        )
+    else:
+        conv_state_token_offset = 0
+    chunk_offset = tl.load(token_chunk_offset_ptr + tl.program_id(1))
+
+    # BLOCK_N elements along the feature-dimension (channel)
+    idx_feats = tl.program_id(0) * BLOCK_N + tl.arange(0, BLOCK_N)
+
+    if idx_seq == pad_slot_id:
+        if launch_pdl:
+            tl.extra.cuda.gdc_launch_dependents()
+        return
+
+    sequence_start_index = tl.load(query_start_loc_ptr + idx_seq)
+    sequence_end_index = tl.load(query_start_loc_ptr + idx_seq + 1)
+    # find the actual sequence length
+    seqlen = sequence_end_index - sequence_start_index
+
+    B_size: tl.constexpr = stride_block_m * BLOCK_M
+
+    if IS_APC_ENABLED:
+        # Handle the case if prefix caching is enabled.
+        # In particular, if prefix caching is enabled, the program write additional cache states to "cache_indices_ptr"
+
+        # Get the length of the completed sequence so far and compute the offset.
+        current_first_index = tl.load(block_idx_first_scheduled_token + idx_seq)
+        current_last_index = tl.load(block_idx_last_scheduled_token + idx_seq)
+        sequence_completed_index = tl.load(num_computed_tokens + idx_seq)
+
+        # Compute the offset where the first stride_block_m-aligned first full block is
+        # Value in "token-space"
+        sequence_completed_offset_token = sequence_completed_index % B_size
+        seq_completed_offset = B_size - sequence_completed_offset_token
+        seq_end_offset = (seqlen - seq_completed_offset) % B_size
+        last_full_block_token_index = sequence_end_index - seq_end_offset
+        # If the sequence without the sequence_offset_index is stride_cache_chunk-aligned, then the last full chunk is the second-to-last one
+        if seq_end_offset == 0:
+            last_full_block_token_index = last_full_block_token_index - B_size
+
+        # Get the number of blocks to be filled for the current sequence
+        # If n_block_to_fill = 0, then only the state at the sequence end is stored
+        n_block_to_fill = current_last_index - current_first_index
+
+        # Get the index of the init block
+        conv_state_init_index = tl.load(initial_state_idx + idx_seq)
+    else:
+        n_block_to_fill = 0
+        current_last_index = 0
+        conv_state_init_index = 0
+        current_first_index = 0
+        last_full_block_token_index = 0
+
+    token_offset = BLOCK_M * chunk_offset
+    segment_len = min(BLOCK_M, seqlen - token_offset)
+
+    # base of the sequence
+    x_base = (
+        x_ptr + sequence_start_index * stride_x_token + idx_feats * stride_x_dim
+    )  # [BLOCK_N,]
+
+    # cache_idx
+    conv_states_input_coord = tl.load(
+        conv_state_indices_ptr + idx_seq * stride_cache_indices + conv_state_init_index
+    ).to(tl.int64)
+
+    if HAS_NULL_BLOCK:  # noqa
+        if conv_states_input_coord == null_block_id:
+            # not processing as this is a null block (padding)
+            if launch_pdl:
+                tl.extra.cuda.gdc_launch_dependents()
+            return
+    conv_states_base = (
+        conv_states_ptr
+        + (conv_states_input_coord * stride_conv_state_seq)
+        + (idx_feats * stride_conv_state_dim)
+    )  # [BLOCK_N,]
+
+    w_base = w_ptr + (idx_feats * stride_w_dim)  # [BLOCK_N,]
+
+    # Does 2 things:
+    # 1. READ prior-block init-state data - [done by every Triton programs]
+    # 2. update conv_state with new data [only by the Triton program handles chunk_offset=0]
+    if chunk_offset == 0:
+        # read from conv_states
+        load_init_state = tl.load(has_initial_states_ptr + idx_seq).to(tl.int1)
+        if load_init_state:
+            # load from conv_states
+            prior_tokens = (
+                conv_states_base
+                + (state_len - 1 + conv_state_token_offset) * stride_conv_state_tok
+            )
+            mask_w = idx_feats < dim
+            if KERNEL_WIDTH == 2:
+                conv_states_ptrs = prior_tokens  # [BLOCK_N]
+                col0 = tl.load(conv_states_ptrs, mask_w, 0.0)
+            if KERNEL_WIDTH == 3:
+                conv_states_ptrs = prior_tokens  # [BLOCK_N]
+                col1 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 1 * stride_conv_state_tok  # [BLOCK_N]
+                col0 = tl.load(conv_states_ptrs, mask_w, 0.0)
+            if KERNEL_WIDTH == 4:
+                conv_states_ptrs = prior_tokens  # [BLOCK_N]
+                col2 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 1 * stride_conv_state_tok  # [BLOCK_N]
+                col1 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 2 * stride_conv_state_tok  # [BLOCK_N]
+                col0 = tl.load(conv_states_ptrs, mask_w, 0.0)
+            if KERNEL_WIDTH == 5:
+                conv_states_ptrs = prior_tokens  # [BLOCK_N]
+                col3 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 1 * stride_conv_state_tok  # [BLOCK_N]
+                col2 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 2 * stride_conv_state_tok  # [BLOCK_N]
+                col1 = tl.load(conv_states_ptrs, mask_w, 0.0)
+                conv_states_ptrs = prior_tokens - 3 * stride_conv_state_tok  # [BLOCK_N]
+                col0 = tl.load(conv_states_ptrs, mask_w, 0.0)
+        else:
+            # prior-tokens are zeros
+            if KERNEL_WIDTH >= 2:  # STRATEGY1
+                # first chunk and does not have prior-token, so just set to 0
+                col0 = tl.zeros((BLOCK_N,), dtype=x_ptr.dtype.element_ty)
+            if KERNEL_WIDTH >= 3:  # STRATEGY1
+                col1 = tl.zeros((BLOCK_N,), dtype=x_ptr.dtype.element_ty)
+            if KERNEL_WIDTH >= 4:  # STRATEGY1
+                col2 = tl.zeros((BLOCK_N,), dtype=x_ptr.dtype.element_ty)
+            if KERNEL_WIDTH >= 5:  # STRATEGY1
+                col3 = tl.zeros((BLOCK_N,), dtype=x_ptr.dtype.element_ty)
+
+        # STEP 2:
+        # here prepare data for updating conv_state
+        if (
+            state_len <= seqlen
+        ):  # SMALL_CACHE=True (only move part of 'x' into conv_state cache)
+            # just read from 'x'
+            # copy 'x' data to conv_state
+            # load only 'x' data (and set 0 before 'x' if seqlen < state_len)
+            idx_tokens_last = (seqlen - state_len) + tl.arange(
+                0, NP2_STATELEN
+            )  # [BLOCK_M]
+            x_ptrs = (
+                x_ptr
+                + ((sequence_start_index + idx_tokens_last) * stride_x_token)[:, None]
+                + (idx_feats * stride_x_dim)[None, :]
+            )  # [BLOCK_M,BLOCK_N,]
+            mask_x = (
+                (idx_tokens_last >= 0)[:, None]
+                & (idx_tokens_last < seqlen)[:, None]
+                & (idx_feats < dim)[None, :]
+            )  # token-index  # token-index  # feature-index
+            loaded_x = tl.load(x_ptrs, mask_x, 0.0)
+            idx_tokens_conv = tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+
+            # Compute the offset where the last block should be written in the conv_states
+            conv_states_output_coord = tl.load(
+                conv_state_indices_ptr
+                + idx_seq * stride_cache_indices
+                + current_last_index
+            ).to(tl.int64)
+
+            conv_states_ptrs_target = (
+                conv_states_ptr
+                + (conv_states_output_coord * stride_conv_state_seq)  # Offset from seq
+                + (idx_feats * stride_conv_state_dim)
+            )[None, :] + (  # [BLOCK_N,]
+                idx_tokens_conv * stride_conv_state_tok
+            )[:, None]
+
+            mask = (idx_tokens_conv < state_len)[:, None] & (idx_feats < dim)[None, :]
+            tl.debug_barrier()  #  NOTE: use this due to bug in Triton compiler
+            tl.store(conv_states_ptrs_target, loaded_x, mask)
+
+        else:
+            if load_init_state:
+                # update conv_state by shifting left, i.e. take last few cols from conv_state + cols from 'x'
+                idx_tokens_conv = tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+
+                conv_states_ptrs_source = (
+                    conv_states_ptr
+                    + (conv_states_input_coord * stride_conv_state_seq)
+                    + (idx_feats * stride_conv_state_dim)[None, :]
+                    + (
+                        (idx_tokens_conv + seqlen + conv_state_token_offset)
+                        * stride_conv_state_tok
+                    )[:, None]
+                )  # [BLOCK_M, BLOCK_N]
+                mask = (
+                    (conv_states_input_coord < num_cache_lines)
+                    & ((idx_tokens_conv + seqlen) < state_len)[:, None]
+                    & (idx_feats < dim)[None, :]
+                )
+                conv_state = tl.load(conv_states_ptrs_source, mask, other=0.0)
+
+                VAL = state_len - seqlen
+
+                x_ptrs = (
+                    x_base[None, :]
+                    + ((idx_tokens_conv - VAL) * stride_x_token)[:, None]
+                )  # [BLOCK_M, BLOCK_N]
+
+                mask_x = (
+                    (idx_tokens_conv - VAL >= 0)[:, None]
+                    & (idx_tokens_conv - VAL < seqlen)[:, None]
+                    & (idx_feats < dim)[None, :]
+                )  # token-index  # token-index  # feature-index
+                loaded_x = tl.load(x_ptrs, mask_x, 0.0)
+
+                tl.debug_barrier()  # need this due to the bug in tl.where not enforcing this when data is the result of another tl.load
+                new_conv_state = tl.where(
+                    mask, conv_state, loaded_x
+                )  # BUG in 'tl.where'  which requires a barrier before this
+                conv_states_ptrs_target = (
+                    conv_states_base
+                    + (idx_tokens_conv * stride_conv_state_tok)[:, None]
+                )  # [BLOCK_M, BLOCK_N]
+                mask = (idx_tokens_conv < state_len)[:, None] & (idx_feats < dim)[
+                    None, :
+                ]
+                tl.store(conv_states_ptrs_target, new_conv_state, mask)
+            else:  # load_init_state == False
+                # update conv_state by shifting left, BUT
+                # set cols prior to 'x' as zeros + cols from 'x'
+                idx_tokens_conv = tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+
+                VAL = state_len - seqlen
+
+                x_ptrs = (
+                    x_base[None, :]
+                    + ((idx_tokens_conv - VAL) * stride_x_token)[:, None]
+                )  # [BLOCK_M, BLOCK_N]
+
+                mask_x = (
+                    (idx_tokens_conv - VAL >= 0)[:, None]
+                    & (idx_tokens_conv - VAL < seqlen)[:, None]
+                    & (idx_feats < dim)[None, :]
+                )  # token-index  # token-index  # feature-index
+                new_conv_state = tl.load(x_ptrs, mask_x, 0.0)
+
+                conv_states_ptrs_target = (
+                    conv_states_base
+                    + (idx_tokens_conv * stride_conv_state_tok)[:, None]
+                )  # [BLOCK_M, BLOCK_N]
+                mask = (idx_tokens_conv < state_len)[:, None] & (idx_feats < dim)[
+                    None, :
+                ]
+                tl.store(conv_states_ptrs_target, new_conv_state, mask)
+
+    else:  # chunk_offset > 0
+        # read prior-token data from `x`
+        load_init_state = True
+        prior_tokens = x_base + (token_offset - 1) * stride_x_token
+        mask_w = idx_feats < dim
+        if KERNEL_WIDTH == 2:
+            conv_states_ptrs = prior_tokens  # [BLOCK_N]
+            col0 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+        if KERNEL_WIDTH == 3:
+            conv_states_ptrs = prior_tokens  # [BLOCK_N]
+            col1 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 1 * stride_x_token  # [BLOCK_N]
+            col0 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+        if KERNEL_WIDTH == 4:
+            conv_states_ptrs = prior_tokens  # [BLOCK_N]
+            col2 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 1 * stride_x_token  # [BLOCK_N]
+            col1 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 2 * stride_x_token  # [BLOCK_N]
+            col0 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+        if KERNEL_WIDTH == 5:
+            # ruff: noqa: F841
+            conv_states_ptrs = prior_tokens  # [BLOCK_N]
+            col3 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 1 * stride_x_token  # [BLOCK_N]
+            col2 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 2 * stride_x_token  # [BLOCK_N]
+            col1 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+            conv_states_ptrs = prior_tokens - 3 * stride_x_token  # [BLOCK_N]
+            col0 = tl.load(conv_states_ptrs, mask_w, 0.0, cache_modifier=".ca")
+
+        # Store intermediate states aligned with stride_block_m
+        # The additional states are cached starting from the last stride_block_m.
+        # For example:
+        # If n_block_to_fill = 0, then only the state at the sequence end is cached and the process below is not involved.
+        # If n_block_to_fill > 0, then the states at the sequence end and at the n_block_to_fill-last
+        # stride_block_m are cached.
+        # For example chunk_offset = n_block_to_fill stores the state at last_full_block
+        if (chunk_offset - 1) < n_block_to_fill:
+            # Store the states at the chunk boundaries from the start of the sequence
+            idx_tokens_last = (
+                last_full_block_token_index
+                - (n_block_to_fill - chunk_offset) * B_size
+                - state_len
+            ) + tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+            x_ptrs = (
+                x_ptr
+                + (idx_tokens_last * stride_x_token)[:, None]
+                + (idx_feats * stride_x_dim)[None, :]
+            )  # [BLOCK_M,BLOCK_N,]
+
+            mask_x = (idx_tokens_last >= 0)[:, None] & (idx_feats < dim)[
+                None, :
+            ]  # token-index  # token-index  # feature-index
+            loaded_x = tl.load(x_ptrs, mask_x, 0.0)
+            idx_tokens_conv = tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+
+            # cache_idx
+            conv_states_output_coord = tl.load(
+                conv_state_indices_ptr
+                + idx_seq * stride_cache_indices
+                + current_first_index
+                + (chunk_offset - 1)
+            ).to(tl.int64)
+
+            conv_states_ptrs_target = (
+                conv_states_ptr
+                + (conv_states_output_coord * stride_conv_state_seq)  # Offset from seq
+                + (idx_feats * stride_conv_state_dim)
+            )[None, :] + (  # [BLOCK_N,]
+                idx_tokens_conv * stride_conv_state_tok
+            )[:, None]
+
+            mask = (idx_tokens_conv < state_len)[:, None] & (idx_feats < dim)[None, :]
+            tl.debug_barrier()  #  NOTE: use this due to bug in Triton compiler
+            tl.store(conv_states_ptrs_target, loaded_x, mask)
+
+    if HAS_BIAS:
+        bias = bias_ptr + idx_feats
+        mask_bias = idx_feats < dim
+        acc_preload = tl.load(bias, mask=mask_bias, other=0.0).to(
+            tl.float32
+        )  # [BLOCK_N]
+    else:
+        acc_preload = tl.zeros((BLOCK_N,), dtype=tl.float32)
+
+    x_base_1d = x_base + token_offset * stride_x_token  # starting of chunk
+
+    # PRE-LOAD WEIGHTS
+    mask_w = idx_feats < dim
+    if KERNEL_WIDTH >= 2:
+        w_ptrs = w_base + (0 * stride_w_width)  # [BLOCK_N] tensor
+        w_col0 = tl.load(w_ptrs, mask_w, other=0.0)
+        w_ptrs = w_base + (1 * stride_w_width)  # [BLOCK_N] tensor
+        w_col1 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 3:
+        w_ptrs = w_base + (2 * stride_w_width)  # [BLOCK_N] tensor
+        w_col2 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 4:
+        w_ptrs = w_base + (3 * stride_w_width)  # [BLOCK_N] tensor
+        w_col3 = tl.load(w_ptrs, mask_w, other=0.0)
+    mask_x_1d = idx_feats < dim
+
+    if launch_pdl:
+        tl.extra.cuda.gdc_launch_dependents()
+
+    for idx_token in range(segment_len):
+        acc = acc_preload
+
+        matrix_w = w_col0
+        matrix_x = col0
+        for j in tl.static_range(KERNEL_WIDTH):
+            if KERNEL_WIDTH == 2:
+                if j == 1:  # KERNEL_WIDTH-1:
+                    matrix_w = w_col1
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 3:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 4:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    matrix_x = col2
+                elif j == 3:
+                    matrix_w = w_col3
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+
+            acc += matrix_x.to(tl.float32) * matrix_w.to(tl.float32)  # [BLOCK_N]
+
+        if KERNEL_WIDTH == 2:
+            col0 = matrix_x
+        elif KERNEL_WIDTH == 3:
+            col0 = col1
+            col1 = matrix_x
+        elif KERNEL_WIDTH == 4:
+            col0 = col1
+            col1 = col2
+            col2 = matrix_x
+
+        if SILU_ACTIVATION:
+            acc = acc / (1 + tl.exp(-acc))
+        mask_1d = (idx_token < segment_len) & (
+            idx_feats < dim
+        )  # token-index  # feature-index
+        o_ptrs = (
+            o_ptr
+            + (sequence_start_index + token_offset + idx_token) * stride_o_token
+            + (idx_feats * stride_o_dim)
+        )
+
+        tl.store(o_ptrs, acc, mask=mask_1d)
+
+
+def causal_conv1d_fn(
+    x: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor | None,
+    conv_states: torch.Tensor,
+    query_start_loc: torch.Tensor,
+    cache_indices: torch.Tensor | None = None,
+    has_initial_state: torch.Tensor | None = None,
+    activation: str | None = "silu",
+    pad_slot_id: int = PAD_SLOT_ID,
+    null_block_id: int = NULL_BLOCK_ID,
+    num_accepted_tokens: torch.Tensor | None = None,
+    block_idx_first_scheduled_token: torch.Tensor | None = None,
+    block_idx_last_scheduled_token: torch.Tensor | None = None,
+    initial_state_idx: torch.Tensor | None = None,
+    num_computed_tokens: torch.Tensor | None = None,
+    block_size_to_align=0,
+    metadata=None,
+    validate_data=False,
+):
+    """support varlen + continuous batching when x is 2D tensor
+
+    x: (dim,cu_seq_len)
+        cu_seq_len = total tokens of all seqs in that batch
+        sequences are concatenated from left to right for varlen
+    weight: (dim, width)
+    conv_states: (...,dim,width - 1) itype
+        updated inplace if cache_indices are not provided
+        [it use `cache_indices` to get the index to the cache of conv_state for that sequence
+
+        conv_state[cache_indices[i]] for seq-i - to be used as initial_state when has_initial_state[i] = True
+             and after that conv_state[cache_indices[i]] need to be shift-left and updated with values from 'x'
+        ]
+    query_start_loc: (batch + 1) int32
+        The cumulative sequence lengths of the sequences in
+        the batch, used to index into sequence. prepended by 0.
+        if
+        x = [5, 1, 1, 1] <- continuous batching (batch=4)
+        then
+        query_start_loc = [0, 5, 6, 7, 8] <- the starting index of the next sequence; while the last value is
+           the ending index of the last sequence
+        [length(query_start_loc)-1 == batch]
+        for example: query_start_loc = torch.Tensor([0,10,16,17]),
+        x.shape=(dim,17)
+    cache_indices: (batch)  int32
+        indicates the corresponding state index,
+        like so: conv_state = conv_states[cache_indices[batch_id]]
+    has_initial_state: (batch) bool
+        indicates whether should the kernel take the current state as initial
+        state for the calculations
+        [single boolean for each sequence in the batch: True or False]
+    bias: (dim,)
+    activation: either None or "silu" or "swish" or True
+    pad_slot_id: int
+        if cache_indices is passed, lets the kernel identify padded
+        entries that will not be processed,
+        for example: cache_indices = [pad_slot_id, 1, 20, pad_slot_id]
+        in this case, the kernel will not process entries at
+        indices 0 and 3
+    block_idx_first_scheduled_token: (batch,), dtype int32
+        The pointer into cache_indices, where the first cache block to be filled is located.
+    block_idx_last_scheduled_token: (batch,), dtype int32
+        The pointer into cache_indices, where the last cache block to be filled is located.
+    initial_state_idx: (batch,), dtype int32
+        The pointer into cache_indices, where the cache block containing the initial state is located.
+    num_computed_tokens: (batch,), dtype int32
+        The number of tokens already completed for each sequence
+    block_size_to_align: int
+        The block size to align the cached states to
+    out: same shape as `x`
+    """
+    if isinstance(activation, bool) and activation:
+        activation = "silu"
+
+    args = None
+    # Store original dtype to cast back at the end
+    original_x_dtype = x.dtype
+    x = x.to(conv_states.dtype)
+    out = torch.empty_like(x)
+    if metadata is not None:
+        nums_dict = metadata.nums_dict
+        args = nums_dict
+        batch_ptr = metadata.batch_ptr
+        token_chunk_offset_ptr = metadata.token_chunk_offset_ptr
+    else:
+        seqlens = query_start_loc.diff().to("cpu")
+        args = seqlens
+        MAX_NUM_PROGRAMS = 1024
+
+        batch_ptr = torch.full(
+            (MAX_NUM_PROGRAMS,), PAD_SLOT_ID, dtype=torch.int32, device=x.device
+        )  # tracking which seq-idx the Triton program is handling
+        token_chunk_offset_ptr = torch.full(
+            (MAX_NUM_PROGRAMS,), PAD_SLOT_ID, dtype=torch.int32, device=x.device
+        )  # tracking BLOCK_M-based index in the sequence the Triton program is handling
+
+    is_channel_last = (x.stride(0) == 1) & (x.stride(1) > 1)
+    dim, cu_seqlen = x.shape
+    _, width = weight.shape
+    state_len = width - 1
+    np2_statelen = triton.next_power_of_2(state_len)
+
+    padded_batch = query_start_loc.size(0) - 1
+    stride_x_dim = x.stride(0)
+    stride_x_token = x.stride(1)
+    stride_w_dim = weight.stride(0)
+    stride_w_width = weight.stride(1)
+    stride_istate_seq = 0
+    stride_istate_dim = 0
+    stride_istate_token = 0
+    num_cache_lines = 0
+    BLOCK_M = 8
+    if conv_states is not None:
+        # extensions to support vLLM:
+        # 1. conv_states is used to replaced initial_states
+        # 2. conv_states serve as a cache with num cache lines can be larger than batch size
+        # 3. mapping from sequence x[idx] to a cache line at index as specified via cache_indices[idx]
+        # 4. computation can be skipped if cache_indices[idx] == pad_slot_id
+        num_cache_lines = conv_states.size(0)
+        assert (
+            num_cache_lines == conv_states.shape[0]
+            and dim == conv_states.shape[1]
+            and width - 1 <= conv_states.shape[2]
+        )
+        stride_istate_seq = conv_states.stride(0)
+        stride_istate_dim = conv_states.stride(1)
+        stride_istate_token = conv_states.stride(2)
+    if out.dim() == 2:
+        stride_o_dim = out.stride(0)
+        stride_o_token = out.stride(1)
+    else:
+        stride_o_dim = out.stride(1)
+        stride_o_token = out.stride(2)
+    stride_cache_indices = cache_indices.stride(0) if cache_indices is not None else 0
+
+    if validate_data:
+        assert x.dim() == 2
+        assert query_start_loc is not None
+        assert query_start_loc.dim() == 1
+        assert x.stride(0) == 1 or x.stride(1) == 1
+        if bias is not None:
+            assert bias.dim() == 1
+            assert dim == bias.size(0)
+        if cache_indices is not None:
+            assert cache_indices.dim() == 1
+            assert padded_batch == cache_indices.size(0)
+        if has_initial_state is not None:
+            assert has_initial_state.size() == (padded_batch,)
+            assert conv_states is not None, (
+                "ERROR: `has_initial_state` is used, which needs also `conv_states`"
+            )
+        assert weight.stride(1) == 1
+        assert (dim, width) == weight.shape
+        assert is_channel_last, "Need to run in channel-last layout"
+        if block_size_to_align is not None and block_size_to_align > 0:
+            assert (block_size_to_align % BLOCK_M) == 0, (
+                "The mamba block size needs to be divisible by the BLOCK_M"
+            )
+        else:
+            block_size_to_align = BLOCK_M
+
+    if metadata is None:
+
+        def num_program(META, seqlens):
+            tot = 0
+
+            mlist = []
+            offsetlist = []  # type: ignore
+
+            nums = -(-seqlens // META["BLOCK_M"])
+
+            tot = nums.sum().item()
+            mlist = np.repeat(np.arange(len(nums)), nums)
+            for idx, num in enumerate(nums):
+                offsetlist.extend(
+                    range(num)
+                )  # chunk-idx if a sequence is split into multiple chunks
+
+            if META["batch_ptr"].nelement() < len(mlist):
+                newlen = len(mlist) + 1
+                META["batch_ptr"].resize_(newlen).fill_(PAD_SLOT_ID)
+                META["token_chunk_offset_ptr"].resize_(newlen).fill_(PAD_SLOT_ID)
+
+            if META["batch_ptr"].nelement() >= len(mlist):
+                META["batch_ptr"][0 : len(mlist)].copy_(
+                    torch.from_numpy(np.array(mlist))
+                )
+                META["token_chunk_offset_ptr"][0 : len(mlist)].copy_(
+                    torch.from_numpy(np.array(offsetlist))
+                )
+
+            META["batch_ptr"] = META["batch_ptr"].to(META["x_ptr"].device)
+            META["token_chunk_offset_ptr"] = META["token_chunk_offset_ptr"].to(
+                META["x_ptr"].device
+            )
+            return tot
+    else:
+
+        def num_program(META, nums_dict):
+            tot = nums_dict[META["BLOCK_M"]]["tot"]
+
+            mlist = nums_dict[META["BLOCK_M"]]["mlist"]
+            mlist_len = nums_dict[META["BLOCK_M"]]["mlist_len"]
+
+            offsetlist = nums_dict[META["BLOCK_M"]]["offsetlist"]
+
+            if nums_dict[META["BLOCK_M"]]["batch_ptr"] is not None:
+                META["batch_ptr"] = nums_dict[META["BLOCK_M"]]["batch_ptr"]
+                META["token_chunk_offset_ptr"] = nums_dict[META["BLOCK_M"]][
+                    "token_chunk_offset_ptr"
+                ]
+            else:
+                if META["batch_ptr"].nelement() < mlist_len:
+                    newlen = mlist_len + 1
+                    META["batch_ptr"].resize_(newlen).fill_(PAD_SLOT_ID)
+                    META["token_chunk_offset_ptr"].resize_(newlen).fill_(PAD_SLOT_ID)
+
+                if META["batch_ptr"].nelement() >= mlist_len:
+                    META["batch_ptr"][0:mlist_len].copy_(mlist)
+                    META["token_chunk_offset_ptr"][0:mlist_len].copy_(offsetlist)
+            return tot
+
+    def grid(META):
+        return (
+            triton.cdiv(dim, META["BLOCK_N"]),
+            num_program(META, args),
+        )
+
+    if batch_ptr.device != x.device:
+        batch_ptr = batch_ptr.to(x.device)
+        token_chunk_offset_ptr = token_chunk_offset_ptr.to(x.device)
+
+    _causal_conv1d_fwd_kernel[grid](
+        # Pointers to matrices
+        x,
+        weight,
+        bias,
+        conv_states,
+        cache_indices,
+        has_initial_state,
+        query_start_loc,
+        batch_ptr,
+        token_chunk_offset_ptr,
+        block_idx_first_scheduled_token,
+        block_idx_last_scheduled_token,
+        initial_state_idx,
+        num_computed_tokens,
+        num_accepted_tokens,
+        out,
+        # Matrix dimensions
+        dim,
+        num_cache_lines,
+        # stride
+        stride_x_dim,
+        stride_x_token,
+        stride_w_dim,
+        stride_w_width,
+        stride_istate_seq,
+        stride_istate_dim,
+        stride_istate_token,
+        stride_cache_indices,
+        stride_o_dim,
+        stride_o_token,
+        block_size_to_align // BLOCK_M,
+        # others
+        pad_slot_id,
+        null_block_id,
+        # META
+        HAS_BIAS=bias is not None,
+        KERNEL_WIDTH=width,
+        SILU_ACTIVATION=activation in ["silu", "swish"],
+        IS_APC_ENABLED=block_idx_last_scheduled_token is not None,
+        IS_SPEC_DECODING=num_accepted_tokens is not None,
+        HAS_NULL_BLOCK=null_block_id is not None,
+        NP2_STATELEN=np2_statelen,
+        # launch_cooperative_grid=True
+        BLOCK_M=BLOCK_M,
+        BLOCK_N=256,
+        num_stages=2,
+        launch_pdl=current_platform.is_arch_support_pdl(),
+    )
+    return out.to(original_x_dtype)
+
+
+@triton.jit(do_not_specialize_on_alignment=["num_cache_lines"])
+def _causal_conv1d_update_kernel(
+    # Pointers to matrices
+    x_ptr,  # (batch, dim, seqlen)
+    w_ptr,  # (dim, width)
+    bias_ptr,
+    conv_state_ptr,
+    conv_state_indices_ptr,
+    num_accepted_tokens_ptr,
+    query_start_loc_ptr,  # (batch + 1)
+    block_idx_last_scheduled_token,  # (batch,)
+    initial_state_idx,  # (batch,)
+    o_ptr,  # (batch, dim, seqlen)
+    # Matrix dimensions
+    batch: int,
+    dim: tl.constexpr,
+    seqlen: tl.constexpr,
+    state_len: tl.constexpr,
+    num_cache_lines,  # added to support vLLM larger cache lines
+    # Strides
+    stride_x_seq: tl.constexpr,
+    stride_x_dim: tl.constexpr,
+    stride_x_token: tl.int64,
+    stride_w_dim: tl.constexpr,
+    stride_w_width: tl.constexpr,
+    stride_conv_state_seq: tl.constexpr,
+    stride_conv_state_dim: tl.constexpr,
+    stride_conv_state_tok: tl.constexpr,
+    stride_state_indices: tl.constexpr,
+    stride_o_seq: tl.constexpr,
+    stride_o_dim: tl.constexpr,
+    stride_o_token: tl.int64,
+    # others
+    null_block_id: tl.constexpr,
+    # Meta-parameters
+    HAS_BIAS: tl.constexpr,
+    KERNEL_WIDTH: tl.constexpr,
+    SILU_ACTIVATION: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+    IS_APC_ENABLED: tl.constexpr,
+    IS_SPEC_DECODING: tl.constexpr,
+    NP2_STATELEN: tl.constexpr,
+    HAS_NULL_BLOCK: tl.constexpr,
+    BLOCK_N: tl.constexpr,
+    launch_pdl: tl.constexpr,
+):
+    if launch_pdl:
+        tl.extra.cuda.gdc_wait()
+
+    # ruff: noqa: E501
+    idx_seq = tl.program_id(0)
+    if idx_seq >= batch:
+        if launch_pdl:
+            tl.extra.cuda.gdc_launch_dependents()
+        return
+
+    # [BLOCK_N,] elements along the feature-dimension (channel)
+    idx_feats = tl.program_id(1) * BLOCK_N + tl.arange(0, BLOCK_N)
+
+    if IS_APC_ENABLED:
+        # Get the state from the initial_state_idx
+        conv_state_init = tl.load(initial_state_idx + idx_seq)
+        current_last_index = tl.load(block_idx_last_scheduled_token + idx_seq)
+    else:
+        conv_state_init = 0
+        current_last_index = 0
+
+    # cache_idx
+    conv_states_input_coord = tl.load(
+        conv_state_indices_ptr + idx_seq * stride_state_indices + conv_state_init
+    ).to(tl.int64)
+
+    if HAS_NULL_BLOCK:  # noqa
+        if conv_states_input_coord == null_block_id:
+            # not processing as this is not the actual sequence
+            if launch_pdl:
+                tl.extra.cuda.gdc_launch_dependents()
+            return
+
+    if IS_VARLEN:
+        query_start_index = tl.load(query_start_loc_ptr + idx_seq).to(tl.int64)
+        query_end_index = tl.load(query_start_loc_ptr + (idx_seq + 1)).to(tl.int64)
+        # revise state_len and seqlen
+        state_len = state_len - (seqlen - (query_end_index - query_start_index))
+        seqlen = query_end_index - query_start_index
+        x_offset = query_start_index * stride_x_token
+        o_offset = query_start_index * stride_o_token
+    else:
+        query_start_index = idx_seq * seqlen
+        query_end_index = query_start_index + seqlen
+        x_offset = idx_seq * stride_x_seq
+        o_offset = idx_seq * stride_o_seq
+
+    if query_start_index == query_end_index:
+        if launch_pdl:
+            tl.extra.cuda.gdc_launch_dependents()
+        return
+
+    if IS_SPEC_DECODING:
+        # The rolling of conv state:
+        #
+        # Before forward, the conv_state is:
+        # [history1, history2, ..., historyM].
+        #
+        # After forward, the conv_state becomes:
+        # [history2, ..., historyM, draft1, draft2, ..., draftN].
+        #
+        # After acceptance, it becomes:
+        #
+        # - accept 1 tokens: [history2, ..., historyM, draft1]
+        # - accept 2 tokens: [history3, ..., historyM, draft1, draft2]
+        # - and so on.
+        conv_state_token_offset = (
+            tl.load(num_accepted_tokens_ptr + idx_seq).to(tl.int64) - 1
+        )
+    else:
+        conv_state_token_offset = 0
+
+    # STEP 1: READ init_state data
+    conv_states_base = (
+        conv_state_ptr
+        + (conv_states_input_coord * stride_conv_state_seq)
+        + (idx_feats * stride_conv_state_dim)
+    )
+    mask_w = idx_feats < dim
+
+    prior_tokens = conv_states_base + conv_state_token_offset * stride_conv_state_tok
+    if KERNEL_WIDTH >= 2:
+        conv_states_ptrs = prior_tokens  # [BLOCK_N]
+        col0 = tl.load(conv_states_ptrs, mask_w, 0.0)
+    if KERNEL_WIDTH >= 3:
+        conv_states_ptrs = prior_tokens + 1 * stride_conv_state_tok  # [BLOCK_N]
+        col1 = tl.load(conv_states_ptrs, mask_w, 0.0)
+    if KERNEL_WIDTH >= 4:
+        conv_states_ptrs = prior_tokens + 2 * stride_conv_state_tok  # [BLOCK_N]
+        col2 = tl.load(conv_states_ptrs, mask_w, 0.0)
+    if KERNEL_WIDTH >= 5:
+        conv_states_ptrs = prior_tokens + 3 * stride_conv_state_tok  # [BLOCK_N]
+        col3 = tl.load(conv_states_ptrs, mask_w, 0.0)
+    if KERNEL_WIDTH >= 6:
+        conv_states_ptrs = prior_tokens + 4 * stride_conv_state_tok  # [BLOCK_N]
+        col4 = tl.load(conv_states_ptrs, mask_w, 0.0)
+
+    # STEP 2: assume state_len > seqlen
+    idx_tokens = tl.arange(0, NP2_STATELEN)  # [BLOCK_M]
+
+    # With speculative decoding, the conv_state updates works in a sliding
+    # window manner, at each forward pass, the tokens are shift by 1, so we
+    # load since idx_tokens + 1.
+    conv_state_ptrs_source = (
+        conv_state_ptr
+        + (conv_states_input_coord * stride_conv_state_seq)
+        + conv_state_token_offset * stride_conv_state_tok
+        + (idx_feats * stride_conv_state_dim)[None, :]
+        + ((idx_tokens + (1 if IS_SPEC_DECODING else seqlen)) * stride_conv_state_tok)[
+            :, None
+        ]
+    )  # [BLOCK_M, BLOCK_N]
+    mask = (
+        (conv_states_input_coord < num_cache_lines)
+        & ((idx_tokens + seqlen) < state_len)[:, None]
+        & (idx_feats < dim)[None, :]
+    )
+    conv_state = tl.load(conv_state_ptrs_source, mask, other=0.0)
+
+    VAL = state_len - seqlen
+    x_base = x_ptr + x_offset + (idx_feats * stride_x_dim)  # [BLOCK_N]
+
+    x_ptrs = (
+        x_base[None, :] + ((idx_tokens - VAL) * stride_x_token)[:, None]
+    )  # [BLOCK_M, BLOCK_N]
+
+    mask_x = (
+        (idx_tokens - VAL >= 0)[:, None]
+        & (idx_tokens - VAL < seqlen)[:, None]
+        & (idx_feats < dim)[None, :]
+    )  # token-index  # token-index  # feature-index
+    loaded_x = tl.load(x_ptrs, mask_x, 0.0)
+    tl.debug_barrier()
+
+    new_conv_state = tl.where(mask, conv_state, loaded_x)
+
+    # Get the state from the initial_state_idx
+    # cache_idx
+    conv_states_offset = tl.load(
+        conv_state_indices_ptr + idx_seq * stride_state_indices + current_last_index
+    ).to(tl.int64)
+    conv_state_ptrs_target = (
+        conv_state_ptr
+        + (conv_states_offset * stride_conv_state_seq)  # Offset from seq
+        + (idx_feats * stride_conv_state_dim)
+    )[None, :] + (  # [BLOCK_N,]
+        idx_tokens * stride_conv_state_tok
+    )[:, None]
+    mask = (idx_tokens < state_len)[:, None] & (idx_feats < dim)[None, :]
+    tl.store(conv_state_ptrs_target, new_conv_state, mask)
+
+    # STEP 3: init accumulator
+    if HAS_BIAS:
+        bias = bias_ptr + idx_feats
+        mask_bias = idx_feats < dim
+        acc_preload = tl.load(bias, mask=mask_bias, other=0.0).to(
+            tl.float32
+        )  # [BLOCK_N]
+    else:
+        acc_preload = tl.zeros((BLOCK_N,), dtype=tl.float32)
+
+    # STEP 4:
+    # PRE-LOAD WEIGHTS
+    # first kernel column, configured for weights to handle BLOCK_N features in range
+    w_base = w_ptr + (idx_feats * stride_w_dim)  # [BLOCK_N,]
+    mask_w = idx_feats < dim
+    if KERNEL_WIDTH >= 2:
+        w_ptrs = w_base + (0 * stride_w_width)  # [BLOCK_N] tensor
+        w_col0 = tl.load(w_ptrs, mask_w, other=0.0)
+        w_ptrs = w_base + (1 * stride_w_width)  # [BLOCK_N] tensor
+        w_col1 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 3:
+        w_ptrs = w_base + (2 * stride_w_width)  # [BLOCK_N] tensor
+        w_col2 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 4:
+        w_ptrs = w_base + (3 * stride_w_width)  # [BLOCK_N] tensor
+        w_col3 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 5:
+        w_ptrs = w_base + (4 * stride_w_width)  # [BLOCK_N] tensor
+        w_col4 = tl.load(w_ptrs, mask_w, other=0.0)
+    if KERNEL_WIDTH >= 6:
+        w_ptrs = w_base + (5 * stride_w_width)  # [BLOCK_N] tensor
+        w_col5 = tl.load(w_ptrs, mask_w, other=0.0)
+
+    x_base_1d = x_base  # starting of chunk [BLOCK_N]
+    mask_x_1d = idx_feats < dim
+
+    # STEP 5: compute each token
+    if launch_pdl:
+        tl.extra.cuda.gdc_launch_dependents()
+
+    for idx_token in tl.range(seqlen):
+        acc = acc_preload
+
+        matrix_w = w_col0
+        matrix_x = col0
+        for j in tl.static_range(KERNEL_WIDTH):
+            if KERNEL_WIDTH == 2:
+                if j == 1:  # KERNEL_WIDTH-1:
+                    matrix_w = w_col1
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 3:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 4:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    matrix_x = col2
+                elif j == 3:
+                    matrix_w = w_col3
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 5:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    matrix_x = col2
+                elif j == 3:
+                    matrix_w = w_col3
+                    matrix_x = col3
+                elif j == 4:
+                    matrix_w = w_col4
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+            elif KERNEL_WIDTH == 6:
+                if j == 1:
+                    matrix_w = w_col1
+                    matrix_x = col1
+                elif j == 2:
+                    matrix_w = w_col2
+                    matrix_x = col2
+                elif j == 3:
+                    matrix_w = w_col3
+                    matrix_x = col3
+                elif j == 4:
+                    matrix_w = w_col4
+                    matrix_x = col4
+                elif j == 5:
+                    matrix_w = w_col5
+                    x_ptrs_1d = x_base_1d + idx_token * stride_x_token  # [BLOCK_N]
+                    matrix_x = tl.load(x_ptrs_1d, mask=mask_x_1d)
+
+            acc += matrix_x.to(tl.float32) * matrix_w.to(tl.float32)  # [BLOCK_N]
+
+        if KERNEL_WIDTH == 2:
+            col0 = matrix_x
+        elif KERNEL_WIDTH == 3:
+            col0 = col1
+            col1 = matrix_x
+        elif KERNEL_WIDTH == 4:
+            col0 = col1
+            col1 = col2
+            col2 = matrix_x
+        elif KERNEL_WIDTH == 5:
+            col0 = col1
+            col1 = col2
+            col2 = col3
+            col3 = matrix_x
+        elif KERNEL_WIDTH == 6:
+            col0 = col1
+            col1 = col2
+            col2 = col3
+            col3 = col4
+            col4 = matrix_x
+
+        if SILU_ACTIVATION:
+            acc = acc / (1 + tl.exp(-acc))
+        mask_1d = (idx_token < seqlen) & (
+            idx_feats < dim
+        )  # token-index  # feature-index
+        o_ptrs = (
+            o_ptr + o_offset + idx_token * stride_o_token + (idx_feats * stride_o_dim)
+        )
+
+        tl.store(o_ptrs, acc, mask=mask_1d)
+
+
+def causal_conv1d_update(
+    x: torch.Tensor,
+    conv_state: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor | None = None,
+    activation: bool | str | None = None,
+    conv_state_indices: torch.Tensor | None = None,
+    num_accepted_tokens: torch.Tensor | None = None,
+    query_start_loc: torch.Tensor | None = None,
+    max_query_len: int = -1,
+    null_block_id: int = NULL_BLOCK_ID,
+    block_idx_last_scheduled_token: torch.Tensor | None = None,
+    initial_state_idx: torch.Tensor | None = None,
+    validate_data=False,
+    out: torch.Tensor | None = None,
+):
+    """
+    x: Input tensor which can take the following shapes:
+
+    - `[batch, dim]` - single token prediction
+    - `[batch, dim, seqlen]` - single or multiple tokens prediction
+    - `[num_tokens, dim]` - continuous batching, where num_tokens is
+        the total tokens of all sequences in that batch
+
+    conv_state: (..., dim, state_len), where state_len >= width - 1
+    weight: (dim, width)
+    bias: (dim,)
+    conv_state_indices: (batch,), dtype int32
+        If not None, the conv_state is a larger tensor along the batch dim,
+        and we are selecting the batch coords specified by conv_state_indices.
+        Useful for a continuous batching scenario.
+    block_idx_last_scheduled_token: (batch,), dtype int32
+        The pointer into conv_state_indices, where the last cache block to be filled is located.
+    initial_state_idx: (batch,), dtype int32
+        The pointer into conv_state_indices, where the cache block containing the initial state is located.
+    num_accepted_tokens: (batch,), dtype int32
+        If not None, it indicates the number of accepted tokens for each
+        sequence in the batch.
+        This is used in speculative decoding, where the conv_state is updated
+        in a sliding window manner.
+    query_start_loc: (batch + 1,) int32
+        If not None, the inputs is given in a varlen fashion and this indicates
+        the starting index of each sequence in the batch.
+    max_query_len: int
+        If query_start_loc is not None, this indicates the maximum query
+        length in the batch.
+    null_block_id: int
+            Block ID used to identify padded entries in
+            conv_state_indices. Block 0 is the null block.
+            for example: conv_state_indices = [null_block_id, 1, 20, null_block_id]
+            in this case, the kernel will not process entries at
+            indices 0 and 3
+    out: optional output tensor with the same shape as `x`. When omitted,
+        the input is overwritten.
+    """
+    if validate_data:
+        assert null_block_id is not None
+        assert x.stride(1) == 1
+    if isinstance(activation, bool):
+        activation = "silu" if activation is True else None
+    elif activation is not None:
+        assert activation in ["silu", "swish"]
+
+    original_x_dtype = x.dtype
+    x = x.to(conv_state.dtype)
+    if out is None:
+        out = x
+    else:
+        if out.shape != x.shape:
+            raise ValueError(
+                f"`out` shape {tuple(out.shape)} must match `x` shape {tuple(x.shape)}."
+            )
+        if out.dtype != original_x_dtype or out.device != x.device:
+            raise ValueError(
+                "`out` must have the same dtype and device as the input `x`."
+            )
+    unsqueeze = query_start_loc is None and x.dim() == 2
+    if unsqueeze:
+        # make it (batch, dim, seqlen) with seqlen == 1
+        x = x.unsqueeze(-1)
+        out = out.unsqueeze(-1)
+    if query_start_loc is None:
+        batch, dim, seqlen = x.shape
+    else:
+        assert conv_state_indices is not None
+        batch = conv_state_indices.size(0)
+        dim = x.size(1)
+        seqlen = max_query_len
+    _, width = weight.shape
+    # conv_state: (..., dim, state_len), where state_len >= width - 1
+    num_cache_lines, _, state_len = conv_state.size()
+
+    if validate_data:
+        assert dim == weight.size(0)
+        assert state_len >= width - 1
+        # when above happens, we don't shift-left to keep any records in conv_state
+        assert dim == conv_state.size(1)
+        if conv_state_indices is None:
+            assert conv_state.size(0) >= batch
+        else:
+            assert batch == conv_state_indices.shape[0], (
+                f"ERROR: conv_state_indices should have shape ({batch},*) but got {conv_state_indices.shape}"
+            )
+
+        assert num_cache_lines >= batch
+        assert weight.stride(1) == 1  # Need this
+
+    stride_w_dim, stride_w_width = weight.stride()
+
+    if query_start_loc is None:
+        # X (batch, dim, seqlen)
+        stride_x_seq, stride_x_dim, stride_x_token = x.stride()
+        stride_o_seq, stride_o_dim, stride_o_token = out.stride()
+    else:
+        # X (dim, cu_seqlen)
+        stride_x_token, stride_x_dim = x.stride()
+        stride_x_seq = 0
+        stride_o_token, stride_o_dim = out.stride()
+        stride_o_seq = 0
+
+    stride_istate_seq, stride_istate_dim, stride_istate_token = conv_state.stride()
+    stride_state_indices = (
+        conv_state_indices.stride(0) if conv_state_indices is not None else 0
+    )
+    if num_accepted_tokens is not None:
+        state_len = width - 1 + (seqlen - 1)  # effective state_len needed
+    else:
+        state_len = width - 1
+    np2_statelen = triton.next_power_of_2(state_len)
+
+    def grid(META):
+        return (
+            batch,
+            triton.cdiv(dim, META["BLOCK_N"]),
+        )
+
+    _causal_conv1d_update_kernel[grid](
+        # Pointers to matrices
+        x,
+        weight,
+        bias,
+        conv_state,
+        conv_state_indices,
+        num_accepted_tokens,
+        query_start_loc,
+        block_idx_last_scheduled_token,
+        initial_state_idx,
+        out,
+        # Matrix dimensions
+        batch,
+        dim,
+        seqlen,
+        state_len,
+        num_cache_lines,
+        # stride
+        stride_x_seq,
+        stride_x_dim,
+        stride_x_token,
+        stride_w_dim,
+        stride_w_width,
+        stride_istate_seq,
+        stride_istate_dim,
+        stride_istate_token,
+        stride_state_indices,
+        stride_o_seq,
+        stride_o_dim,
+        stride_o_token,
+        # others
+        null_block_id,
+        # META
+        HAS_BIAS=bias is not None,
+        KERNEL_WIDTH=width,
+        SILU_ACTIVATION=activation in ["silu", "swish"],
+        IS_VARLEN=query_start_loc is not None,
+        IS_APC_ENABLED=block_idx_last_scheduled_token is not None,
+        IS_SPEC_DECODING=num_accepted_tokens is not None,
+        NP2_STATELEN=np2_statelen,
+        HAS_NULL_BLOCK=null_block_id is not None,
+        BLOCK_N=256,
+        launch_pdl=current_platform.is_arch_support_pdl(),
+    )
+    if unsqueeze:
+        out = out.squeeze(-1)
+    return out.to(original_x_dtype)
+
+
+if current_platform.is_cpu():
+    from vllm.model_executor.layers.mamba.ops.cpu.causal_conv1d import (
+        causal_conv1d_fn_cpu,
+        causal_conv1d_update_cpu,
+    )
+
+    causal_conv1d_fn = causal_conv1d_fn_cpu  # type: ignore
+    causal_conv1d_update = causal_conv1d_update_cpu  # type: ignore

bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_o.py

@@ -0,0 +1,190 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Songlin Yang, Yu Zhang
+#
+# This file contains code copied from the flash-linear-attention project.
+# The original source code was licensed under the MIT license and included
+# the following copyright notice:
+# Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
+
+# ruff: noqa: E501
+
+
+import torch
+
+from vllm.triton_utils import tl, triton
+
+from .index import prepare_chunk_indices
+from .op import exp
+from .utils import FLA_CHUNK_SIZE, check_shared_mem, is_nvidia_hopper
+
+BKV_LIST = [64, 128] if check_shared_mem() else [32, 64]
+NUM_WARPS = [2, 4] if is_nvidia_hopper else [2, 4, 8]
+
+
+@triton.heuristics(
+    {
+        "USE_G": lambda args: args["g"] is not None,
+        "IS_VARLEN": lambda args: args["cu_seqlens"] is not None,
+    }
+)
+@triton.autotune(
+    configs=[
+        triton.Config({"BK": BK, "BV": BV}, num_warps=num_warps, num_stages=num_stages)
+        for BK in BKV_LIST
+        for BV in BKV_LIST
+        for num_warps in NUM_WARPS
+        for num_stages in [2, 3, 4]
+    ],
+    key=["H", "K", "V", "BT"],
+)
+@triton.jit(do_not_specialize=["T"])
+def chunk_fwd_kernel_o(
+    q,
+    k,
+    v,
+    h,
+    g,
+    o,
+    cu_seqlens,
+    chunk_indices,
+    scale,
+    T,
+    H: tl.constexpr,
+    Hg: tl.constexpr,
+    K: tl.constexpr,
+    V: tl.constexpr,
+    BT: tl.constexpr,
+    BK: tl.constexpr,
+    BV: tl.constexpr,
+    USE_G: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+):
+    i_v, i_t, i_bh = tl.program_id(0), tl.program_id(2), tl.program_id(1)
+    i_b, i_h = i_bh // H, i_bh % H
+
+    if IS_VARLEN:
+        i_tg = i_t
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+        NT = tl.cdiv(T, BT)
+    else:
+        NT = tl.cdiv(T, BT)
+        i_tg = i_b * NT + i_t
+        bos, eos = i_b * T, i_b * T + T
+
+    # offset calculation
+    q += (bos * Hg + i_h // (H // Hg)) * K
+    k += (bos * Hg + i_h // (H // Hg)) * K
+    v += (bos * H + i_h) * V
+    o += (bos * H + i_h) * V
+    h += (i_tg * H + i_h).to(tl.int64) * V * K
+
+    b_o = tl.zeros([BT, BV], dtype=tl.float32)
+    b_A = tl.zeros([BT, BT], dtype=tl.float32)
+
+    for i_k in range(tl.cdiv(K, BK)):
+        p_q = tl.make_block_ptr(
+            q, (T, K), (Hg * K, 1), (i_t * BT, i_k * BK), (BT, BK), (1, 0)
+        )
+        p_k = tl.make_block_ptr(
+            k, (K, T), (1, Hg * K), (i_k * BK, i_t * BT), (BK, BT), (0, 1)
+        )
+        p_h = tl.make_block_ptr(
+            h, (V, K), (K, 1), (i_v * BV, i_k * BK), (BV, BK), (1, 0)
+        )
+        # [BT, BK]
+        b_q = tl.load(p_q, boundary_check=(0, 1))
+        # [BK, BT]
+        b_k = tl.load(p_k, boundary_check=(0, 1))
+        # [BV, BK]
+        b_h = tl.load(p_h, boundary_check=(0, 1))
+
+        # [BT, BK] @ [BK, BV] -> [BT, BV]
+        b_o += tl.dot(b_q, tl.trans(b_h))
+        # [BT, BK] @ [BK, BT] -> [BT, BT]
+        b_A += tl.dot(b_q, b_k)
+
+    if USE_G:
+        g += bos * H + i_h
+        p_g = tl.make_block_ptr(g, (T,), (H,), (i_t * BT,), (BT,), (0,))
+        b_g = tl.load(p_g, boundary_check=(0,))
+        b_o = b_o * exp(b_g)[:, None]
+        b_A = b_A * exp(b_g[:, None] - b_g[None, :])
+
+    o_t = i_t * BT + tl.arange(0, BT)
+    m_t = o_t < T
+    m_A = (o_t[:, None] >= o_t[None, :]) & (m_t[:, None] & m_t)
+    b_A = tl.where(m_A, b_A, 0)
+
+    p_v = tl.make_block_ptr(
+        v, (T, V), (H * V, 1), (i_t * BT, i_v * BV), (BT, BV), (1, 0)
+    )
+    p_o = tl.make_block_ptr(
+        o, (T, V), (H * V, 1), (i_t * BT, i_v * BV), (BT, BV), (1, 0)
+    )
+    b_v = tl.load(p_v, boundary_check=(0, 1))
+
+    # to fix mma -> mma layout conversion
+    # already solved by triton v3.2 or higher
+    b_o = b_o * scale + tl.dot(b_A.to(b_v.dtype), b_v) * scale
+    tl.store(p_o, b_o.to(p_o.dtype.element_ty), boundary_check=(0, 1))
+
+
+def chunk_fwd_o(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    h: torch.Tensor,
+    g: torch.Tensor | None = None,  # cumsum of log decay
+    scale: float | None = None,
+    cu_seqlens: torch.Tensor | None = None,
+    chunk_indices: torch.Tensor | None = None,
+    chunk_size: int = FLA_CHUNK_SIZE,
+    core_attn_out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    B, T, Hg, K, V = *q.shape, v.shape[-1]
+    H = v.shape[-2]
+    BT = chunk_size
+    if chunk_indices is None and cu_seqlens is not None:
+        chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
+    NT = triton.cdiv(T, BT) if cu_seqlens is None else len(chunk_indices)
+    if scale is None:
+        scale = k.shape[-1] ** -0.5
+
+    if core_attn_out is not None:
+        assert core_attn_out.numel() >= v.numel(), (
+            f"core_attn_out too small: {core_attn_out.numel()} < {v.numel()}"
+        )
+        o = core_attn_out[: v.numel()].view(*v.shape)
+    else:
+        o = torch.empty_like(v)
+
+    def grid(meta):
+        return (triton.cdiv(V, meta["BV"]), B * H, NT)
+
+    chunk_fwd_kernel_o[grid](
+        q,
+        k,
+        v,
+        h,
+        g,
+        o,
+        cu_seqlens,
+        chunk_indices,
+        scale,
+        T=T,
+        H=H,
+        Hg=Hg,
+        K=K,
+        V=V,
+        BT=BT,
+    )
+    return o

bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/chunk_scaled_dot_kkt.py

@@ -0,0 +1,176 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Songlin Yang, Yu Zhang
+#
+# This file contains code copied from the flash-linear-attention project.
+# The original source code was licensed under the MIT license and included
+# the following copyright notice:
+# Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
+# ruff: noqa: E501
+
+import torch
+
+from vllm.platforms import current_platform
+from vllm.triton_utils import tl, triton
+
+from .index import prepare_chunk_indices
+from .op import exp
+from .utils import FLA_CHUNK_SIZE
+
+# On RDNA (gfx11xx/gfx12xx) WMMA only
+# accepts 16-bit/int inputs, so a widened (e.g. fp32) tl.dot is lowered to a
+# software matmul (~190x amdgcn-stage blowup). There we cast both operands down
+# to k's native storage dtype (bf16/fp16) so fast WMMA is used instead.
+_CAST_DOT_TO_K_DTYPE = False
+if current_platform.is_rocm():
+    from vllm.platforms.rocm import on_gfx1x
+
+    _CAST_DOT_TO_K_DTYPE = on_gfx1x()
+
+
+@triton.heuristics(
+    {
+        "USE_G": lambda args: args["g"] is not None,
+        "IS_VARLEN": lambda args: args["cu_seqlens"] is not None,
+    }
+)
+@triton.autotune(
+    configs=[
+        triton.Config({"BK": BK}, num_warps=num_warps, num_stages=num_stages)
+        for BK in [32, 64, 128]
+        for num_warps in [2, 4, 8]
+        for num_stages in [2, 3, 4]
+    ],
+    key=["H", "K", "BT", "IS_VARLEN"],
+)
+@triton.jit(do_not_specialize=["T"])
+def chunk_scaled_dot_kkt_fwd_kernel(
+    k,
+    beta,
+    g,
+    A,
+    cu_seqlens,
+    chunk_indices,
+    T,
+    H: tl.constexpr,
+    Hg: tl.constexpr,
+    K: tl.constexpr,
+    BT: tl.constexpr,
+    BK: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+    USE_G: tl.constexpr,
+    CAST_DOT_TO_K_DTYPE: tl.constexpr,
+):
+    i_t, i_bh = tl.program_id(1), tl.program_id(0)
+    i_b, i_h = i_bh // H, i_bh % H
+    if IS_VARLEN:
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+    else:
+        bos, eos = i_b * T, i_b * T + T
+    o_t = i_t * BT + tl.arange(0, BT)
+    m_t = o_t < T
+
+    p_beta = tl.make_block_ptr(
+        beta + bos * H + i_h, (T,), (H,), (i_t * BT,), (BT,), (0,)
+    )
+    b_beta = tl.load(p_beta, boundary_check=(0,))
+
+    b_A = tl.zeros([BT, BT], dtype=tl.float32)
+    for i_k in range(tl.cdiv(K, BK)):
+        p_k = tl.make_block_ptr(
+            k + (bos * Hg + i_h // (H // Hg)) * K,
+            (T, K),
+            (Hg * K, 1),
+            (i_t * BT, i_k * BK),
+            (BT, BK),
+            (1, 0),
+        )
+        b_k = tl.load(p_k, boundary_check=(0, 1))
+        b_kb = b_k * b_beta[:, None]
+        if CAST_DOT_TO_K_DTYPE:
+            # RDNA: force operands to k's native dtype so WMMA is used.
+            b_A += tl.dot(b_kb.to(b_k.dtype), tl.trans(b_k))
+        else:
+            # Keep the promoted precision of the beta-scaled operand (WGMMA/MFMA).
+            b_A += tl.dot(b_kb, tl.trans(b_k).to(b_kb.dtype))
+
+    if USE_G:
+        p_g = tl.make_block_ptr(g + bos * H + i_h, (T,), (H,), (i_t * BT,), (BT,), (0,))
+        b_g = tl.load(p_g, boundary_check=(0,))
+        b_g_diff = b_g[:, None] - b_g[None, :]
+        b_A = b_A * exp(b_g_diff)
+
+    m_A = (o_t[:, None] > o_t[None, :]) & (m_t[:, None] & m_t)
+    b_A = tl.where(m_A, b_A, 0)
+    p_A = tl.make_block_ptr(
+        A + (bos * H + i_h) * BT, (T, BT), (BT * H, 1), (i_t * BT, 0), (BT, BT), (1, 0)
+    )
+    tl.store(p_A, b_A.to(p_A.dtype.element_ty), boundary_check=(0, 1))
+
+
+def chunk_scaled_dot_kkt_fwd(
+    k: torch.Tensor,
+    g: torch.Tensor | None = None,
+    beta: torch.Tensor | None = None,
+    cu_seqlens: torch.Tensor | None = None,
+    chunk_indices: torch.Tensor | None = None,
+    chunk_size: int = FLA_CHUNK_SIZE,
+    output_dtype: torch.dtype = torch.float32,
+) -> torch.Tensor:
+    r"""
+    Compute beta * K * K^T.
+
+    Args:
+        k (torch.Tensor):
+            The key tensor of shape `[B, T, H, K]`.
+        beta (torch.Tensor):
+            The beta tensor of shape `[B, T, H]`.
+        g (torch.Tensor):
+            The cumulative sum of the gate tensor of shape `[B, T, H]`. Default: `None`.
+        cu_seqlens (torch.Tensor):
+            The cumulative sequence lengths of the input tensor.
+            Default: None
+        chunk_indices (torch.Tensor):
+            Pre-computed chunk indices. If None and cu_seqlens is provided,
+            computed internally. Default: None
+        chunk_size (int):
+            The chunk size. Default: 64.
+        output_dtype (torch.dtype):
+            The dtype of the output tensor. Default: `torch.float32`
+
+    Returns:
+        beta * K * K^T of shape `[B, T, H, BT]` where `BT` is the chunk size.
+    """
+    # This kernel is slightly different from fla to support Q/K with different head numbers.
+    # In fla, Q/K always have the same head number, so Hg is always equal to H.
+    B, T, Hg, K = k.shape
+    H = beta.shape[-1]
+    BT = chunk_size
+    if chunk_indices is None and cu_seqlens is not None:
+        chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
+    NT = triton.cdiv(T, BT) if cu_seqlens is None else len(chunk_indices)
+
+    A = torch.empty(B, T, H, BT, device=k.device, dtype=output_dtype)
+    chunk_scaled_dot_kkt_fwd_kernel[(B * H, NT)](
+        k=k,
+        g=g,
+        beta=beta,
+        A=A,
+        cu_seqlens=cu_seqlens,
+        chunk_indices=chunk_indices,
+        T=T,
+        H=H,
+        Hg=Hg,
+        K=K,
+        BT=BT,
+        CAST_DOT_TO_K_DTYPE=_CAST_DOT_TO_K_DTYPE,
+    )
+    return A

bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/solve_tril.py

@@ -0,0 +1,558 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Songlin Yang, Yu Zhang
+#
+# This file contains code copied from the flash-linear-attention project.
+# The original source code was licensed under the MIT license and included
+# the following copyright notice:
+# Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
+# ruff: noqa: E501
+
+import os
+
+import torch
+
+from vllm.triton_utils import tl, triton
+
+from .index import prepare_chunk_indices
+from .op import make_tensor_descriptor
+from .utils import input_guard, is_amd, is_tma_supported
+
+FLA_TRIL_PRECISION = os.environ.get("FLA_TRIL_PRECISION", "ieee")
+ALLOWED_TRIL_PRECISIONS = ["ieee", "tf32"] if is_amd else ["ieee", "tf32", "tf32x3"]
+assert FLA_TRIL_PRECISION in ALLOWED_TRIL_PRECISIONS, (
+    f"FLA_TRIL_PRECISION must be one of {ALLOWED_TRIL_PRECISIONS}, but got {FLA_TRIL_PRECISION}"
+)
+
+
+@triton.heuristics({"IS_VARLEN": lambda args: args["cu_seqlens"] is not None})
+@triton.autotune(
+    configs=[
+        triton.Config({}, num_warps=num_warps, num_stages=num_stages)
+        for num_warps in [1, 2, 4, 8]
+        for num_stages in [2, 3, 4, 5]
+    ],
+    key=["BT"],
+)
+@triton.jit(do_not_specialize=["T"])
+def solve_tril_16x16_kernel(
+    A,
+    Ai,
+    cu_seqlens,
+    chunk_indices,
+    T,
+    H: tl.constexpr,
+    BT: tl.constexpr,
+    USE_TMA: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+    DOT_PRECISION: tl.constexpr,
+):
+    i_t, i_bh = tl.program_id(0), tl.program_id(1)
+    i_b, i_h = i_bh // H, i_bh % H
+    if IS_VARLEN:
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+    else:
+        bos, eos = i_b * T, i_b * T + T
+    o_i = tl.arange(0, 16)
+    m_A = o_i[:, None] > o_i[None, :]
+    m_I = o_i[:, None] == o_i[None, :]
+
+    A = A + (bos * H + i_h) * BT
+    Ai = Ai + (bos * H + i_h) * 16
+
+    offset = (i_t * 16) % BT
+    if not USE_TMA:
+        p_A = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * 16, offset), (16, 16), (1, 0)
+        )
+        # [16, 16]
+        b_A = tl.load(p_A, boundary_check=(0, 1)).to(tl.float32)
+    else:
+        desc = make_tensor_descriptor(A, [T, BT], [H * BT, 1], [16, 16])
+        desc_o = make_tensor_descriptor(Ai, [T, 16], [H * 16, 1], [16, 16])
+        b_A = desc.load([i_t * 16, offset]).to(tl.float32)
+    b_A = -tl.where(m_A, b_A, 0)
+
+    for i in range(2, min(16, T - i_t * 16)):
+        # [16]
+        b_a = -tl.load(A + (i_t * 16 + i) * H * BT + o_i + offset)
+        b_a = b_a + tl.sum(b_a[:, None] * b_A, 0)
+        b_A = tl.where((o_i == i)[:, None], b_a, b_A)
+    b_A += m_I
+    if not USE_TMA:
+        p_Ai = tl.make_block_ptr(
+            Ai, (T, 16), (H * 16, 1), (i_t * 16, 0), (16, 16), (1, 0)
+        )
+        tl.store(
+            p_Ai,
+            b_A.to(p_Ai.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+    else:
+        desc_o.store([i_t * 16, 0], b_A.to(desc_o.dtype, fp_downcast_rounding="rtne"))
+
+
+@triton.heuristics({"IS_VARLEN": lambda args: args["cu_seqlens"] is not None})
+@triton.autotune(
+    configs=[
+        triton.Config({}, num_warps=num_warps, num_stages=num_stages)
+        for num_warps in [1, 2, 4, 8]
+        for num_stages in [2, 3, 4, 5]
+    ],
+    key=["H", "BT", "IS_VARLEN"],
+)
+@triton.jit(do_not_specialize=["T"])
+def merge_16x16_to_32x32_inverse_kernel(
+    A,
+    Ai,
+    cu_seqlens,
+    chunk_indices,
+    T,
+    H: tl.constexpr,
+    BT: tl.constexpr,
+    USE_TMA: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+    DOT_PRECISION: tl.constexpr,
+):
+    i_t, i_bh = tl.program_id(0), tl.program_id(1)
+    i_b, i_h = i_bh // H, i_bh % H
+    if IS_VARLEN:
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+    else:
+        bos, eos = i_b * T, i_b * T + T
+
+    o_i = tl.arange(0, 16)
+    m_A = o_i[:, None] > o_i[None, :]
+    m_I = o_i[:, None] == o_i[None, :]
+    A += (bos * H + i_h) * BT
+    Ai += (bos * H + i_h) * BT
+
+    if not USE_TMA:
+        p_A_11 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT, 0), (16, 16), (1, 0)
+        )
+        p_A_22 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 16, 16), (16, 16), (1, 0)
+        )
+        b_Ai_11 = tl.load(p_A_11, boundary_check=(0, 1)).to(tl.float32)
+        b_Ai_22 = tl.load(p_A_22, boundary_check=(0, 1)).to(tl.float32)
+    else:
+        desc = make_tensor_descriptor(A, [T, BT], [H * BT, 1], [16, 16])
+        desc_o = make_tensor_descriptor(Ai, [T, BT], [H * BT, 1], [16, 16])
+        b_Ai_11 = desc.load([i_t * BT + 0, 0]).to(tl.float32)
+        b_Ai_22 = desc.load([i_t * BT + 16, 16]).to(tl.float32)
+
+    # [16, 16]
+    b_Ai_11 = -tl.where(m_A, b_Ai_11, 0)
+    b_Ai_22 = -tl.where(m_A, b_Ai_22, 0)
+
+    for i in range(2, min(16, T - i_t * BT)):
+        b_a_11 = -tl.load(A + (i_t * BT + i) * H * BT + o_i)
+        b_a_11 += tl.sum(b_a_11[:, None] * b_Ai_11, 0)
+        b_Ai_11 = tl.where((o_i == i)[:, None], b_a_11, b_Ai_11)
+    for i in range(16 + 2, min(32, T - i_t * BT)):
+        b_a_22 = -tl.load(A + (i_t * BT + i) * H * BT + o_i + 16)
+        b_a_22 += tl.sum(b_a_22[:, None] * b_Ai_22, 0)
+        b_Ai_22 = tl.where((o_i == i - 16)[:, None], b_a_22, b_Ai_22)
+
+    b_Ai_11 += m_I
+    b_Ai_22 += m_I
+
+    if not USE_TMA:
+        p_A_21 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 16, 0), (16, 16), (1, 0)
+        )
+        b_A_21 = tl.load(p_A_21, boundary_check=(0, 1)).to(tl.float32)
+    else:
+        b_A_21 = desc.load([i_t * BT + 16, 0]).to(tl.float32)
+
+    b_Ai_21 = -tl.dot(
+        tl.dot(b_Ai_22, b_A_21, input_precision=DOT_PRECISION),
+        b_Ai_11,
+        input_precision=DOT_PRECISION,
+    )
+
+    if not USE_TMA:
+        p_Ai_11 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT, 0), (16, 16), (1, 0)
+        )
+        p_Ai_21 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 16, 0), (16, 16), (1, 0)
+        )
+        p_Ai_22 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 16, 16), (16, 16), (1, 0)
+        )
+        tl.store(
+            p_Ai_11,
+            b_Ai_11.to(p_Ai_11.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_22,
+            b_Ai_22.to(p_Ai_22.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_21,
+            b_Ai_21.to(p_Ai_21.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+    else:
+        desc_o.store(
+            [i_t * BT + 0, 0], b_Ai_11.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 16, 0], b_Ai_21.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 16, 16], b_Ai_22.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+
+
+@triton.heuristics({"IS_VARLEN": lambda args: args["cu_seqlens"] is not None})
+@triton.autotune(
+    configs=[
+        triton.Config({}, num_warps=num_warps, num_stages=num_stages)
+        for num_warps in [2, 4, 8]
+        for num_stages in [2, 3, 4, 5]
+    ],
+    key=["H", "BT", "IS_VARLEN"],
+)
+@triton.jit(do_not_specialize=["T"])
+def merge_16x16_to_64x64_inverse_kernel(
+    A,
+    Ai,
+    cu_seqlens,
+    chunk_indices,
+    T,
+    H: tl.constexpr,
+    BT: tl.constexpr,
+    USE_TMA: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+    DOT_PRECISION: tl.constexpr,
+):
+    i_t, i_bh = tl.program_id(1), tl.program_id(0)
+    i_b, i_h = i_bh // H, i_bh % H
+    if IS_VARLEN:
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+    else:
+        bos, eos = i_b * T, i_b * T + T
+
+    o_i = tl.arange(0, 16)
+    m_A = o_i[:, None] > o_i[None, :]
+    m_I = o_i[:, None] == o_i[None, :]
+    A += (bos * H + i_h) * BT
+    Ai += (bos * H + i_h) * BT
+
+    if not USE_TMA:
+        p_A_11 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT, 0), (16, 16), (1, 0)
+        )
+        p_A_22 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 16, 16), (16, 16), (1, 0)
+        )
+        p_A_33 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 32, 32), (16, 16), (1, 0)
+        )
+        p_A_44 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 48, 48), (16, 16), (1, 0)
+        )
+        b_Ai_11 = tl.load(p_A_11, boundary_check=(0, 1)).to(tl.float32)
+        b_Ai_22 = tl.load(p_A_22, boundary_check=(0, 1)).to(tl.float32)
+        b_Ai_33 = tl.load(p_A_33, boundary_check=(0, 1)).to(tl.float32)
+        b_Ai_44 = tl.load(p_A_44, boundary_check=(0, 1)).to(tl.float32)
+    else:
+        desc = make_tensor_descriptor(A, [T, BT], [H * BT, 1], [16, 16])
+        desc_o = make_tensor_descriptor(Ai, [T, BT], [H * BT, 1], [16, 16])
+        b_Ai_11 = desc.load([i_t * BT + 0, 0]).to(tl.float32)
+        b_Ai_22 = desc.load([i_t * BT + 16, 16]).to(tl.float32)
+        b_Ai_33 = desc.load([i_t * BT + 32, 32]).to(tl.float32)
+        b_Ai_44 = desc.load([i_t * BT + 48, 48]).to(tl.float32)
+
+    # [16, 16]
+    b_Ai_11 = -tl.where(m_A, b_Ai_11, 0)
+    b_Ai_22 = -tl.where(m_A, b_Ai_22, 0)
+    b_Ai_33 = -tl.where(m_A, b_Ai_33, 0)
+    b_Ai_44 = -tl.where(m_A, b_Ai_44, 0)
+
+    for i in range(2, min(16, T - i_t * BT)):
+        b_a_11 = -tl.load(A + (i_t * BT + i) * H * BT + o_i)
+        b_a_11 += tl.sum(b_a_11[:, None] * b_Ai_11, 0)
+        b_Ai_11 = tl.where((o_i == i)[:, None], b_a_11, b_Ai_11)
+    for i in range(16 + 2, min(32, T - i_t * BT)):
+        b_a_22 = -tl.load(A + (i_t * BT + i) * H * BT + o_i + 16)
+        b_a_22 += tl.sum(b_a_22[:, None] * b_Ai_22, 0)
+        b_Ai_22 = tl.where((o_i == i - 16)[:, None], b_a_22, b_Ai_22)
+    for i in range(32 + 2, min(48, T - i_t * BT)):
+        b_a_33 = -tl.load(A + (i_t * BT + i) * H * BT + o_i + 32)
+        b_a_33 += tl.sum(b_a_33[:, None] * b_Ai_33, 0)
+        b_Ai_33 = tl.where((o_i == i - 32)[:, None], b_a_33, b_Ai_33)
+    for i in range(48 + 2, min(64, T - i_t * BT)):
+        b_a_44 = -tl.load(A + (i_t * BT + i) * H * BT + o_i + 48)
+        b_a_44 += tl.sum(b_a_44[:, None] * b_Ai_44, 0)
+        b_Ai_44 = tl.where((o_i == i - 48)[:, None], b_a_44, b_Ai_44)
+    b_Ai_11 += m_I
+    b_Ai_22 += m_I
+    b_Ai_33 += m_I
+    b_Ai_44 += m_I
+
+    if not USE_TMA:
+        p_A_21 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 16, 0), (16, 16), (1, 0)
+        )
+        p_A_31 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 32, 0), (16, 16), (1, 0)
+        )
+        p_A_32 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 32, 16), (16, 16), (1, 0)
+        )
+        p_A_41 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 48, 0), (16, 16), (1, 0)
+        )
+        p_A_42 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 48, 16), (16, 16), (1, 0)
+        )
+        p_A_43 = tl.make_block_ptr(
+            A, (T, BT), (H * BT, 1), (i_t * BT + 48, 32), (16, 16), (1, 0)
+        )
+        b_A_21 = tl.load(p_A_21, boundary_check=(0, 1)).to(tl.float32)
+        b_A_31 = tl.load(p_A_31, boundary_check=(0, 1)).to(tl.float32)
+        b_A_32 = tl.load(p_A_32, boundary_check=(0, 1)).to(tl.float32)
+        b_A_41 = tl.load(p_A_41, boundary_check=(0, 1)).to(tl.float32)
+        b_A_42 = tl.load(p_A_42, boundary_check=(0, 1)).to(tl.float32)
+        b_A_43 = tl.load(p_A_43, boundary_check=(0, 1)).to(tl.float32)
+    else:
+        b_A_21 = desc.load([i_t * BT + 16, 0]).to(tl.float32)
+        b_A_31 = desc.load([i_t * BT + 32, 0]).to(tl.float32)
+        b_A_32 = desc.load([i_t * BT + 32, 16]).to(tl.float32)
+        b_A_41 = desc.load([i_t * BT + 48, 0]).to(tl.float32)
+        b_A_42 = desc.load([i_t * BT + 48, 16]).to(tl.float32)
+        b_A_43 = desc.load([i_t * BT + 48, 32]).to(tl.float32)
+
+    b_Ai_21 = -tl.dot(
+        tl.dot(b_Ai_22, b_A_21, input_precision=DOT_PRECISION),
+        b_Ai_11,
+        input_precision=DOT_PRECISION,
+    )
+    b_Ai_32 = -tl.dot(
+        tl.dot(b_Ai_33, b_A_32, input_precision=DOT_PRECISION),
+        b_Ai_22,
+        input_precision=DOT_PRECISION,
+    )
+    b_Ai_43 = -tl.dot(
+        tl.dot(b_Ai_44, b_A_43, input_precision=DOT_PRECISION),
+        b_Ai_33,
+        input_precision=DOT_PRECISION,
+    )
+
+    b_Ai_31 = -tl.dot(
+        b_Ai_33,
+        tl.dot(b_A_31, b_Ai_11, input_precision=DOT_PRECISION)
+        + tl.dot(b_A_32, b_Ai_21, input_precision=DOT_PRECISION),
+        input_precision=DOT_PRECISION,
+    )
+    b_Ai_42 = -tl.dot(
+        b_Ai_44,
+        tl.dot(b_A_42, b_Ai_22, input_precision=DOT_PRECISION)
+        + tl.dot(b_A_43, b_Ai_32, input_precision=DOT_PRECISION),
+        input_precision=DOT_PRECISION,
+    )
+    b_Ai_41 = -tl.dot(
+        b_Ai_44,
+        tl.dot(b_A_41, b_Ai_11, input_precision=DOT_PRECISION)
+        + tl.dot(b_A_42, b_Ai_21, input_precision=DOT_PRECISION)
+        + tl.dot(b_A_43, b_Ai_31, input_precision=DOT_PRECISION),
+        input_precision=DOT_PRECISION,
+    )
+
+    if not USE_TMA:
+        p_Ai_11 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT, 0), (16, 16), (1, 0)
+        )
+        p_Ai_22 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 16, 16), (16, 16), (1, 0)
+        )
+        p_Ai_33 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 32, 32), (16, 16), (1, 0)
+        )
+        p_Ai_44 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 48, 48), (16, 16), (1, 0)
+        )
+        p_Ai_21 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 16, 0), (16, 16), (1, 0)
+        )
+        p_Ai_31 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 32, 0), (16, 16), (1, 0)
+        )
+        p_Ai_32 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 32, 16), (16, 16), (1, 0)
+        )
+        p_Ai_41 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 48, 0), (16, 16), (1, 0)
+        )
+        p_Ai_42 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 48, 16), (16, 16), (1, 0)
+        )
+        p_Ai_43 = tl.make_block_ptr(
+            Ai, (T, BT), (H * BT, 1), (i_t * BT + 48, 32), (16, 16), (1, 0)
+        )
+        tl.store(
+            p_Ai_11,
+            b_Ai_11.to(p_Ai_11.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_22,
+            b_Ai_22.to(p_Ai_22.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_33,
+            b_Ai_33.to(p_Ai_33.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_44,
+            b_Ai_44.to(p_Ai_44.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_21,
+            b_Ai_21.to(p_Ai_21.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_31,
+            b_Ai_31.to(p_Ai_31.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_32,
+            b_Ai_32.to(p_Ai_32.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_41,
+            b_Ai_41.to(p_Ai_41.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_42,
+            b_Ai_42.to(p_Ai_42.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+        tl.store(
+            p_Ai_43,
+            b_Ai_43.to(p_Ai_43.dtype.element_ty, fp_downcast_rounding="rtne"),
+            boundary_check=(0, 1),
+        )
+    else:
+        desc_o.store(
+            [i_t * BT + 0, 0], b_Ai_11.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 16, 16], b_Ai_22.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 32, 32], b_Ai_33.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 48, 48], b_Ai_44.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 16, 0], b_Ai_21.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 32, 0], b_Ai_31.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 32, 16], b_Ai_32.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 48, 0], b_Ai_41.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 48, 16], b_Ai_42.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+        desc_o.store(
+            [i_t * BT + 48, 32], b_Ai_43.to(desc_o.dtype, fp_downcast_rounding="rtne")
+        )
+
+
+@input_guard
+def solve_tril(
+    A: torch.Tensor,
+    cu_seqlens: torch.Tensor | None = None,
+    chunk_indices: torch.Tensor | None = None,
+    output_dtype: torch.dtype = torch.float,
+) -> torch.Tensor:
+    """
+    Compute the inverse of the matrix I + A
+    A should be strictly lower triangular, i.e., A.triu() == 0.
+
+    Args:
+        A (torch.Tensor):
+            [B, T, H, BT], where BT should only be 16, 32, or 64.
+        cu_seqlens (torch.Tensor):
+            The cumulative sequence lengths of the input tensor. Default: `None`.
+        chunk_indices (torch.Tensor):
+            Pre-computed chunk indices. Default: `None`.
+        output_dtype (torch.dtype):
+            The dtype of the output tensor. Default: `torch.float`.
+            If `None`, the output dtype will be the same as the input dtype.
+
+    Returns:
+        (I + A)^-1 with the same shape as A
+    """
+    assert A.shape[-1] in [16, 32, 64]
+    output_dtype = A.dtype if output_dtype is None else output_dtype
+
+    B, T, H, BT = A.shape
+    if chunk_indices is None and cu_seqlens is not None:
+        chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
+    NT = len(chunk_indices) if cu_seqlens is not None else triton.cdiv(T, BT)
+
+    Ai = torch.zeros_like(A, dtype=output_dtype)
+    if BT == 16:
+        merge_fn = solve_tril_16x16_kernel
+    elif BT == 32:
+        merge_fn = merge_16x16_to_32x32_inverse_kernel
+    elif BT == 64:
+        merge_fn = merge_16x16_to_64x64_inverse_kernel
+
+    merge_fn[(B * H, NT) if BT == 64 else (NT, B * H)](
+        A=A,
+        Ai=Ai,
+        cu_seqlens=cu_seqlens,
+        chunk_indices=chunk_indices,
+        T=T,
+        H=H,
+        BT=BT,
+        USE_TMA=is_tma_supported,
+        DOT_PRECISION=FLA_TRIL_PRECISION,
+    )
+    return Ai

bundle/plugin-site/ornith_g256/optimized/pipeline/candidate/wy_fast.py

@@ -0,0 +1,158 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# SPDX-FileCopyrightText: Songlin Yang, Yu Zhang
+#
+# This file contains code copied from the flash-linear-attention project.
+# The original source code was licensed under the MIT license and included
+# the following copyright notice:
+# Copyright (c) 2023-2025, Songlin Yang, Yu Zhang
+
+# ruff: noqa: E501
+
+import torch
+
+from vllm.triton_utils import tl, triton
+
+from .index import prepare_chunk_indices
+
+
+@triton.heuristics({"IS_VARLEN": lambda args: args["cu_seqlens"] is not None})
+@triton.autotune(
+    configs=[
+        triton.Config({}, num_warps=num_warps, num_stages=num_stages)
+        for num_warps in [2, 4, 8]
+        for num_stages in [2, 3, 4]
+    ],
+    key=["H", "K", "V", "BT", "BK", "BV", "IS_VARLEN"],
+)
+@triton.jit(do_not_specialize=["T"])
+def recompute_w_u_fwd_kernel(
+    k,
+    v,
+    beta,
+    w,
+    u,
+    A,
+    g,
+    cu_seqlens,
+    chunk_indices,
+    T,
+    H: tl.constexpr,
+    Hg: tl.constexpr,
+    K: tl.constexpr,
+    V: tl.constexpr,
+    BT: tl.constexpr,
+    BK: tl.constexpr,
+    BV: tl.constexpr,
+    IS_VARLEN: tl.constexpr,
+):
+    i_t, i_bh = tl.program_id(1), tl.program_id(0)
+    i_b, i_h = i_bh // H, i_bh % H
+    if IS_VARLEN:
+        i_n, i_t = (
+            tl.load(chunk_indices + i_t * 2).to(tl.int32),
+            tl.load(chunk_indices + i_t * 2 + 1).to(tl.int32),
+        )
+        bos, eos = (
+            tl.load(cu_seqlens + i_n).to(tl.int32),
+            tl.load(cu_seqlens + i_n + 1).to(tl.int32),
+        )
+        T = eos - bos
+    else:
+        bos, eos = i_b * T, i_b * T + T
+    p_beta = tl.make_block_ptr(
+        beta + bos * H + i_h, (T,), (H,), (i_t * BT,), (BT,), (0,)
+    )
+    p_g = tl.make_block_ptr(g + (bos * H + i_h), (T,), (H,), (i_t * BT,), (BT,), (0,))
+    p_A = tl.make_block_ptr(
+        A + (bos * H + i_h) * BT, (T, BT), (H * BT, 1), (i_t * BT, 0), (BT, BT), (1, 0)
+    )
+    b_beta = tl.load(p_beta, boundary_check=(0,))
+    b_A = tl.load(p_A, boundary_check=(0, 1))
+    b_g = tl.exp(tl.load(p_g, boundary_check=(0,)))
+
+    for i_v in range(tl.cdiv(V, BV)):
+        p_v = tl.make_block_ptr(
+            v + (bos * H + i_h) * V,
+            (T, V),
+            (H * V, 1),
+            (i_t * BT, i_v * BV),
+            (BT, BV),
+            (1, 0),
+        )
+        p_u = tl.make_block_ptr(
+            u + (bos * H + i_h) * V,
+            (T, V),
+            (H * V, 1),
+            (i_t * BT, i_v * BV),
+            (BT, BV),
+            (1, 0),
+        )
+        b_v = tl.load(p_v, boundary_check=(0, 1))
+        b_vb = (b_v * b_beta[:, None]).to(b_v.dtype)
+        b_u = tl.dot(b_A, b_vb, allow_tf32=False)
+        tl.store(p_u, b_u.to(p_u.dtype.element_ty), boundary_check=(0, 1))
+
+    for i_k in range(tl.cdiv(K, BK)):
+        p_k = tl.make_block_ptr(
+            k + (bos * Hg + i_h // (H // Hg)) * K,
+            (T, K),
+            (Hg * K, 1),
+            (i_t * BT, i_k * BK),
+            (BT, BK),
+            (1, 0),
+        )
+        p_w = tl.make_block_ptr(
+            w + (bos * H + i_h) * K,
+            (T, K),
+            (H * K, 1),
+            (i_t * BT, i_k * BK),
+            (BT, BK),
+            (1, 0),
+        )
+        b_k = tl.load(p_k, boundary_check=(0, 1))
+        b_kb = (b_k * b_beta[:, None] * b_g[:, None]).to(b_k.dtype)
+        b_w = tl.dot(b_A, b_kb)
+        tl.store(p_w, b_w.to(p_w.dtype.element_ty), boundary_check=(0, 1))
+
+
+def recompute_w_u_fwd(
+    k: torch.Tensor,
+    v: torch.Tensor,
+    beta: torch.Tensor,
+    g_cumsum: torch.Tensor,
+    A: torch.Tensor,
+    cu_seqlens: torch.Tensor | None,
+    chunk_indices: torch.Tensor | None = None,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    B, T, Hg, K, V = *k.shape, v.shape[-1]
+    H = v.shape[-2]
+    BT = A.shape[-1]
+
+    if chunk_indices is None and cu_seqlens is not None:
+        chunk_indices = prepare_chunk_indices(cu_seqlens, BT)
+    NT = triton.cdiv(T, BT) if cu_seqlens is None else len(chunk_indices)
+    BK = 64
+    BV = 64
+    u = torch.empty_like(v)
+    w = k.new_empty(B, T, H, K)
+    recompute_w_u_fwd_kernel[(B * H, NT)](
+        k=k,
+        v=v,
+        beta=beta,
+        w=w,
+        u=u,
+        A=A,
+        g=g_cumsum,
+        cu_seqlens=cu_seqlens,
+        chunk_indices=chunk_indices,
+        T=T,
+        H=H,
+        Hg=Hg,
+        K=K,
+        V=V,
+        BT=BT,
+        BK=BK,
+        BV=BV,
+    )
+    return w, u